Skip to content

test(gemma4): hold a moved pick to a near tie, not to equality - #1136

Open
sysuTobo wants to merge 8 commits into
pegainfer-project:mainfrom
sysuTobo:fix/gemma4-near-tie-picks
Open

sysuTobo wants to merge 8 commits into
pegainfer-project:mainfrom
sysuTobo:fix/gemma4-near-tie-picks

Conversation

@sysuTobo

@sysuTobo sysuTobo commented Oct 8, 2026

Copy link
Copy Markdown
Contributor

What

serve_oracle::compare_row required an identical argmax between two arms whose raw-logit drift the same callers bound by a calibrated ceiling — 2.0 in incremental_serving_matches_recompute — and whose difference neutral_scale, the function right above it, measures at 0.31 to 5.75 raw logits on the same rows. The rule cannot hold at depth, and on a clean L20 it did not:

prefill pos 2 ours:   top5 [(594, 13.6875), (497, 13.625 ), (236747, 13.5625), 236745, 575]
prefill pos 2 theirs: top5 [(497, 13.875 ), (594, 13.6875), (236747, 13.6875), 236745, 575]

The same five ids on both sides and the same tail; the top two 0.0625 apart in one row and 0.1875 in the other, the rows agreeing to 0.3125 everywhere else. That is reduction-order drift across a near tie, not a moved distribution — and it is deterministic on that box (green elsewhere), i.e. a rule that fails wherever the top two happen to sit inside the drift.

Change

The rule is now the one the rest of the repo applies (lane_gates_tp::distribution_gap, and the HF fixtures with their own k): a flipped pick must be a near tie — each pick inside the other row's top-NEAR_TIE_TOP_K (8). A pick outside that is a different distribution and still fails, and both rows are printed when it happens. One function, so all six call sites get it, the replacement-kernel gates included.

Receipts

12B, one L20, PEGAINFER_GEMMA4_GOLDEN = the in-repo gemma4-12b-hf-golden.safetensors, this commit's own code:

