-
Notifications
You must be signed in to change notification settings - Fork 2.7k
[None][perf] Emission-assisted GVR top-K decode for the DeepSeek V4 indexer #16953
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Merged
Merged
Changes from all commits
Commits
Show all changes
128 commits
Select commit
Hold shift + click to select a range
45e16a2
[None][perf] GVR top-k: block-skip for the R0 M-ary count pass and Ph…
siyidNV 7eb1d88
[None][fix] GVR top-k: make block_max a trailing defaulted launch param
siyidNV dc42f23
[None][fix] GVR block-skip: support unaligned cluster slice starts
siyidNV 36a542f
[None][perf] GVR block-skip: rung-tightening active-list build
siyidNV 086eaba
[None][fix] GVR block-skip: guard list capacity/id width, bound tail …
siyidNV 01bc507
[None][perf] GVR block-skip: host dispatch gate below N=200k
siyidNV e9f4865
[None][perf] GVR block-skip: skip-aware launch policy and dispatch gates
siyidNV 344b082
[None][feat] GVR: emu references for epilogue seed counts and candidates
siyidNV 289553f
[None][feat] GVR: waterfall L1 admission from external epilogue counts
siyidNV 1fe3238
[None][perf] GVR ext counts: skip Phase 1 (preIdx gather) entirely
siyidNV 44e9671
[None][perf] GVR ext counts v2: compose with the stock count/skip path
siyidNV 604c379
[None][feat] GVR: xstate closed-loop writeback at Phase 4 exit
siyidNV f9428fc
[None][feat] GVR: waterfall L2 direct-to-P4 from pre-collected pairs
siyidNV 076323c
[None][perf] GVR: slope-adaptive seed rung derivation (host helper)
siyidNV 0e3470d
[None][perf] GVR ext counts: per-row dynamic routing to the stock path
siyidNV b30d476
[None][perf] GVR waterfall: cluster_size > 1 support for ext/L2/xstate
siyidNV 6cf89c1
[None][perf] GVR L2 direct: two-parameter (K, K_max) in-list filtering
siyidNV b11f4a0
[None][perf] GVR L2: count-only smem-list selection (SoA candidates)
siyidNV 5677541
[None][perf] GVR L2 v4: emitter-counted lines, known-counts admission
siyidNV cac0898
[None][perf] GVR L2 v5: bucketed segments, mapped-prefix copy, sample…
siyidNV 3ca9731
[None][perf] GVR P4: in-place class compaction + three-tier boundary …
siyidNV da75cd8
[None][feat] GVR self_scan: fused self-contained closed-loop top-k mode
siyidNV 7970473
[None][feat] GVR self_scan stage 2: block-max skip, tightest-line sin…
siyidNV 5796021
[None][perf] GVR self_scan skip: two-pass compact-then-gather block scan
siyidNV 6b117ca
[None][perf] GVR self_scan skip: split pass-2 into load-then-claim ph…
siyidNV aa6c275
[None][perf] GVR self_scan: software-pipeline the dense scan rounds
siyidNV 26df9e6
[None][perf] GVR self_scan: cp.async deep-pipeline the dense scan
siyidNV 2253600
[None][perf] GVR self_scan: pin the CTA to 1024 threads in the wrapper
siyidNV 8ab4f74
[None][perf] GVR self_scan: pair-step the cp.async scan pipeline
siyidNV 5970b3e
[None][feat] GVR ext_rungs: two-pass variant B (count external rungs …
siyidNV 26608ef
[None][perf] GVR ext counts: pack lines and counts into one seed row
siyidNV 09699c9
[None][perf] GVR: open the K>512 skip gate for parked ext counts; pub…
siyidNV 993082c
[None][fix] GVR ext_rungs: cluster-side count publish + runtime rung …
siyidNV ed64e6c
[None][fix] GVR ext counts: full line validity on both admission prev…
siyidNV 9b611a9
[None][perf] GVR ext counts: single-column parked count + merged admi…
siyidNV fc0137a
[None][perf] GVR wrapper: 512 threads for small-K list-hit rows
siyidNV a6728fa
[None][perf] GVR ext counts: fused count-free take for short rows
siyidNV c1c5ada
[None][perf] FP4 indexer: port fused block-metadata emission (P0)
siyidNV b5421cc
[None][perf] FP4 indexer: emit per-row seed counts in the GEMM epilog…
siyidNV 57f1540
[None][perf] FP4 indexer: drop seed-count register budget cut
siyidNV b7f4de3
[None][perf] FP4 indexer: L2 candidate pre-collect in the GEMM epilogue
siyidNV 68ef232
[None][fix] FP4 indexer: packed seed-row emission + counts-only thres…
siyidNV 5d1812d
[None][chore] GVR: device-residency stamps in P4_TAIL_DBG
siyidNV aec6c66
[None][perf] GVR P4: eager position fetch + warp0 bin search
siyidNV 71f3d70
[None][perf] GVR P4: walk-fused range + hist zero for list rows
siyidNV badea6c
[None][fix] GVR P4: revert warp0 bin search to the 3-step block search
siyidNV 9bf8b0c
[None][perf] GVR: closed-loop adaptive block-skip gate
siyidNV 71d569a
[None][fix] GVR: drop the block-skip pass-rate veto, keep col-6 count
siyidNV c10ae8c
[None][feat] GVR: host-side routing module (v1 constants)
siyidNV 07cac99
[None][feat] FP4 indexer: bucketed candidate emission (v5 SoA contract)
siyidNV 321f8eb
[None][feat] GVR: emission-assisted tiers on the production top-k op
siyidNV 6d652af
[None][feat] DSA: emission-assisted GVR decode wiring (env opt-in)
siyidNV 3a4f15d
[None][fix] cute-dsl ops: keep optional emission tensors out of mutat…
siyidNV 143453c
[None][fix] GVR top-k: rescue degenerate preIdx seeds instead of iden…
siyidNV 5390d7f
[None][fix] DSA GVR ext wiring: live seq domains, block_max width, du…
siyidNV af9df05
[None][fix] GVR: xstate anchor visibility barrier + defensive contrac…
siyidNV 58ebe34
[None][fix] test: seed-row col 6 is the adaptive-skip pass count, not…
siyidNV 8af7c59
[None][perf] Fused indexer->top-k handshake: PDL early release + per-…
siyidNV bc2621b
Revert "[None][perf] Fused indexer->top-k handshake: PDL early releas…
siyidNV b3410e2
[None][perf] GVR top-k: take the PDL wait after the prologue, not at …
siyidNV 6949d7f
[None][perf] GVR P4: filter the fine recursion by value range, not by…
siyidNV d66fa12
[None][perf] GVR P4: lane-parallel warp bin sums in both bin searches
siyidNV bdeabeb
[None][perf] GVR P4: classify the scatter by value range too
siyidNV 9ede841
Merge branch 'main' into perf/gvr-emission-topk
siyidNV bba4dd4
[None][fix] disable the two P4 range-test fast paths (not fp32-exact)
siyidNV 66cc032
[None][perf] GVR routing: stay on the stock kernel below 8k context
siyidNV d2ab3c1
[None][fix] GVR: keep the stock kernel path on upstream behaviour
siyidNV c65dc31
[None][chore] GVR: condense the narrative comment blocks
siyidNV d697e3d
[None][fix] GVR: drop the const-flag branch shells from the stock path
siyidNV e3db787
[None][chore] GVR: put the P1r degenerate rescue behind a const_expr
siyidNV 1fe0fee
[None][fix] GVR: keep range-test and debug scalars off the stock path
siyidNV 84f8a36
[None][fix] GVR: let the assist path keep the compacted tail repair
siyidNV cf87bf0
[None][perf] GVR routing: pick the tier from the shape-dependent tax
siyidNV 52973e6
[None][fix] GVR routing: cap the list tier at batch 1, measured
siyidNV a048dfa
[None][perf] GVR routing: gate the list tier on emitted volume
siyidNV dc10a1e
[None][perf] GVR routing: widen the list tier, hand over to rungs sooner
siyidNV 75f89e6
[None][perf] GVR: turn the block-skip prefix on one doubling later
siyidNV cda9b03
[None][chore] GVR: make the coarse histogram width a knob
siyidNV 738ce68
[None][perf] GVR routing: fix the block-skip and cluster-split gates
siyidNV 78227d5
[None][fix] GVR: keep the P1 refine rescue off the stock path
siyidNV 1e70d74
[None][perf] GVR routing: hand the mid-row band back to the stock kernel
siyidNV 0982dc1
Revert "[None][fix] GVR: keep the P1 refine rescue off the stock path"
siyidNV 15f1c45
[None][perf] GVR routing: refit the tiers on kernel-only emission cost
siyidNV 445bb7d
[None][perf] GVR: park the tight lines for the list tier
siyidNV f8e6871
[None][fix] GVR: stop dividing the routing length by the compress rat…
siyidNV 428c25c
[None][perf] GVR: resolve P4's coarse bin search lane-parallel
siyidNV ed8ea66
[None][fix] GVR: pass the stream where the FP4 indexer's __call__ exp…
siyidNV 8c147a7
[None][perf] GVR: give K=512 the compacted boundary-class repair
siyidNV d1c4cd5
[None][perf] GVR: one candidate pass and a lane-parallel fold in the …
siyidNV 97639c2
[None][perf] GVR routing: narrow the weak band after the phase-4 work
siyidNV 2eb7d92
[None][perf] GVR: drain the last two serial folds out of phase 4
siyidNV 26f4bc6
[None][perf] GVR: stop paying the big-class tail machinery on a 2-way…
siyidNV f378bf8
[None][perf] GVR: let the scatter hand the tie class to the tail repair
siyidNV 7f823de
[None][perf] GVR: rank the tail class with the block, not one warp
siyidNV d17a58a
[None][perf] GVR routing: the weak band no longer holds at n_comp 65536
siyidNV 9466f59
[None][perf] GVR: drop phase 4's fine level, at compile time
siyidNV 5383607
[None][perf] GVR: collect the dead weight the fine level left behind
siyidNV 8d3cd3c
Merge branch 'main' into perf/gvr-emission-topk
siyidNV 1bd86c0
[None][test] give the GVR decode unittest entry 120 minutes on B300
siyidNV fda4f39
[None][fix] correct GVR ext list-tier consumption and gate its state
siyidNV fb9643e
[None][chore] GVR routing exclusivity, entry guards, comment slimming
siyidNV 1ec6f2f
[None][test] split the heavy CuTe DSL files out of the B200 attention…
siyidNV 0620720
[None][fix] address GVR emission review findings
siyidNV c48cacf
[None][chore] commit the GVR per-step perf grid driver (provenance)
siyidNV c4de505
[None][chore] size list-tier candidate buffers by the routing cap
siyidNV d1c5cd7
Merge branch 'main' of https://github.com/NVIDIA/TensorRT-LLM into fi…
siyidNV 412a657
Merge branch 'main' of https://github.com/NVIDIA/TensorRT-LLM into fi…
siyidNV ba1d698
[None][fix] restore the plateau terminal dropped in the main merge
siyidNV 2a0c11f
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV c58812d
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV 6c581a7
[None][fix] fp4 paged MQA logits: make_fragment -> make_rmem_tensor
siyidNV f24be00
[None][chore] restore the .claude subtree clobbered in the main merge
siyidNV 8f16b4a
[None][fix] emission-tier op calls must bypass the tiered fast path
siyidNV f9a74a6
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV 8efc753
[None][feat] house emission-assisted GVR inside the TopK module
siyidNV 4091764
[None][test] cover emission-state parking for reused decode slots
siyidNV a29d92e
[None][chore] address perf-script and test-list review comments
siyidNV d9fc773
[None][chore] drop the perf experiment dir, move gvr_emission next to…
siyidNV be40f85
[None][test] cut the new GVR top-K tests' JIT cost by 29%
siyidNV c19e85a
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV 2a923f6
[None][fix] complete the #18094 port: phase 3's post-loop plateau ter…
siyidNV 875ce23
[None][chore] ruff-format: drop the blank lines left by the phase-3 t…
siyidNV b9513c8
Merge branch 'main' into port
siyidNV c15f3d9
Merge remote-tracking branch 'origin/main' into port
siyidNV 2bf2dae
[None][fix] Gate GVR emission on the FP4 indexer scoring path
siyidNV 0d527a7
[None][fix] Do not arm GVR emission when Top-K row reordering is active
siyidNV eb81cc7
[None][fix] Decline the GVR emission tier under torch.compile tracing
siyidNV b73d4cd
Merge remote-tracking branch 'origin/main' into HEAD
siyidNV File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Large diffs are not rendered by default.
Oops, something went wrong.
Oops, something went wrong.
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.