[None][perf] Emission-assisted GVR top-K decode for the DeepSeek V4 indexer - #16953
[None][perf] Emission-assisted GVR top-K decode for the DeepSeek V4 indexer#16953siyidNV wants to merge 121 commits into
Conversation
…ase-3 collect Port the block-skip consumer from the skip-finegrain development chain onto the R0 (op#26) architecture, measured-optimal configuration only (grain 32, int16 active list = 16KB SMEM, strided coalesced 3-barrier build, UN=2 software-pipelined compact scan). The active-block list is built once per row at the loosest rung (lossless for every rung count and the collect); phase3's compact stream-write replays the count pass's per-thread walk order so prefix-sum positions stay exact; a list-current flag pairs the two and is cleared on any dense fallback. Misaligned slice starts fall through to the dense path. Contract: workspace/epilogue_topk_interface.md. Correctness: REAL-data smoke (flash 16k/64k/256k/1024k + pro 64k/1024k, dense + skip arms, exactness contract) all pass incl. hit_rate=0.08. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The block-skip bounds tensor initially landed as a required positional in __call__, breaking every pre-existing compile path that does not pass it (16457's equivalence tests: 'Missing required argument'). Move it after stream with a None default so legacy callers are untouched; the wrapper passes it positionally last (the TVM-FFI env-stream launch takes no runtime stream arg). Verified standalone: main equivalence family 768 passed / 0 failed, launch_autoconfig 4/4, real-data smoke (flash+pro up to 1M) all green. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The active list previously required 32-aligned slice starts (runtime
guard falling back to dense), which silently disabled the skip on every
cluster slicing whose N/cs is not a multiple of 32 — including the
launch policy's cs=8 picks at the 512k/1024k rungs. The list now covers
FULL blocks only (first-full-block ceil in the build); the sub-block
head region of an unaligned slice is counted by all threads in a
strided scalar pass ordered BEFORE the list walk, and Phase 3's compact
write replays the same head-then-list per-thread order, so prefix-sum
positions stay exact. Boundary blocks shared with a neighbouring CTA
appear only in that CTA's head region — no double count, no gap.
Verified: skip arms exact at cs in {1,2,4,8} on flash+pro real cells
(64k/512k/1024k, unaligned N=262127 and 131075 slices).
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The active list was built at the loosest rung over all M columns. On
low-hit-rate rows the lowest sample-quantile rung sits far below the
K-th value (real pro 1024k: retained fraction 0.63 at that rung vs
0.08 at the final threshold), so the list barely skipped anything.
Build now iterates (cs==1): if the list exceeds CAP = 3/4 kC blocks,
DROP that rung — dropping is always correct (the rung is merely an
unmeasured probe; its partial counts are excluded from the admission
argmin and the fallback bracket seeding via a dropped-rung mask) — and
rebuild at the next tighter threshold, bounded by M-1 extra ~2-5us
builds. At cs>1 per-CTA list lengths differ (the drop decision would
diverge across the cluster), so the plain loosest-rung build stays.
Real-data 1024k, cs1, warm-L2 directional: pro 28.8 -> 16.4us (skip
ratio 1.16x -> 2.05x), flash 18.5 -> 12.8us (1.73x -> 2.49x).
Correctness: cs {1,2,4,8} x flash+pro x {64k,512k,1024k} all exact.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
…loads Review fixes for the compact machinery: - skip_ok now also requires nb_slice <= SKIP_MAX_BLOCKS and absolute block id < 32768 (int16 list entries); wider/higher slices fall back to the dense walk losslessly instead of silently truncating counts and the collect (or wrapping ids negative at cluster_size > 1). - Both compact walks vector-load only fully in-bounds chunks; the slice-end straddle goes through the scalar path, so an unaligned row no longer reads past the row/allocation. - Ctor rejects enable_block_skip without enable_r0 (dead 16KB SMEM). - emu_block_max defaults to records='positional' (the shipped kernel is grain 32; 'rotate' is a grain-128 fold fixture) and the wrapper asserts block_max covers every 32-position record of the row. Validated: capacity boundary (8192/8193/9375 blocks), N=1.05M at cs=4/8, unaligned exact-size rows (N=1000/65535/65529), real-data flash/pro 64k/1024k — all exact with planted tail winners. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The compact walk only wins on long rows (cold-L2 protocol: >= 2.18x at N=262k, 4-14% loss at N <= 131k). Gate block_max shape-based (no device sync) behind skip_min_n=200_000: below it the wrapper drops to the dense arms. Protocol after gating: flash/pro 256k/512k cells all 0.99-1.02x, 1024k wins intact (flash BS1 15.63us 2.18x, BS1024 4.18x; pro BS1024 3.15x). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
pick_config gains has_block_max: with bounds available and N >= 200k the policy pins cluster_size = 1 - the compact list + rung tightening (cs1-only) beat the row-split configs outright once the bounds prune the scan (cold protocol, real data: BS1 1.21x, BS64 2.12x, BS1024 4.16x vs the stock picks; splitting shrinks each CTA slice below the skip break-even and disables tightening). The wrapper dispatch gains a second gate next to skip_min_n: K > 512 at num_rows < 8 keeps the stock path - the acceptance band is proportionally tighter (kC/K = 6 vs 10), the bounds prune less, and the row-split configs win (pro 262k BS1: skip 21.3-21.6us at cs1/cs8 vs stock cs8 19.7us). Cold protocol vs op26 at its own launch policy, 24 cells: zero regressions; flash 1024k 1.21/2.12/4.16x (BS 1/64/1024), pro 1024k 1.68/3.15x (BS 64/1024), everything below the gates identical to stock. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
emu_seed_counts / emu_cand implement the A-side products per the epilogue<->topk buffer contract v2 so the consumer waterfall can be developed and tested against torch references before the fused indexer lands. Real-data coverage probe: 7/8 cells have a seed count inside [K, kC]; prev-kth drifts too loose at long context, so the L2 collect threshold should be the middle rung / xstate-adaptive. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
use_ext_counts: rung thresholds AND their exact counts arrive from the
indexer epilogue (seed_thr/seed_counts [rows, 3], interface v2), so
P1b and the M-ary R0 count pass are skipped. The seeded refine routes
both cases: an in-band rung is re-measured once (building the
per-thread hand-off Phase 3 requires) and accepted; a full miss seeds
log-falsi with the external brackets. cs==1, requires fb_fix and a
3-slot rung config (the wrapper pins 2 qfracs + vseed; the values are
irrelevant since P1b never runs).
Real-data validation (flash/pro x 16k..1024k, thresholds {prev-kth,
q35, q85} + emu counts): 8/8 exact on stock/ext/ext+skip arms incl.
the pro-1M full-miss bracket cell. Directional: +9-10% at 1M (P1b +
M-count saved), small-N slightly negative (the waterfall routes those
to the L2 direct path instead). Next: skip P1 under ext (outer
brackets from xstate) and the L2 direct-to-P4 branch.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
With external epilogue counts the only surviving P1 products are the [v_lo, v_hi] outer bracket and the scalar-state init; the ext rungs provide the bracket directly (host contract: t_0 < t_2, finite, all rows valid) and tid0 initializes the scalars. A miss whose target falls outside [t_0, t_2] recovers through the refine loop's 8x bracket expansion, same as the stock fail-soft. Real data 8/8 exact unchanged; directional gains vs stock R0 improve to 1.15x at flash 256k / 1.14x at flash 1M (from 1.04x/1.09x with P1 still running). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
v1 routed external-count admission through the dense seeded refine and forfeited the compact-walk win (flash 1M ext 34.7us vs skipR0 15.6us cold). v2 only skips P1b: the stock M-ary pass runs on the ext rungs, so list build, rung tightening, per-thread hand-off and classify compose unchanged. When an ext count is already in [K, kC] the admitted threshold is parked in ALL rung slots (v2b) — the M-ary pass degenerates to one compact single-threshold count and classify admits it; a miss keeps the distinct rungs as measured brackets. Real data 8/8 exact (stock/ext/ext+skip). Cold protocol: flash 64k 1.21x/1.25x (BS1/BS1024, the slim-admission cell); flash 1M ext+skip 16.1us vs skipR0 15.6us (composition recovered); pro-1M miss rows still pay multi-count+refine (0.7x) — routing sends those to stock skipR0 via xstate feedback (next step). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
emit_xstate writes the per-row loop state (interface v2 layout [rows, 8]: [0] valid, [1] kth proxy, [2] accepted threshold, [3] cand_count from the pre-P4 snapshot — P4 repurposes the s_iscalars slots) at the cs==1 Phase-4 exit; degenerate identity rows write valid=0. The next step derives its seed rung group from these fields. Real-data validation: exactness unchanged; state fields exact (cand_count == count_ge(threshold): flash 991/633, pro 1854/2354); same-step reseed from the written state admits in-band with the exact count on ALL cells — including pro 1M, whose static rung group missed entirely (the temporal rung fixes the miss AND slims flash 1M admission 1290 -> 633). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
use_ext_cand: epilogue-collected (value, index) pairs land straight in
smem_keys/vals via a cooperative sentinel-skipping SMEM-atomic load —
no P1, no counting, no Phase-3 scan. Eligibility (void == 0, claimed
<= cand_cap, collect rung count in [K, kC]) is a CTA-uniform register
predicate, so the dynamic skip of the P2/P3 slab stays convergent;
ineligible rows fall through to the ext-counts path unchanged.
Real data: 6 cells x {ext, l2, forced-void fallback} all exact. The
direct path makes top-k O(cand_count), independent of N: eligible rows
cost ~12.2us warm from 16k through 1M (flash 1M: 2.84x vs the ext
count path, below even the cold skipR0 15.6us). With the 0.89-0.97
chain in-band rates, ~90% of production rows hit this floor when the
epilogue emits cand.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
derive_seed_rungs places the next step's guard rungs a fixed number of count-OCTAVES from the previous accepted threshold, using the local slope of log2(count) vs threshold estimated from the previous step's own 3 rung measurements (log-linearity is the same property log-falsi exploits). Fixed spreads face a two-sided trap: too narrow misses drift, too wide puts the guard rungs themselves out of band — no single value wins both models (best fixed: pro 0.97/flash 0.92 vs flash-tuned 0.82/0.97). Real-chain kernel validation (V4-Pro/Flash multi-step captures): in-band admission pro 0.89 -> 0.99, flash 0.96, all steps exact. Combined with the L2 direct arm this routes ~97%+ of production rows to the O(cand_count) floor (cold protocol: flash 1M BS1024 30.9us = 12.1x, BS1 8.7us = 4.0x; pro 256k 1.5-1.8x). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The ext knobs were compile-time, so a row whose epilogue rungs all miss [K, kC] (or an xstate-invalid row, t_0 = +FLT_MAX) still paid the ext bracket-refine — measurably worse than stock (pro-1M cold: ext-miss 51us vs stock-skip 21us). Routing is now a per-row runtime predicate read from the ext counts themselves (CTA-uniform loads, so the dynamic branches with barriers inside stay convergent): in-band rows keep the ext fast path (skip P1 + P1b), miss/invalid rows run the full stock path (P1 + P1b + vseed + count) including the block-skip machinery. Warm validation: pro 1M ext+skip 43.5us (0.77x) -> 18.8us (1.80x); in-band cells unchanged (flash 1M 2.40x, pro 256k 1.07x); mixed-row chains exact with in-band 0.99/0.97 (pro/flash, adaptive rungs). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The BS=1 mid/long-row cells stayed on op26 because the waterfall fast paths were cs==1-only while op26's pick_config splits a single row across cs=4/8 CTAs. The pre-collected pairs are O(cand_count), so row splitting buys the direct path nothing: at cs > 1 the LEADER loads the pairs alone (take_cand is cluster-uniform - every CTA reads the same per-row control words) and peers publish zero local candidates for the DSMEM gather; ineligible/invalid rows fall through to the native stock path at op26's own cluster split. xstate writes at the leader's Phase-4 exit; the ext count pass composes with the existing cs>1 cluster merge unchanged. Validation: bl2 cells exact at cs=1/4/8 including forced-void fallback; cs1 smokes and the adaptive-rung chains unchanged (in-band 0.99). Cold protocol with the production arm (op26 launch config + ext inputs + in-kernel routing), vs op26 baseline: flash 256k BS1/64 1.24x/1.59x, flash 512k 1.33x/1.42x, flash 1M BS64 3.69x, pro 256k 1.20-1.73x - the former regression cells flip to wins; pro 512k/1M BS1 static-rung misses route to stock (adaptive xstate rungs take them direct in the closed loop, 1.4-1.6x steady-state). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
Collect the pre-collected candidate list at the LOOSEST seed rung and admit it whenever it is complete (claimed <= K_max) and any rung counts >= K; the filter rung (count closest to K from above) is applied on the fly while loading the pairs, so P4 sees the thinnest covering set. kC leaves the admission vocabulary and remains only as the physical smem capacity guard. Correctness: C(t_lo) >= C(t_filt) >= K implies true top-K subset of list subset of filtered set; list truncation (claim order is value-blind) remains the only fatal case and falls back. - K_max = 24576, set by a four-chain search on real captures (incl. 320k/640k long decode): 16K->24K gains 8pp direct-hit rate, 24K->32K only 0.1pp (band-limited, not capacity-limited). - Loader: 4x-unrolled latency-overlapped walk with ballot-batched smem claims (loop exit must stay warp-uniform: ragged exits deadlock the warp collectives) and un-nested value loads. Device-level (nsys kern-sum, cold L2) on a 160k real chain: the naive walk ran 0.64x vs the block-skip arm; this form reaches 1.05x at full loosest-rung coverage (eligibility 1.00). - Straddle refine (cs=1): when no rung count lands in [K, kC] but the list is complete, one 256-bin histogram pass over the list finds an in-band edge and the filtered load proceeds; smem overflow demotes to the fallback. 640k chain: straddle steps 30 -> 14-16us, device mean 1.41x -> 1.73x vs block-skip. - Byte-parity routing keeps fat lists (2*claimed*cs >= N) on the fallback: measured both ways, walking them is slower at every cs. Validation (B200): four admission modes exact (direct/filter/refine/ fallback) at cs=1 and 18/18 exact at cs=1/4/8 on real V4 bundles; four real decode chains (160k/132k/320k/640k) all-step exact with wall ratios 0.97/1.00/1.04/1.21x and device-level 1.05x (160k) / 1.73x (640k) vs the block-skip arm. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
Replace the rung-based in-list-filter admission with the count-only scheme: the candidate list is SoA (score column + position column, sentinel score -inf), collected at a single loose line, and admitted purely by entry count (K + 64 <= claimed <= K_max = 24576; the 64 is the emitter sentinel-pad bound, so the live count provably covers K). Rung admission, filter-line selection, straddle refine and the parity gate are all deleted - the seed-count columns are no longer consumed on the list path (the GEMM-side L1 pass becomes deletable, -3.2% emission tax). - THIN list (fits kC): every entry lands AT ITS LIST INDEX in the candidate buffers - no ballots, no smem atomics (128 serialized same-address atomics per trip measured ~1.1us/1k entries), no warp-uniform loop constraint. - FAT list: atomic-free copy of the score column into a dedicated 96KB smem region (sentinels sanitized to t_lo - 1), a zooming smem histogram (3 rounds, NBL^3 resolution - value-linear bins collapse on long-tailed logits) finds an edge whose exact count lands in [K, kC] (lands ~1030 for K=1024), survivors compact with one merged-ballot atomic per warp per trip. The vals slots carry LIST INDICES (no second cold gmem pass over the position column); a post-P4 repair swaps the K winners' positions with fully-parallel gathers. - Closed loop: xstate[1] publishes the exact k-th (output slot K-1 of the rank-ordered scatter), xstate[2] the ~3K-crossing anchor from the round-0 histogram. Host policy picks the anchor field per domain (tight k-th for short/stable rows, wide 3K edge for volatile long rows - the exact-k-th anchor alone shrinks the next down-guard target to 4K and slope noise then undershoots K, forcing ~26us fallbacks). GVR_P4_TAIL_DBG compiles per-phase clock64 stamps into the spare xstate slots. Validation (B200): 24/24 exact across cs=1/4/8 and the straddle- threshold suite on real V4 bundles; per-row cold device phases: thin walk 1.5-3us, fat stage+zoom+compact ~1.1us/1k entries, Phase 4 flat 5.5-6.5us. Real-chain device-level vs the block-skip arm: 640k 1.42x (fallback steps are C(t_lo) < K undershoots - a host anchor-policy matter), 160k 0.93x. Kernel-only chain means trade 5-20% vs the previous rung-based commit at B=8 in exchange for the interface collapse; the deleted L1 emission pass dominates E2E at large batch. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The emitter (indexer GEMM epilogue; emulated host-side in the bench
harness) now counts the two tighter lines while writing the SoA list -
two extra compares per EMITTED element only, against the full-row L1
pass this replaces - and the control words widen to {n0, void, n1,
n2}. The topk side enters with every count known and the whole list
path collapses to a scalar state machine:
- some line's count lands in the acceptance band [K, B*]: cut at the
TIGHTEST such line, ONE filtered gmem pass straight into the
candidate buffers (positions deferred as list indices; the position
column is gathered only for the K winners after Phase 4). Counts
and load predicates are the same comparison, so line cuts need no
overflow net at all.
- the band is straddled or overshot by every line: a zooming
histogram over the gmem list CLAMPED between the two known bracket
lines finds an in-band edge (narrow domain - no long-tail bin
collapse; the all-above case takes one max pass first).
- void, or n0 < K + 64 (the emitter sentinel bound, proving live
coverage of K): fallback.
The dedicated smem staging region is deleted (frees 96-128KB; the
kernel's smem drops back to the pre-list footprint), and B* / kC
become constructor knobs (accept_cap, kc_override) for the band
search. Closed loop publishes the exact k-th (rank-ordered output
slot K-1) and the loosest in-band line as the anchor.
Line placement is a searched host policy (derive_seed_lines_v4):
count targets (t0, t1, t2), grid-searched on real chains =
(4096, 3584, 1536) for short domains / (12288, 5120, 2048) for long;
physical kC stays 5120 (8192 measured no gain).
Validation (B200): five admission states each exercised exact
(hit-t2/hit-t1/bracketed-histogram/above-t2-histogram/fallback),
24/24 exact at cs=1/4/8; real-chain device-level vs the block-skip
arm: 160k 1.09x (first config of this lineage to beat the rung-based
1.05x), 640k 1.39x (residual: 2-3 volatile rows/step whose collection
count escapes any placement - a host anchor-policy iteration item).
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
…d histogram Emitter writes the candidate list into three fixed segments (>=t2 / [t1,t2) / [t0,t1), caps B*/B*/rest, spill to the looser segment on overflow), so a line cut only ever reads the dense mapped prefix of the segments above it: the hit path becomes a pure copy (no value filter, no ballots, no atomics) and the histogram path walks mapped indices. When all three lines overshoot B*, the bracket segment's own prefix doubles as an unbiased sample: the histogram runs on it at the sample rate with 1.25x-scaled fire targets, and the exact post-load count net absorbs the sampling noise. Device-level cold-chain results vs the block-skip arm (B=1): flash 132k 1.58x, pro 160k 1.57x, pro 320k 1.54x, pro 640k 1.98x (fastest steps 8-14us). Exactness smokes pass for cs=1/4/8 including forced straddle/void routings. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
…repair Sub-phase clock64 instrumentation (GVR_P4_SUB_DBG) showed the P4 rank-scatter core costs only ~0.6us/k candidates; the chain-observed ~1.9us/k came from the exact-tail boundary repair: the tiny-tie fast path ran an O(need x class) serial select on thread0 (~10us on real rows with need ~100 x class ~100), and bigger classes re-scanned every candidate per radix level behind ~20 block barriers. The repair is now: (1) a block-wide pure-tie check over the straddle class (bit-equal class needs no repair at all - the scatter's arrival fill is already value-set exact); (2) mixed classes are compacted IN PLACE into smem_keys/vals[0..class) with a register-buffered two-phase pass (warp-aggregated slot claims), so every later step scales with the class, never the candidate count; (3) class <= 128 takes an exact warp0 pairwise-rank rewrite, larger classes a block-parallel 4-level MSB radix over the compacted pairs with a warp0 shuffle-scan digit search (3 block barriers per level instead of 5). The full-candidate radix fallback is gone from the fast-tail variant. Device-level cold chains, 640k B=1: step mean 13.7 -> 12.6us (1.97x -> 2.13x vs block-skip; the previously slowest window improves 17.35 -> 12.6us as five 19-21us serial-repair victims drop to 9.3-10.8us); 640k B=8 19.5us (1.60x). Exactness: 18/18 microbench cells including forced tie/outlier stressors, smokes cs=1/4/8 plus forced straddle/void routings all bit-exact. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
New self_scan mode: the kernel itself streams the row ONCE against the
three closed-loop seed lines and buckets candidates on the fly - no
external emitter, no indexer-side changes, no gmem candidate values.
One CTA per row, four phases: (0) scan-bucket - VALUES land in on-chip
segments (A/B/C at bases 0/B*/2B*, values-only 4B/entry, spill to the
looser segment, cursor totals ARE the line counts), POSITIONS stream to
a write-only gmem column reusing the cand_idx slot; (1) the v5 cut
state machine unchanged (a line cut compacts winning segment runs to
the smem prefix and fills smem_vals with segment coordinates, so P4,
the tail repair and the deferred K-gather run verbatim); ineligible
rows take the stock in-kernel fallback.
Scan-loop lessons baked in (each measured): per-element warp ballots
serialize every load (~1.8us/k); 16-wide register lists spill at 1024
threads (64 regs/thread ceiling) - values re-read from the load
fragments, positions derived arithmetically, classes recomputed;
warp-collective claim prefixes cap in-flight loads at 2/warp (ncu:
0.19% memory throughput) - final form claims passers with per-element
smem atomics, which do not synchronize the warp and hide under the
read stream (0.13us/k comp).
Exactness: 25-cell REPORT-S4 dataset x B in {1,2,4,8} = 100/100
bit-exact (flash/pro/v32 incl. K=2048, tiny-N and straddle fallbacks).
Perf vs PR16457 tip (same node, cold kernel-sum): geomean 0.64-0.71x,
short rows 0.8-1.05x, long rows 0.4-0.8x - the single-CTA read wall by
design; stage 2 (block-max skip) attacks the read itself.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
…gle band
Phase 0 gains a block-skip variant (enable_block_skip + self_scan):
per-32-position maxima from the GEMM tail gate whole blocks out of the
scan. Measured design pivot: skipping against the LOOSE collection
line can never pay (n0/N ~5-12% density -> ~80-99% of blocks contain a
passer; benched 0.16-0.39x), so the skip mode collects a SINGLE BAND
against the TIGHTEST line (density 0.4-0.8% -> 12-22% pass): only
segment A fills, the cursor keeps exact attempt counts, and the v5
state machine runs unchanged fed n0 == n1 == n2 - a cut lands on t2
(common), the sample-hist path absorbs over-B* rows (the A prefix
stays a value-blind sample), under-K rows take the stock fallback.
The small-batch block_max gate in the wrapper is bypassed for
self_scan (stage 2 owns its own skip economics).
Exactness: 25-cell REPORT-S4 x B in {1,2,4,8} = 100/100 bit-exact,
plus forced under-K fallback cells. Perf state (B=1 vs PR16457 tip,
same node): long rows improve markedly over the dense scan (flash
512k 33.9 -> 24.5us = 0.87x of tip; 1024k 46.7 -> 38.7; pro 1M 52 ->
44) while short/mid rows should route to the dense scan (host picks
by expected block pass rate). Known remaining work, measured and
documented: the block loop is still latency-bound on the bmax stream
(8 scalar loads/warp round); a lane-per-block + ballot variant was
tried and loses at high pass rates - loop shape per density regime is
the open optimization, along with a t2-only closed-loop line-derive
for chains.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The skip scan is restructured into two passes that both run at the tuned dense-loop shape: (1) DENSE-vector-scan the block-max array itself (1/32 of the row, 128-bit vectors - the bmax row base is only 16B aligned) and compact the PASSING BLOCK IDS into the idle C segment (single-band mode never fills C; ids store exactly as floats); (2) walk the compact list, eight listed blocks per warp round issued back-to-back - every element read is useful and the loads pipeline. A list overflow (pass rate too high for skipping to ever pay) falls back to a dense full scan of the row inside the same phase. This removes both latency walls the one-pass shapes hit (8-scalar bmax rounds; serial per-block walks): flash 512k drops 25 -> 20us and BEATS the PR16457 tip (1.06x) - first cell where the fused self-contained kernel wins outright; flash 1024k 47 -> 26us (0.72x of tip), pro 1M 53 -> 36us, v32 128k+ 28us. Dense/skip best-of geomean 0.69 -> 0.73-0.75x across the 25-cell REPORT-S4 dataset, all 100 cells bit-exact. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
…ases The pass-2 gather interleaved each block load with its smem atomic claims; atomics are memory-ordered, so the compiler could not overlap the next block load and the eight-block round degenerated into a serial latency chain (phase-0 stamp: 17.7us at flash-1024k against a ~5us budget). Loading all eight listed blocks into registers first and claiming afterwards restores the in-flight parallelism: phase 0 drops to 10.4us and the 25-cell table moves decisively - flash 512k 1.39x over the PR16457 tip, 1024k parity (19us), 256k 0.93x; v32 64k parity, 128k+ 0.90-0.92x; pro 1M 0.85x. Dense/skip best-of geomean 0.73 -> 0.84-0.86x, still 100/100 bit-exact. Remaining gap concentrates in the mid-row dense regime (64-128k, 0.63-0.72x), where the dense scan's single-CTA latency wall stands (cp.async staging is the known next lever). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
Preload the next round's two vectors into shadow fragments before the current round's atomic claims (the pass-2 lesson applied to the dense loop). Measured neutral-to-slightly-positive (phase-0 38.7 -> 37.3us at flash-1024k): unlike pass 2 the dense loop's wall is not the cross-round atomic ordering - documented for the record; the next dense-lane lever is cp.async/smem staging. Final 25-cell state (dense/skip best-of vs PR16457 tip, B=1..8 geomean 0.84-0.86x, 100/100 bit-exact): flash 512k 1.39x / 1024k 1.00x / 256k 0.93x; v32 64k 1.00x / 128k+ 0.90x; pro 1M 0.85x; remaining gap concentrated at the 64-128k dense regime. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
Replace the register-preload dense scan with an LDGSTS staging pipeline: each thread streams one 16B vector per step into a private slot-major smem slot (no data registers, no scoreboard stall until the wait), keeping stage_slots rounds in flight. The staging buffer aliases smem_vals - written only after phase 0, with every non-empty cp.async group drained inside the loop - so depth 2 costs zero smem; trimming cap_c to <= 16384 frees 32KB of keys for depth 4. flash-1024k phase0 37.3 -> 34.8us; exactness unchanged (fused and skip smoke 6/6, 25-cell real-data sweep 50/50 bit-exact). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The short-row 512-thread heuristic is tuned for the stock multi-pass
kernel; under self_scan it silently halved the warp count of every
N_dec < 65536 cell and cost ~5us/cell in the phase-0 scan (flash-128k
p0 14.4 -> 9.4us at 1024 threads). Route self_scan to 1024 threads
unconditionally.
25-cell x B{1,2,4,8} same-node sweep vs PR16457 tip: best-of geomean
0.838 -> 0.864 (B8 0.881), 100/100 bit-exact.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
The dense scan is instruction-issue bound (pcsamp: no_instructions +
fixed-latency wait dominate; long_scoreboard is 6%), so each pipeline
step now processes two 16B vectors per thread - loop, wait, commit and
address arithmetic amortize over 8 elements while the in-flight byte
count stays at 2 pairs x 32B across the 4 staging slots.
The pair shape needs all 4 slot rows, and the 64KB staging fits the
CTA budget only with the C segment trimmed, so self_scan now defaults
cap_c to 16384 and rejects anything larger (validated bit-exact across
the 25-cell x B{1,2,4,8} sweep).
flash p0 (warm, 1024 threads): 512k 17.8 -> 16.8us, 1024k 33.5 ->
32.4us; smoke 6/6 exact.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
…in-kernel) New ext_rungs mode: the host supplies only the three closed-loop rung THRESHOLDS (previous-step xstep lines); the kernel counts them itself through the stock R0 multi-count pass and admits the tightest rung with count in [K, kC], then collects and refines as usual. This is the fully self-contained two-pass shape: no emission of any kind, pass 1 = one fused 3-rung count (cluster-merge and block-skip compose unchanged), pass 2 = the stock single-line collect. Versus use_ext_counts (variant A) the only delta is where the counts come from; P1's preIdx gather and the P1b quantile rung derivation are both skipped (the seed lines carry the bracket). Smoke: 15/15 bit-exact across cs=1/4, block_max skip, and the thin (all rungs below K) and fat (all counts above kC) miss paths. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
|
CI is green on This PR should now be ready to merge pending the JIRA ticket number in the title. |
| # reused slots cold-start the emission closed loop; stale | ||
| # lines only mis-place cuts - counts are re-measured | ||
| # in-kernel, so exactness never rides on this reset | ||
| self.top_k.reset_gvr_emission_rows( |
There was a problem hiding this comment.
This resets only the current prefill rows. When an earlier generation finishes, the remaining requests can compact to different generation slots without passing this block, while _gvr_emission_state.xstate remains positional. That gives a live request another request's finite thresholds. Please remap or invalidate emission state when generation-slot occupants change, and cover continuous-batching turnover that shifts an active request between slots.
There was a problem hiding this comment.
Thanks - you're right that the state is positional and that compaction bypasses this block. I dug into what a stale row can actually cause, and ran the scenario:
Exactness does not ride on the lines. The consumer's admission test reads cand_ctl, which the emitter writes this step (claimed/n1/n2 are atomic counts of what actually crossed each line on the current query), not any host-side prediction:
void_c == 0 and claimed_c >= top_k + 64 and claimed_c <= list_cap
A stale line only changes how many candidates get collected. Too tight and claimed < K+64 fails the gate, so the row falls back to the in-kernel full scan; too loose and the list is larger than needed and the cut line is chosen from the measured counts anyway. Either way the selection is exact - which is also why xstate-invalid rows (cold start) are safe: they park onto the stock path by construction.
Adversarial check on B200 (batch 4, N=131072, K=2048): warm the closed loop until every slot carries finite lines, then hand each slot a different request's score distribution with no reset (weights re-drawn at 3x scale, so the stale lines are badly mismatched):
warm lines=[123.9, 61.0, 114.8, 144.6] claimed=[2701, 18440, 10742, 3210] exact
post-churn lines=[ 83.2, 98.3, 131.8, 118.4] claimed=[92861, 100057, 4016, 43392] exact
post-churn+1 claimed=[6709, 5706, 4541, 6171] exact
Top-K is exact at every step; the mismatched lines show up as inflated candidate counts (92K collected where ~3K is typical) and the closed loop re-converges the next step. So the cost of slot turnover is a transient perf dip, not a wrong result.
That said, I agree the transient is worth removing and that relying on this reasoning is fragile. The clean fix is to key the emission state to request identity the same way the prior would need to be, since gvr_prior_indices is positional in exactly the same way and has the same turnover exposure. I'd rather do that as a follow-up covering both pieces of state together, with a continuous-batching turnover test, than bolt a partial remap onto this PR. Does that work for you? If you'd prefer it in this PR I can add an invalidate-on-turnover hook here instead.
There was a problem hiding this comment.
Thanks for digging in — I walked the same path in gvr_topk_decode.py, and the admission really is count-driven: usable gates on void_c == 0 plus claimed_c measured this step, and the cut is picked from n1_c/n2_c against [K, B*], with the clamped-histogram bracket underneath. A mismatched line moves how much gets collected, not what gets selected, so I agree this is a transient collection cost rather than a wrong Top-K. Combined with the path being opt-in behind TRTLLM_GVR_EMISSION, I'm happy to let it go.
A follow-up that keys both xstate and gvr_prior_indices to request identity together sounds better than a partial remap here — they share the same positional exposure, and splitting them would leave the weaker half in place. Please do include the continuous-batching turnover test there.
Not blocking this PR.
There was a problem hiding this comment.
Thanks for walking the kernel path yourself - agreed on all points.
Follow-up scope, so it doesn't get lost: key xstate (and the rest of the emission closed-loop state) plus gvr_prior_indices to request identity in one change, with a continuous-batching turnover test that shifts an active request between generation slots and asserts both exactness and that the closed loop does not inherit a foreign row.
I also pushed 40917644 here: a unit test asserting reset rows park on non-finite lines while untouched slots keep their state, so the invariant this discussion relies on is now covered in-tree.
Continuous-batching turnover can hand a request another request's positional emission state. Exactness never rides on the seed lines - the consumer admits on the candidate counts the emitter measures for the current query, so a stale line only changes how many candidates are collected - but a reset row must park on non-finite lines so it restarts on the stock path instead of inheriting thresholds. Verified on B200 alongside this: after handing every warm slot a different request's score distribution with no reset, top-K stays exact and the closed loop re-converges within one step (candidate counts spike to ~92K, back to ~6K next step). Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
|
/bot run |
|
PR_Github #68625 [ run ] triggered by Bot. Commit: |
|
PR_Github #68625 [ run ] completed with state
|
|
/bot run |
|
PR_Github #68675 [ run ] triggered by Bot. Commit: |
|
PR_Github #68675 [ run ] completed with state |
| - unittest/_torch/attention --ignore=unittest/_torch/attention/sparse/test_cute_dsl_fp8_paged_mqa_logits.py --ignore=unittest/_torch/attention/sparse/test_cute_dsl_fp4_paged_mqa_logits.py --ignore=unittest/_torch/attention/sparse/test_cute_dsl_gvr_topk_decode.py | ||
| - unittest/_torch/attention/sparse/test_cute_dsl_fp8_paged_mqa_logits.py | ||
| - unittest/_torch/attention/sparse/test_cute_dsl_fp4_paged_mqa_logits.py | ||
| - unittest/_torch/attention/sparse/test_cute_dsl_gvr_topk_decode.py TIMEOUT (120) |
There was a problem hiding this comment.
Does it require nearly 120 mins to complete in the CI testing? It doesn't sound a proper testing in the pre-merge.
CC @yuxianq
There was a problem hiding this comment.
No, according to @longcheng-nv 's measurement, it should be ~15 mins.
There was a problem hiding this comment.
I measured it rather than guess: 45m22s wall clock for the whole file (261 tests, shared B200 node), so the 120 cap is closer to right than generous - dropping it to ~20 would make CI flake.
The cost is dominated by CuTe DSL kernel compilation, not execution. ext_list[hit] and ext_closed_loop[tier_shape0] are ~117s each, while their same-group siblings finish in 0.01s once that config is already compiled. So it is a compile-heavy test rather than a two-hour-running one. Happy to split the file if the wall clock itself is the concern.
There was a problem hiding this comment.
@siyidNV Before this PR it is ~15mins, which means you add 30 mins JIT to this file. @longcheng-nv has taken a lot of effort to reduce the JIT time of this file, please also try to reduce the JIT time of new tests to just cover necessary cases.
There was a problem hiding this comment.
Fixed in a29d92e - only the file that needs its own timeout is split out now:
- unittest/_torch/attention --ignore=.../test_cute_dsl_gvr_topk_decode.py
- unittest/_torch/attention/sparse/test_cute_dsl_gvr_topk_decode.py TIMEOUT (120)
On the timeout itself: I measured the file at 45m22s wall clock (261 tests, shared B200), so the 120 cap is closer to right than generous. The cost is CuTe DSL kernel compilation rather than execution - ext_list[hit] and ext_closed_loop[tier_shape0] are ~117s each while their same-group siblings finish in 0.01s once that config is compiled.
| #!/usr/bin/env python3 | ||
| # ruff: noqa | ||
| # Measurement harness committed verbatim for provenance; bench idioms | ||
| # (loop-scoped buffers, del/rebind) trip static analysis. | ||
| # f58: B×N 四张表(flash/pro × 算数均值/最小值),逐步配对 vs PR16457, | ||
| # 我们的臂按线上路由(plan_emission)选取。 |
There was a problem hiding this comment.
Should we cleanup these AI-generated Chinese comments here (and probably other lines)?
There was a problem hiding this comment.
Done in a29d92e - the Chinese comments and print strings in f58_an.py and perstep.py are now English, and I swept the whole folder for non-ASCII while I was there (em dash, multiplication sign, element-of, arrow were used as typography in English prose). All six files in the folder are ASCII-clean now.
| - unittest/_torch/attention --ignore=unittest/_torch/attention/sparse/test_cute_dsl_fp8_paged_mqa_logits.py --ignore=unittest/_torch/attention/sparse/test_cute_dsl_fp4_paged_mqa_logits.py --ignore=unittest/_torch/attention/sparse/test_cute_dsl_gvr_topk_decode.py | ||
| - unittest/_torch/attention/sparse/test_cute_dsl_fp8_paged_mqa_logits.py | ||
| - unittest/_torch/attention/sparse/test_cute_dsl_fp4_paged_mqa_logits.py |
There was a problem hiding this comment.
Why is --ignore necessary here? The ignored cases are added back immediately.
There was a problem hiding this comment.
You're right - two of the three were ignored and added straight back with identical semantics, so only the file that needs its own timeout is split out now (a29d92e):
- unittest/_torch/attention --ignore=.../test_cute_dsl_gvr_topk_decode.py
- unittest/_torch/attention/sparse/test_cute_dsl_gvr_topk_decode.py TIMEOUT (120)
There was a problem hiding this comment.
Early feedback: I think we should clean up this script folder to remove non-unicode characters and formalize these perf scripts before merge.
There was a problem hiding this comment.
Done in a29d92e.
- Non-ASCII: gone from all six files in the folder.
- SPDX headers: the four perf scripts were missing them; added.
# ruff: noqa: narrowed. It was only ever warranted inab_steps.py, where the explicitdelof the per-grid buffers at the end of each loop body makes every closure reference read as possibly-unbound - that file now pinsF821and says why.f58_an.pyandperstep.pywere hiding an unused import andlas a variable name behind the same blanket escape; both are fixed properly and now pass ruff clean.
Reviewer feedback on the committed measurement harnesses: - Translate the remaining Chinese comments and print strings to English (f58_an.py, perstep.py). - Drop every non-ASCII character from the folder: em dash, multiplication sign, element-of and arrow were used as typography in English prose. - Add the SPDX header the four perf scripts were missing. - Narrow the blanket `# ruff: noqa`. It is only warranted in ab_steps.py, where the explicit `del` of the per-grid buffers at the end of each loop body makes every closure reference read as possibly-unbound; that file now pins `F821` and says so. f58_an.py and perstep.py were hiding an unused import and `l` as a variable name behind the same escape - both are fixed instead, and both files now pass ruff clean. - Stop ignoring two test files from the attention sweep only to add them straight back; only the file that needs its own timeout is split out. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
|
@longcheng-nv - merging main brought #18094 in, and this branch now fails 5 tests Where the two changes meet. #18094 rewrote the retry-shrink inside
Questions:
What I already tried, so nobody repeats it: hand-splicing the conflict hunks Separately, on the test-list timeout question above: I measured the whole file at |
| LIST_PARK_LINE = 1.0e30 | ||
|
|
||
|
|
||
| class GvrEmissionState: |
There was a problem hiding this comment.
This file only contains one class GvrEmissionState, and GvrEmissionState is only used by tensorrt_llm/_torch/modules/top_k.py, should we move it into top_k.py instead? Or move it into tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/.
There was a problem hiding this comment.
Moved to cute_dsl_kernels/blackwell/top_k/ in d9fc773 - your second option.
Reasoning for that over folding it into top_k.py: the constants are the kernel's buffer-geometry contract (the segment bases at 0 / LIST_SEG_A / 2*LIST_SEG_A have to match the CUTLASS emitter exactly), and gvr_routing.py - which picks the tier those buffers get emitted for - already lives there and was already being imported across the package boundary. Putting 272 lines of kernel geometry into a torch module felt like the wrong direction. The move also drops the attention_backend -> cute_dsl_kernels reverse dependency; the import is now local to the kernel package.
Import sites updated: modules/top_k.py and three in test_cute_dsl_gvr_topk_decode.py.
There was a problem hiding this comment.
This tests/scripts/cute_dsl_kernels/top_k/perf/ seems like an experiment dir, should we remove it instead? You can keep your perf experiment scripts in local.
There was a problem hiding this comment.
Agreed - removed the whole tests/scripts/cute_dsl_kernels/top_k/perf/ directory in d9fc773. You're right that it belongs in a local tree; the numbers it backs are already in the commit messages and the PR description.
This also supersedes the cleanup I did on those same files in a29d92e (translating comments, stripping non-ASCII, adding SPDX, narrowing the ruff escape) - deleting them is the better answer to the question three of you were circling.
One file kept deliberately: tests/scripts/cute_dsl_kernels/top_k/run_gvr_topk.py, one level up. It is a standalone driver plus a parameterised pytest correctness sweep over dtype x K x N x seed x next_n, not an experiment script. Happy to move that too if you'd rather it lived elsewhere.
… the kernel Two structural review points. Remove tests/scripts/cute_dsl_kernels/top_k/perf/. It is an experiment directory - three reviewers landed on the same folder - and the numbers it backs are already recorded in the commit messages and the PR description, so the scripts belong in a local tree rather than in-repo. This also reverts the compliance work done on those four files in a29d92e (translating comments, stripping non-ASCII, adding SPDX, narrowing the ruff escape); deleting them is the better answer to the same question. run_gvr_topk.py stays: it sits one level up and is a standalone driver plus a parameterised pytest sweep, not an experiment script, and its non-ASCII cleanup still stands. Move gvr_emission.py from attention_backend/sparse/ into cute_dsl_kernels/blackwell/top_k/. The constants in it are the kernel's buffer geometry contract - the segment bases at 0 / LIST_SEG_A / 2*LIST_SEG_A have to match the CUTLASS emitter exactly - and gvr_routing.py, which picks the tier those buffers are emitted for, already lives there and was already imported from across the package boundary. Folding it into modules/top_k.py instead would move 272 lines of kernel geometry into a torch module, which is the wrong direction. The move also removes the attention_backend -> cute_dsl_kernels reverse dependency; the import is now local to the kernel package. Import sites updated: modules/top_k.py and three in test_cute_dsl_gvr_topk_decode.py. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
|
@siyidNV Read your pushed head ( 1. Yes, the repair assumes dense counts, and the list is invalid for it on two separate grounds. The build keeps blocks with 2. 3. Keep the tests. 4. Suspend it. The repair only runs on the non-converged tail, and near a plateau the active list tends toward all blocks anyway, so a skip-aware repair buys nothing and reintroduces the bookkeeping that just failed. What I'd do instead of merging text — port #18094 semantically:
One thing I found while reading that the port has to include: the leader fb_fix fail-soft (the If it's easier I can push the port as a commit onto your branch; it's my defect family. |
|
@yuxianq You're right and my reply just above missed your point - I quoted the 45m22s as Working on it now. The JIT cost is driven by the number of distinct kernel configs the new So the 61 new cases are not the problem; the compile-time axes are:
I am measuring which of those configs actually compile separately and how much each costs, |
The cost of this file is dominated by CuTe DSL kernel compiles, not by
execution or by the case count: of the 60 cases this PR adds, 52 together run
in under a second, and the whole cost sits in eight cases that each instantiate
a fresh kernel config.
Two of those compiles were avoidable.
`ext_list` and `ext_closed_loop`'s list shape already agree on every launch
knob - top_k 512, N 131072, fp32, cluster_size 1, num_threads 512 - but
`emit_xstate` is part of the runner's compile key and only the latter passed an
`xstate` buffer, so the same kernel was built twice. `ext_list` now passes one
too; it costs nothing and additionally covers the xstate publish.
`ext_counts` swept top_k over {512, 1024, 2048}. The four modes it exercises
are count-based admission paths, which do not depend on K, while K is a
compile-time parameter - so the sweep bought two extra compiles and no extra
coverage. One K now.
Measured on a shared B200, new tests in isolation, cold:
before 471s (7m51s), 60 cases
after 333s (5m33s), 52 cases
`ext_closed_loop[list]` drops from 115.1s to 0.02s, which is the shared compile
landing. `degenerate_preidx`'s four configs are left alone: 16-bit has its own
path in this kernel and compress_ratio 1 and 4 are both production values, so
that 80s is coverage rather than waste.
For context on the file as a whole, same node, same protocol: 45m02s with these
tests, 38m21s with them deselected.
Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
|
@yuxianq Cut in be40f85. Numbers first, then a correction to my earlier ones. What the new tests cost, and after the cut
The cost is entirely kernel compiles, not execution and not the case count: of
I left Correction: I cannot reproduce the 30 minutes Same node, same cold protocol, whole file:
So the delta I measure is 6m41s, and the file was already 38 minutes without The likeliest explanation is the machine: mine is a shared B200 and every If you have the CI-side per-test timings for this file I would like to see them - On the timeout: I will leave TIMEOUT (120) for now rather than tune it to my |
|
@siyidNV 5m33s additional time is good for me, since the total time is ~45m, why do we need |
|
@yuxianq I dug into why my numbers are 2.5x yours, because if the 30 minutes is 1. The compile is single-threaded. One list-tier kernel compile: 117.07s wall 2. The kernel cache is per-process. Together those two give a machine-independent way to compare, which I think is
3. Instrumented count for the whole file, on my node:
So ~41 compiles at roughly 30-60s average is the entire cost; execution is Two things I cannot check from here:
If you can share the CI-side per-test timings or the compile count for this file Independently of all that, the cut in be40f85 stands on its own - it removes |
|
@longcheng-nv @yuxianq Found it, and the 38 minutes was my mistake in how I
and For what it is worth, 42m29s over 9 shards is ~4.7 min of shard time, which sits What I verified while chasing this, in case it is useful:
So the cost model is I will leave |
|
@longcheng-nv This is exactly what I needed - thank you, and yes, you read the Taking the port myself since it is my branch; I will ping you to review rather Your point 4 explains my 5 -> 83 result completely. I spliced Plan, following your recipe:
Gate: One question on 5: should the leader arm keep its recount at the undershoot side |
NVIDIA#18094 rewrote the same Phase-2/3 region this branch extends. Rather than merge the text - which mixes two state conventions, since NVIDIA#18094's phase 3 is written against main's unified _run_phases where the leader's own retry copy is gone - the kernel file is taken wholesale from this branch and NVIDIA#18094 is ported semantically, per longcheng-nv's recipe on the PR: - Port the three order-key helpers (f32_order_key_signed, order_key_signed_to_f32, order_key_mid_f32). - Replace Phase 3's overflow-only 10-iter shrink with the two-sided 48-iter bisection: anchor the untested bracket end at a float extreme and bisect on the signed order-key image, which collapses provably. - Clear the block-skip active list at the top of Phase 3's done != 1 block. Any dense re-count invalidates it on two grounds: the list is a superset only at or above the rung its build probe kept, and the repair anchors below that; and the compact stream-write replays the list walk against the smem_ptcnt of its matching compact pass, which a dense re-count overwrites. Clearing once at block entry keeps the done == 1 hot path on its compact write. - Stamp done = 2 on the leader's fail-soft arm instead of done = 1. It used to recount at the undershoot side and ship a -1-padded row as a "non-convergence encoding", which also hid the row from Phase 3 because done == 1 never enters the repair. The recount is dropped; the bisection measures anyway. - Keep _run_phases and the P1r rescue as they are: the rescue and NVIDIA#18094's synthetic bracket give the same exact answers on the rows both cover. Text-merging this file previously took the suite from 5 failures to 83, which is the two-state-convention problem above.
…au terminal The earlier port took NVIDIA#18094's two-sided bisection but stopped at the loop. The 48 lines after it are load-bearing and were missing: - On collapse with count < kK, fall back to s_thr[1] and re-count; val_lo admits >= kK by construction. - Re-check adjacency: if count > kCC and the bracket is already adjacent, take s_thr[2] and stamp done = 3 - the plateau terminal - then re-count. That second step is what relu_sparse_plateau was failing on. Without the done = 3 stamp the Phase-4 plateau fill never fires, so a ReLU-sparse row shipped its n_pos sure winners and padded the rest with -1 ("3 winners, 2045 pads"). It is also why the three earlier attempts missed: they all edited the loop, the leader's terminal code, or the active-list flag, and the missing piece sat directly below the loop. Whole file on B200, serial, cold: 268 passed, 1 xpassed, 0 failed (48m14s). The gate longcheng-nv named - hostile_hint, relu_sparse_plateau, mtp_hostile_hint - passes, as do plateau_terminal and degenerate_preidx. Method-level diff of base -> main vs base -> this branch shows NVIDIA#18094 touches exactly two methods, _run_phases and phase3_collect_candidates, plus three new free functions; both methods are also ones this branch changed. Only phase 3 is ported here - _run_phases stays as it is, per longcheng-nv: the P1r rescue covers what the synthetic bracket covers. Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>
|
@longcheng-nv Ported and green - thanks, your recipe was right and the one piece Whole file, B200, serial, cold: 268 passed, 1 xpassed, 0 failed (48m14s). What I had been missing. I ported the two-sided bisection and stopped at the # Undershoot at collapse: val_lo admits >= kK by construction.
if s_iscalars[0] < cutlass.Int32(kK):
... s_thr[0] = s_thr[1]; re-count
mid_chk, adj_chk = order_key_mid_f32(s_thr[1], s_thr[2])
if s_iscalars[0] > cutlass.Int32(kCC) and adj_chk:
... s_thr[0] = s_thr[2]; s_iscalars[1] = 3 # plateau terminalWithout that On your point 1, the active-list clear: with the port complete, phase 3 Point 5: the leader arm still stamps What finally cracked it was diffing base->main against base->this-branch at |
9 participants
Summary
Emission-assisted GVR top-K decode for the DeepSeek V4 sparse-attention indexer: the FP4 indexer GEMM epilogue now emits selection hints (per-block maxima / packed seed-count rows / a bucketed candidate list) that the GVR top-K kernel consumes through new opt-in fast paths, replacing most of its threshold-search and full-row scan work.
Shipped-path numbers (production routing, emission tax included): geomean 1.296x vs the #16457 baseline / 1.387x vs this PR's own stock kernel, worst cell 1.000 (no routed cell regresses on mean). Per-step over the full grid (B 1..1024 x raw ISL 8k..256k, all layers x all decode steps of real captures, 102k paired steps): mean 1.44-1.8x per model, zero mean-regressing cells, per-step regressions 0.14% of steps (all at raw ISL <= 64k; 128k+ has none). Unrouted kernel-capability corners reach 12.2x (table below, marked for reachability). All measured cells exact.
What's in the change
Indexer emission epilogue (
fp4_paged_mqa_logits.py, +1450)[rows, 8]— three threshold lines +count(>= line)accumulated with fp32 atomics in the epilogue (<<1% tax).{n0, void, n1, n2}control words. Measured emission tax: L1 +0.2-4.4%, L2 +9-14% of the indexer GEMM, flat in batch.GVR top-K consumption tiers (
gvr_topk_decode.py, +3603)wf: known-counts admission over the bucketed list — the tightest in-band line is a pure scalar lookup, the hit path degenerates to a filtered prefix copy straight into P4 rank selection; histogram-over-list and full fallback below.va: seed-count rows replace the preIdx gather and P2 threshold search on hit.vb: closed-loop three-line rungs (zero-emission variant, fallback tier).radix_lensfor thecute_dsl_indexer_topk_decodebranch atcompress_ratio > 1: the op takes 1-D request-level lens, but the old code passed the 2-Dkv_lens_cuda_2dslice; on the FP4-DSL path the live compressed lens aregen_indexer_kv_lens_cuda_runtime. Declared here per review so it is bisectable.[0, K)instead of computing the top-K. Found during real-model bring-up (42/231 dumped rows wrong = 21 layers x 2 sequences, first decode step each). Newphase1r_data_reseedrebuilds the refine bracket from the row itself (restores thecount(>= v_lo) >= Kinvariant), keeping the identity shortcut only where it is provably exact (all-tied row orN <= K). Non-degenerate rows pay nothing.Host routing + production wiring (new
gvr_routing.py, newgvr_emission.py,dsa.py,cute_dsl_custom_ops.py)plan_emission/pick_config: (B, N)-based tier selection (candidate-list tier only where it is net-positive: N >= 64k, B <= 4).GvrEmissionState: emission buffer lifecycle, device-side seed-row updates (CUDA-graph safe), prev-topK feedback loop.TRTLLM_GVR_EMISSION=1and composes with the existinguse_cute_dsl_topkrouting from [None][feat] top-k: route decode to CuTe DSL GVR top-k in e2e #16420 — default-path behavior is unchanged except for the two declared default-path fixes above.Tests
test_cute_dsl_fp4_paged_mqa_logits.py(+707).test_cute_dsl_gvr_topk_decode.py.xfaildocumenting a pre-existing corner inherited from the current kernel (reproduces on the unmodified upstream kernel): when the k-th tie class alone exceeds the candidate capacity, the selected value multiset is still exact but the index list can contain duplicate/unwritten slots. Requires >kC bit-identical scores at the boundary; never observed on real captures.Performance report
Protocol: real DeepSeek V4 captures (V4-Flash 21 indexer layers, V4-Pro 30 layers), all usable decode steps per layer, batch = row replication, nsys cold-L2 kernel-only timing on B200. Baseline = GVR kernel at the #16457 tip (identical to what main carries today). 486 cells, every cell exact (tie-aware score-multiset check).
wfkernel capability grid (UNROUTED: forced list tier;plan_emissionreaches only the B <= 4 columns at N >= 64k — the B >= 8 columns document kernel headroom, not shipped behavior) — speedup vs baseline:Geomeans over the full 9x9 grid (all layers x all steps):
Numbers above are kernel-only; the emission tax (L1 +0.2-4.4%, L2 +9-14% of the indexer GEMM) is charged on the indexer side and is why routing only enables the list tier at N >= 64k, B <= 4 — net accounting stays positive everywhere routed (headline geomeans above are tax-inclusive).
Routing caveats stated explicitly:
plan_emissionseeskv_cache_manager.max_seq_len // cr, an engine-lifetime constant (CUDA-graph capture bakes the tier in); per-step actual-length routing needs length-bucketed graphs and is deferred to a follow-up. Short rows in a long-max engine therefore run assist machinery the planner would refuse at their true length; the in-kernel validity/admission guards keep that exact.SKIP_MAX_BLOCKS = 8192(smem active-list budget) bounds the skip walk to N_local <= 262144; wins at 512k/1M are pure list/counts effects.Reproduction: the grid driver (per-step paired protocol, per-arm launch code, aggregation) is committed under
tests/scripts/cute_dsl_kernels/top_k/in this PR.Correctness validation
TRTLLM_GVR_EMISSION=1, the production-path selections of every indexer layer across all captured decode steps (231 rows) are score-multiset identical totorch.topk. This acceptance run is also what exposed (and now guards, via the new unit battery) the degenerate-preIdx bug fixed here.Review revision (2026-08-11)
All inline findings from the 2026-08-11 review round are addressed in the follow-up commits:
claimed_ccan no longer admit a starved list. Regression test:ext_list[starved](400 real candidates, pads lift the claim into the admission band).update_seed_rowsnow slope-fits log2(count) vs threshold from the previous step's (lines, counts) — the same construction as the harness'sderive_seed_lines_v4— and places lines at K-relative target counts; the list tier two-point-fits through the published exact k-th. No-fit rows fall back to multiplicative guards (~9% of kth, vs the degenerate 2e-4 pin). Chained test:ext_closed_loop(3 steps, k-th drift 0.05 between steps, list + counts tiers).[rows, 3]seed buffer is passed for the rungs tier (the packed-row column view is non-contiguous); count telemetry comes from the kernel's rung-count publish.heuristic_prev_topk(prev_topk from the last context token, xstate zeroed -> validity guard -> exact stock first step). Positional-identity assumption and why exactness survives churn is now stated in code; full slot-keying (no churn signal exists today) is follow-up scope.== 3 or == 8, ext-tiers-require-enable_r0and list-capacity constructor rejects, trace-time flag/tensor contract errors, block-skip single-bandvoidcontract,cand_ctlwidth comments, themutates_argslimitation re-verified on the pinned torch (IndexError fires when a declared-mutable Optional arg is None at call time - i.e. on every hint-less call), so the eager/CUDA-graph-only contract is now stated precisely in both op docstrings;TRTLLM_GVR_EMISSIONdocumented.TRTLLM_GVR_EXT->TRTLLM_GVR_EMISSION,gvr_ext.py->gvr_emission.py,GvrExtState->GvrEmissionState(review naming feedback).🤖 Generated with Claude Code