Repository navigation
Conversation
`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
left a comment
There was a problem hiding this comment.
Thanks for tracing this; pos 2 is a real tie. Two gaps in the new rule, inline.
| 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}"); | ||
| } |
There was a problem hiding this comment.
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?
| /// 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. |
There was a problem hiding this comment.
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>
|
两条都在 1. 互相 top-8 不是平局(
规则的两半各有一个不需要设备的单测(这个文件里唯一不 2. greedy walk 翻之后比的是不同文本( 为此 读数三条单测在 pod 上(CPU 路径,commit 0.0625 / 0.1875 对 1.0 通过(余量 5–16×),门禁自己的线是 2.0。文档里那句 provenance(0.0625 / 0.1875)由这次运行重新核过。 |
FeathBow
left a comment
There was a problem hiding this comment.
Thanks, both points are in, RowGap reads well, and the new unit tests are worth keeping. Three follow ups, inline.
| /// 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; |
There was a problem hiding this comment.
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.
| 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", |
There was a problem hiding this comment.
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?
| /// 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`. |
There was a problem hiding this comment.
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>
…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>
FeathBow
left a comment
There was a problem hiding this comment.
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.
| /// 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( |
There was a problem hiding this comment.
This repeats continue_greedy except for where each step's token comes from; one function taking an optional walk to follow would do.
| 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 |
There was a problem hiding this comment.
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>
|
Fixed( 验证:三个单测、 那两个 TileLang gate 跑不了:这个 build 没有生成 kernel(门禁自己断言拒绝),而且 |
FeathBow
left a comment
There was a problem hiding this comment.
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; |
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
Hmm. This part still needs verification.
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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.
`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
left a comment
There was a problem hiding this comment.
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) |
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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() { |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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).
| /// 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 |
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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>
What
serve_oracle::compare_rowrequired an identical argmax between two arms whose raw-logit drift the same callers bound by a calibrated ceiling — 2.0 inincremental_serving_matches_recompute— and whose differenceneutral_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: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 ownk): 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-repogemma4-12b-hf-golden.safetensors, this commit's own code: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
incremental_serving_matches_recomputefailing deterministically on one box, at single rank. The engine was not at fault, and the gate is untouched by fix(gemma4): drive the TP ranks concurrently so long prompts do not stall #1133, so the fix belongs here rather than there.lane_gates_logprobs, via fix(gemma4): build the scored-prompt gate's prompt from real text #1132); this is the same move forcompare_row.