prefill pos 2: near-tie flip 594 vs 497
prefill pos 2: max |dlogit| 0.3125   pos 4: 1.03125   pos 8: 1.21875
prefill: 9 positions, max |dlogit| 1.21875      (the gate's own ceiling is 2.0)
test result: ok. 1 passed; 0 failed; 0 ignored

the four decode steps after the walk pass too. The replacement-kernel siblings (the_replacement_global_kernel_matches_the_incumbent, the_folded_pool_matches_the_split_one) need a TileLang build, which this container does not have, so they were not run here — they take the same rule through the same function.

Where this came from

`serve_oracle::compare_row` asserted an identical argmax between two arms whose
raw-logit drift the same callers bound by a calibrated ceiling (2.0 here), and
whose difference `neutral_scale` right above measures at 0.31..5.75 on the same
rows. That cannot hold at depth, and it did not: on the `short` prompt the
incremental gate stopped at `prefill pos 2`, where the decode row's top two sat
0.0625 apart (594 at 13.6875 against 497 at 13.625) and the recomputed row's
0.1875 (497 at 13.875 against 594 at 13.6875) — the same top five ids on both
sides, agreeing to 0.3125 everywhere else, i.e. reduction-order drift across a
near tie rather than a moved distribution. With the repo's own rule the gate
runs to the end: nine positions, max |dlogit| 1.21875 against its 2.0 ceiling,
then the four decode steps.

So the rule is now the fleet-wide one — a flipped pick must be a near tie, each
pick inside the other row's top-`NEAR_TIE_TOP_K`; a pick outside that is a
different distribution and still fails, and both rows are printed when it
happens. Six call sites share the one function, so the replacement-kernel gates
get the same rule.

Signed-off-by: bohaotu <bohaotu@pony.ai>

@FeathBow FeathBow left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for tracing this; pos 2 is a real tie. Two gaps in the new rule, inline.

Comment on lines 666 to +677
let (a, b) = (argmax(ours), argmax(theirs));
assert_eq!(a, b, "{what}: argmax diverged ({a} vs {b})");
if a != b {
describe_row(&format!("{what} ours"), ours);
describe_row(&format!("{what} theirs"), theirs);
assert!(
pick_inside_top(theirs, a, NEAR_TIE_TOP_K) && pick_inside_top(ours, b, NEAR_TIE_TOP_K),
"{what}: a pick outside the other row's top-{NEAR_TIE_TOP_K} ({a} vs {b}), which is a \
different distribution rather than the reduction-order drift the callers' ceiling \
covers"
);
eprintln!("{what}: near-tie flip {a} vs {b}");
}

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Mutual top 8 does not mean a tie: [6, 0, -20, ...] against [0, 6, -20, ...] passes here, under the decode gates' 12.0. Could a flip require the two picks to sit within a small fixed width of each other in both rows?

Comment thread pegainfer-gemma4/src/serve_oracle.rs Outdated
Comment on lines +648 to +652
/// The two arms run different launch shapes or kernels (one-token decode against
/// a whole-prompt prefill, an incumbent against its replacement), which this
/// engine does not promise bit-equal: the callers bound the raw-logit drift by a
/// calibrated ceiling, and `neutral_scale` right above measures what that
/// difference is worth on the same rows — 0.31 to 5.75 raw logits.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The two kernel gates walk each arm greedily, so after a flip the later rows compare different text. Could they stop at the first flip, as distribution_gap does?

…both top-ks

Mutual top-k membership let a flip through when the two rows preferred different
tokens by six nats: each pick being the other row's runner-up says nothing about
how far apart they read, and the callers' 12.0 drift line does not see it
either. The rule now also requires the two picks to sit within NEAR_TIE_WIDTH of
each other in each row's own reading — a difference between two entries of one
row, so the same quantity as the TP gates' 1.0 on a shared token's gap, against
the 0.0625 and 0.1875 a real flip reads.

The two kernel gates that walk each arm greedily now stop at the row where a
pick moves: each arm fed its own pick to the next step, so every later pair of
rows compares two different contexts — the same reason distribution_gap stops at
a flip — and each cell prints how many of its rows were compared.

The rule's halves carry unit tests that need no device, since membership alone
passing a six-nat flip is exactly what a test has to hold down.

Signed-off-by: bohaotu <bohaotu@pony.ai>
The membership half is what `lane_gates_tp`'s `distribution_gap` and the
fixtures' `top1_share` floor hold; the width is added in this file, where a
raw-logit flip inside a 12.0 line can sit six nats wide. Claiming the width as
the repo's existing rule would have been wrong about those two callers.

Signed-off-by: bohaotu <bohaotu@pony.ai>
@sysuTobo

sysuTobo commented Oct 8, 2026 •

Copy link
Copy Markdown
Contributor Author

两条都在 1d1be75a(规则 7f6a8ce7,第二条只是把文档改成说清"哪一半是本文件新加的")。

1. 互相 top-8 不是平局(serve_oracle.rs:677)—— 补上了量级那一半。 你的反例正是漏掉的那块:[6, 0, …] 对 [0, 6, …],两边 pick 都是对方的次高,成员判据通过,而两行各自偏好不同的 token 差 6 nat —— 调用方的 12.0 漂移线也看不到它(谓词在比较量级之前就短路了)。现在规则是两半:

  • 每个 pick 在对方的 top-NEAR_TIE_TOP_K 内(原判据);
  • 且两个 pick 在每一行自己的读数里相距不超过 NEAR_TIE_WIDTH(两行都读)。

NEAR_TIE_WIDTH = 1.0 的来由:这个 margin 是同一行两个 entry 之差,所以在 raw logits(这里)与 log_softmax 上是同一个量,而 lane_gates_tp 的 LOGBROB_LINE 对"两轮都保留的同一个 token 的 logprob 差"给的也是 1.0 —— 同一个量、同一个数、同一条理由。真实那次翻转读 0.0625 / 0.1875,你的反例读 6。

规则的两半各有一个不需要设备的单测(这个文件里唯一不 #[ignore] 的测试):6 nat 的翻转必须失败、只在其中一行是平局也必须失败、宽度内的翻转要被报告为翻转。

2. greedy walk 翻之后比的是不同文本(serve_oracle.rs:652)—— 按 distribution_gap 的做法在首次翻转处停下。 两个 kernel 门禁(the_replacement_global_decode_matches_the_incumbent、the_folded_pool_matches_the_split_one)现在把翻转那一行的 gap 先折进读数再 break;每个 cell 的打印行加上 "N of 16 rows compared",这样"提前停下"在读数里是可见的 —— 否则它读起来像量完了整条 walk。incremental_serving_matches_recompute 与 neutral_scale 不动:它们每一行都有自己的 recompute,比较对象始终是同一段文本,翻转不会让后面的行变成别的东西。

为此 compare_row 改返回 RowGap { worst, flipped }:要"翻没翻"的调用方从已经算出它的那一处取,而不是自己再 argmax 一遍(同一事实两个真相源)。

读数

三条单测在 pod 上(CPU 路径,commit 7f6a8ce7)全 ok。incremental_serving_matches_recompute(12B,一张 L20,仓内 PEGAINFER_GEMMA4_GOLDEN)这次把那次翻转本身打了出来,283 s,ok:

prefill pos 2 ours:   top5 [(594, 13.6875), (497, 13.625), (236747, 13.5625), …]
prefill pos 2 theirs: top5 [(497, 13.875), (594, 13.6875), (236747, 13.6875), …]
prefill pos 2: near-tie flip 594 vs 497, margins 0.0625 / 0.1875
prefill: 9 positions, max |dlogit| 1.21875

0.0625 / 0.1875 对 1.0 通过(余量 5–16×),门禁自己的线是 2.0。文档里那句 provenance(0.0625 / 0.1875)由这次运行重新核过。

@FeathBow FeathBow left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks, both points are in, RowGap reads well, and the new unit tests are worth keeping. Three follow ups, inline.

Comment thread pegainfer-gemma4/src/serve_oracle.rs Outdated
Comment on lines +650 to +655
/// The margin is a difference between two entries of one row, so it is the same
/// quantity on raw logits (here) as on log_softmax: `lane_gates_tp`'s
/// `LOGBROB_LINE` reads the same 1.0 on a shared token's gap for the same
/// reason. The flip that motivated the rule read 0.0625 and 0.1875, and the
/// shape the width rejects reads 6.
const NEAR_TIE_WIDTH: f32 = 1.0;

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LOGBROB_LINE bounds one token's gap between two runs; this bounds two tokens' gap within one row, so its calibration does not carry over. Could the line rest on the measured 0.0625 / 0.1875 and the 6 it must reject, and could the replacement and folded pool gates run once under it? Only the incremental gate has so far.

Comment thread pegainfer-gemma4/src/serve_oracle.rs Outdated
Comment on lines +900 to +912
for (i, (a, b)) in incumbent.iter().zip(&replacement).enumerate() {
let gap = compare_row(a, b, &format!("prompt {prompt} at {len}, decode step {i}"));
if gap > worst.0 {
worst = (gap, i);
let row = compare_row(a, b, &format!("prompt {prompt} at {len}, decode step {i}"));
compared = i + 1;
if row.worst > worst.0 {
worst = (row.worst, i);
}
if row.flipped {
break;
}
}
eprintln!(
"prompt {prompt} at {len} tokens: floor 0, neutral scale {scale}, \
replacement |dlogit| {} at step {}",
replacement |dlogit| {} at step {}, {compared} of {STEPS} rows compared",

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

continue_greedy returns the prefill row plus STEPS decode rows, so a full walk prints 17 of 16, and in the folded pool gate a flip at the prefill row reports a decode gap of 0 with no decode row compared. Decoding the replacement on the incumbent's walked tokens would compare every row and avoid both; otherwise could the count and the row 0 case be fixed?

Comment thread pegainfer-gemma4/src/serve_oracle.rs Outdated
Comment on lines +637 to +640
/// How far a moved pick may sit from the other arm's: inside the other row's
/// top-`NEAR_TIE_TOP_K` is reduction-order noise, outside it is a different
/// distribution. The fleet-wide rule — `lane_gates_tp`'s `distribution_gap` and
/// the HF fixtures apply the same shape, with the fixture's own `k`.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This still calls top 8 reduction order noise, and DRIFT_LINE's doc still says the argmax holds on every compared row; neither matches the rule now. Could both describe the current rule?

…the incumbent's tokens

`continue_greedy` returns the prompt's row plus STEPS decode rows, so a full cell
printed 17 of 16 and a flip at the prompt row left a cell with no decode row
measured. Both go away when the second arm is fed the first arm's tokens
(`continue_forced`): every row compares the same text, nothing is skipped, and the
count printed per cell is how many rows moved a pick — which is what
`RowGap::flipped` is kept for.

The width's doc no longer leans on `LOGBROB_LINE`'s calibration (that bounds one
token's gap between two runs; this bounds two tokens' within one row). It rests
on the flips this rule has read — 0.0625 and 0.1875 — and the 6 it must reject.
The top-8 and DRIFT_LINE docs now describe the rule that is actually there.

Signed-off-by: bohaotu <bohaotu@pony.ai>
sysuTobo added a commit to sysuTobo/pegainfer that referenced this pull request Oct 9, 2026
…them

NCCL's abort waits for a captured graph that references the communicator being
aborted, and a rank stopped inside a collective cannot release its own — its
graphs live in an arena its own thread holds, and that release is exactly what
the abort is waiting for. With a captured collective in play the failure path has
no order that works: the abort waits for the release, the release runs after the
join, and the join waits for the rank that abort was to release. The same cycle
runs against rank 0's graphs when an extra rank fails first. Graphs are refused
at a tensor-parallel launch rather than left to hang there, and capture is a
measured wash under TP (146.0 vs 143.7 tok/s at c8), so nothing priced above zero
is given up.

Two smaller items from the same round. Rank 0's device is bound before any peer
is spawned: a failure there would otherwise return from the scope while the peers
were already running, and they would wait on a collective rank 0 never issued
while the abort sat behind the join. And the prefix gate counts the rows whose
top-8 share no token — they contribute nothing to `worst`, so a widespread error
could pass it on one shared row elsewhere; a quarter of the rows is the bound.

The width half moves here from pegainfer-project#1136: distribution_gap requires the two picks
within NEAR_TIE_WIDTH in both runs' own readings, with the rationale corrected —
the folded worst_pick against LOGBROB_LINE already catches the six-nat shape, so
the width's own catch is the band between them (margins 0.9 and 1.05 with folds
of 0.975), which a unit test now fixes.

Signed-off-by: bohaotu <bohaotu@pony.ai>
@sysuTobo
sysuTobo requested a review from FeathBow October 9, 2026 14:43

@FeathBow FeathBow left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks, the forced walk fixes the row count and the row 0 case. The replacement decode and folded pool gates changed shape here and are the ones that hold NEAR_TIE_WIDTH at depth; the incremental gate's short prompt does not stand in for them, and CI skips the ignored GPU tests, so one run of each would close it. Two cleanups inline.

Comment thread pegainfer-gemma4/src/serve_oracle.rs Outdated
/// this arm's, so both arms read the same text on every row and nothing a
/// comparison does can leave one behind. That is what lets the two kernel gates
/// compare *every* row of the walk instead of stopping where a pick moves.
fn continue_forced(

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This repeats continue_greedy except for where each step's token comes from; one function taking an optional walk to follow would do.

Comment thread pegainfer-gemma4/src/serve_oracle.rs Outdated
serve.tilelang_global_attn = true;
let (replacement, _) = continue_greedy(&ctx, &serve, &tokens, STEPS);
let mut worst = (0.0f32, 0usize);
// The replacement walks the incumbent's own tokens, so every row — the

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The forced walk is explained here, in continue_forced and in both gate docs, and the six nat example in NEAR_TIE_TOP_K, NEAR_TIE_WIDTH and compare_row; once each is enough, and NEAR_TIE_WIDTH can drop the LOGBROB_LINE paragraph. RowGap's doc still says each arm feeds its own pick, which no caller does now.

…rule

`continue_greedy` and `continue_forced` differ only in where each step's token
comes from, so they become one `continue_walk` taking an optional walk to follow;
the forced call sites follow the incumbent's tokens with `Some`.

The rule is explained once, where it is applied: `NEAR_TIE_WIDTH` drops its
`LOGBROB_LINE` paragraph (its calibration is a different quantity), `NEAR_TIE_TOP_K`
drops the six-nat example, and `compare_row` keeps the rule, the counterexample
and the 0.0625 / 0.1875 reading the width was set from. `RowGap`'s doc no longer
claims each arm feeds its own pick, which no caller does now.

Verified on the sm_89 pair: the three device-free units, the incremental serving
gate (12.2 s) and the prefix-restore gate (17.0 s) all pass. The two TileLang
gates this reshaping touches still cannot run here — this build carries no
generated kernel, and the family needs sm_90.

Signed-off-by: bohaotu <bohaotu@pony.ai>
@sysuTobo

sysuTobo commented Oct 10, 2026 •

Copy link
Copy Markdown
Contributor Author

Fixed(f9a73754):continue_greedy/continue_forced 合成一个 continue_walk(…, follow: Option<&[u32]>),forced 侧传 Some;解释收敛到一处 —— NEAR_TIE_WIDTH 删掉 LOGBROB_LINE 那段,NEAR_TIE_TOP_K 删掉六 nat 例子,规则/反例/0.0625-0.1875 只留在 compare_row,RowGap 的 doc 改成 "with one arm walking and the other following its tokens (continue_walk)"。

验证:三个单测、incremental_serving_matches_recompute(12.2 s)、prefix_restore_matches_cold_path(17.0 s)全过。

那两个 TileLang gate 跑不了:这个 build 没有生成 kernel(门禁自己断言拒绝),而且 build.rs 里 GEMMA4_TILELANG 写着 min_sm: 90,我没有 sm_90 硬件。

@sysuTobo
sysuTobo requested a review from FeathBow October 10, 2026 07:45

@FeathBow FeathBow left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks, this reads much better. One small request inline.

/// from are in [`compare_row`]. It transfers between the raw-logit gates here and
/// `lane_gates_tp`'s log-softmax ones because the margin is a difference inside
/// one row.
const NEAR_TIE_WIDTH: f32 = 1.0;

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This 1.0 comes from one flip in a nine token prompt, and in #1133 the same 1.0 fails at 1024 tokens. The two TileLang gates run at up to 3000 tokens and have not run yet. Could the comment here say the width is unverified for them?

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hmm. This part still needs verification.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

To put it more directly, the question is why 1.0 is not checked at that depth. The incumbent walk and neutral_scale need no TileLang kernel, and they already pass every row through compare_row at 512, 1500 and 3000 tokens. Could that part run on your card, so the width is measured where these gates use it?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ran that part on an L20: the walk and neutral_scale over the sweep cells. neutral_scale's comparison — decode row against whole-prompt recompute — breaks the rule at depth: at prompt 0 / 512 tokens the flipped pick leaves the other row's top-8, and on the golden's edge text the two picks sit 3.5 apart against the width of 1.0. That comparison is a printed magnitude, not one the gates judge, so 133931e splits the two: neutral_scale now reads a raw gap and asserts nothing, and the rule stays in compare_row on the gates' own two arms. The part that runs without TileLang is committed as the_drift_floor_is_measured_over_the_decode_sweep — all nine cells on one GPU, floors 0.5 to 21.75.

The width at 512/1500/3000 is only exercised inside the two TileLang gates, and I have no access to an sm_90a card, so I cannot run them here. NEAR_TIE_WIDTH's doc now says plainly that the width is unverified at the sweep's depths.

@FeathBow FeathBow self-assigned this Oct 11, 2026
`NEAR_TIE_WIDTH`'s provenance is one flip on the nine-token `short` prompt,
but the two generated-kernel gates sweep 512/1500/3000 tokens and neither
has run, so the width is unverified there. Running the part that needs no
generated kernel does not measure it either: the incumbent walk's rows
through `neutral_scale` at the sweep's first cell trip the rule's membership
half before the width — a flipped pick outside the other row's top-8 — on
the cycled fixture prompts and on the golden's own `edge` text alike
(`max |dlogit| 3.5` on the latter). The width is not what binds at that
depth; a depth gate has to start from the membership.

Signed-off-by: bohaotu <bohaotu@pony.ai>
`neutral_scale` prints the magnitude an accepted implementation difference
reaches on the rows the decode gates sweep, but it computed that through
`compare_row`, so it applied the near-tie rule as well — to a comparison the
gates do not judge (an incumbent decode row against a recompute of the same
context). At the sweep's depths that comparison's flips leave the other row's
top-8, so `the_replacement_global_decode_matches_the_incumbent` and
`the_folded_pool_matches_the_split_one` panicked inside their own floor print,
before reaching the arms they exist to compare. The rule stays on the arms a
gate compares (`compare_row`); the floor reads a raw gap now.

`the_drift_floor_is_measured_over_the_decode_sweep` runs that half where no
generated kernel is needed. Measured on two L20s over the nine cells at
512/1500/3000: 0.5 to 21.75 raw logits — 21.75 at prompt 0 / 512, and one cell
above `DRIFT_LINE` itself — so the doc's sixteen-cell 0.31..5.75 no longer
describes this sweep.

Signed-off-by: bohaotu <bohaotu@pony.ai>

@FeathBow FeathBow left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for running the sweep, and the width doc now says plainly where 1.0 is unverified. The two TileLang decode gates still have not run, so the width at 512, 1500 and 3000 tokens stays open. Three points inline.

ours.iter()
.zip(theirs)
.map(|(x, y)| (x - y).abs())
.fold(0.0f32, f32::max)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

f32::max drops a NaN when the other operand is a number, so raw_gap(&[1.0], &[f32::NAN]) reads 0.0. Now that neutral_scale calls this instead of compare_row, a NaN recompute row reads as no drift, and the new test's is_finite check passes on it. Could raw_gap keep the finite check that compare_row has?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in 24c4db2: raw_gap now carries the same finiteness assert compare_row has, plus a unit test that raw_gap(&[1.0], &[f32::NAN]) panics instead of reading 0.0. Ran it on an L20 — the assert fires at serve_oracle.rs:671 and the three row-rule units still pass.

/// reads are on the record in [`neutral_scale`]'s doc.
#[test]
#[ignore = "requires a Gemma 4 checkpoint, a GPU, and --test-threads=1"]
fn the_drift_floor_is_measured_over_the_decode_sweep() {

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These nine cells never reach compare_row, so they check neither the top 8 membership nor the 1.0 margin. A green run here says the floor is finite and nothing about the width. The width at these depths still needs one run of each TileLang gate.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Agreed, and the gate's doc now says it: a green run reads a finite floor and nothing about the width — these nine cells never reach compare_row, so neither the top-8 nor the 1.0 is exercised; the width at these depths takes a run of each TileLang gate (24c4db2).

Comment thread pegainfer-gemma4/src/serve_oracle.rs Outdated
/// calibrated ceiling, and `neutral_scale` right above measures what that
/// difference is worth on the same rows — 0.31 to 5.75 raw logits.
///
/// A pick that moves inside that drift is therefore the arithmetic, not this

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Passing the margin does not show that a moved pick is arithmetic, so "is therefore the arithmetic" says more than the check does. The 0.31 to 5.75 here also predates the sweep that neutral_scale now cites at 0.5 to 21.75. Why the floor measures without the rule is explained four times, in NEAR_TIE_WIDTH, raw_gap, neutral_scale and the new test. Could it stay in neutral_scale only?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

All three in 24c4db2: the "is therefore the arithmetic" claim is gone (passing the ceiling does not show the flip was arithmetic), the 0.31-to-5.75 quote is dropped, and the why-measure-without-the-rule explanation stays in neutral_scale only — NEAR_TIE_WIDTH, raw_gap and the floor gate point there.

…n one place

`f32::max` returns the other operand for a NaN, so `raw_gap` — the helper the
previous commit introduced for the floor's measurement — folded a NaN row to
zero and read as no drift, and the new floor gate's `is_finite` check would
have passed on it. It now asserts finiteness the way `compare_row` does, with
a unit test pinning that a NaN row does not read as no drift.

Three doc corrections from the same review. `compare_row` no longer says a pick
inside the callers' ceiling "is therefore the arithmetic" (the magnitude does
not settle that) and no longer quotes the retired 0.31-to-5.75 sweep, which
the floor now cites at 0.5 to 21.75. Why the floor measures without the rule
was explained four times, in `NEAR_TIE_WIDTH`, `raw_gap`, `neutral_scale` and
the floor gate; it stays in `neutral_scale`, and the other three point there.
The floor gate's doc now says what a green run does not show: its nine cells
never reach `compare_row`, so the width at those depths still needs a run of
each generated-kernel gate.

Verified on an L20: `cargo fmt --check` clean, and the four row-rule units pass
with the new guard firing at serve_oracle.rs:671.

Signed-off-by: bohaotu <bohaotu@pony.ai>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

gemma4 gemma4 model crate (pegainfer-gemma4)

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants