Skip to content
Merged
Show file tree
Hide file tree
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 Jul 20, 2026
7eb1d88
[None][fix] GVR top-k: make block_max a trailing defaulted launch param
siyidNV Jul 20, 2026
dc42f23
[None][fix] GVR block-skip: support unaligned cluster slice starts
siyidNV Jul 20, 2026
36a542f
[None][perf] GVR block-skip: rung-tightening active-list build
siyidNV Jul 20, 2026
086eaba
[None][fix] GVR block-skip: guard list capacity/id width, bound tail …
siyidNV Jul 20, 2026
01bc507
[None][perf] GVR block-skip: host dispatch gate below N=200k
siyidNV Jul 20, 2026
e9f4865
[None][perf] GVR block-skip: skip-aware launch policy and dispatch gates
siyidNV Jul 21, 2026
344b082
[None][feat] GVR: emu references for epilogue seed counts and candidates
siyidNV Jul 20, 2026
289553f
[None][feat] GVR: waterfall L1 admission from external epilogue counts
siyidNV Jul 20, 2026
1fe3238
[None][perf] GVR ext counts: skip Phase 1 (preIdx gather) entirely
siyidNV Jul 20, 2026
44e9671
[None][perf] GVR ext counts v2: compose with the stock count/skip path
siyidNV Jul 20, 2026
604c379
[None][feat] GVR: xstate closed-loop writeback at Phase 4 exit
siyidNV Jul 20, 2026
f9428fc
[None][feat] GVR: waterfall L2 direct-to-P4 from pre-collected pairs
siyidNV Jul 20, 2026
076323c
[None][perf] GVR: slope-adaptive seed rung derivation (host helper)
siyidNV Jul 21, 2026
0e3470d
[None][perf] GVR ext counts: per-row dynamic routing to the stock path
siyidNV Jul 21, 2026
b30d476
[None][perf] GVR waterfall: cluster_size > 1 support for ext/L2/xstate
siyidNV Jul 21, 2026
6cf89c1
[None][perf] GVR L2 direct: two-parameter (K, K_max) in-list filtering
siyidNV Jul 22, 2026
b11f4a0
[None][perf] GVR L2: count-only smem-list selection (SoA candidates)
siyidNV Jul 23, 2026
5677541
[None][perf] GVR L2 v4: emitter-counted lines, known-counts admission
siyidNV Jul 23, 2026
cac0898
[None][perf] GVR L2 v5: bucketed segments, mapped-prefix copy, sample…
siyidNV Jul 24, 2026
3ca9731
[None][perf] GVR P4: in-place class compaction + three-tier boundary …
siyidNV Jul 25, 2026
da75cd8
[None][feat] GVR self_scan: fused self-contained closed-loop top-k mode
siyidNV Jul 25, 2026
7970473
[None][feat] GVR self_scan stage 2: block-max skip, tightest-line sin…
siyidNV Jul 25, 2026
5796021
[None][perf] GVR self_scan skip: two-pass compact-then-gather block scan
siyidNV Jul 25, 2026
6b117ca
[None][perf] GVR self_scan skip: split pass-2 into load-then-claim ph…
siyidNV Jul 25, 2026
aa6c275
[None][perf] GVR self_scan: software-pipeline the dense scan rounds
siyidNV Jul 25, 2026
26df9e6
[None][perf] GVR self_scan: cp.async deep-pipeline the dense scan
siyidNV Jul 25, 2026
2253600
[None][perf] GVR self_scan: pin the CTA to 1024 threads in the wrapper
siyidNV Jul 25, 2026
8ab4f74
[None][perf] GVR self_scan: pair-step the cp.async scan pipeline
siyidNV Jul 25, 2026
5970b3e
[None][feat] GVR ext_rungs: two-pass variant B (count external rungs …
siyidNV Jul 26, 2026
26608ef
[None][perf] GVR ext counts: pack lines and counts into one seed row
siyidNV Jul 26, 2026
09699c9
[None][perf] GVR: open the K>512 skip gate for parked ext counts; pub…
siyidNV Jul 26, 2026
993082c
[None][fix] GVR ext_rungs: cluster-side count publish + runtime rung …
siyidNV Jul 26, 2026
ed64e6c
[None][fix] GVR ext counts: full line validity on both admission prev…
siyidNV Jul 26, 2026
9b611a9
[None][perf] GVR ext counts: single-column parked count + merged admi…
siyidNV Jul 26, 2026
fc0137a
[None][perf] GVR wrapper: 512 threads for small-K list-hit rows
siyidNV Jul 26, 2026
a6728fa
[None][perf] GVR ext counts: fused count-free take for short rows
siyidNV Jul 26, 2026
c1c5ada
[None][perf] FP4 indexer: port fused block-metadata emission (P0)
siyidNV Jul 20, 2026
b5421cc
[None][perf] FP4 indexer: emit per-row seed counts in the GEMM epilog…
siyidNV Jul 20, 2026
57f1540
[None][perf] FP4 indexer: drop seed-count register budget cut
siyidNV Jul 20, 2026
b7f4de3
[None][perf] FP4 indexer: L2 candidate pre-collect in the GEMM epilogue
siyidNV Jul 20, 2026
68ef232
[None][fix] FP4 indexer: packed seed-row emission + counts-only thres…
siyidNV Jul 27, 2026
5d1812d
[None][chore] GVR: device-residency stamps in P4_TAIL_DBG
siyidNV Jul 27, 2026
aec6c66
[None][perf] GVR P4: eager position fetch + warp0 bin search
siyidNV Jul 27, 2026
71f3d70
[None][perf] GVR P4: walk-fused range + hist zero for list rows
siyidNV Jul 27, 2026
badea6c
[None][fix] GVR P4: revert warp0 bin search to the 3-step block search
siyidNV Jul 27, 2026
9bf8b0c
[None][perf] GVR: closed-loop adaptive block-skip gate
siyidNV Jul 27, 2026
71d569a
[None][fix] GVR: drop the block-skip pass-rate veto, keep col-6 count
siyidNV Jul 27, 2026
c10ae8c
[None][feat] GVR: host-side routing module (v1 constants)
siyidNV Jul 27, 2026
07cac99
[None][feat] FP4 indexer: bucketed candidate emission (v5 SoA contract)
siyidNV Jul 27, 2026
321f8eb
[None][feat] GVR: emission-assisted tiers on the production top-k op
siyidNV Jul 27, 2026
6d652af
[None][feat] DSA: emission-assisted GVR decode wiring (env opt-in)
siyidNV Jul 28, 2026
3a4f15d
[None][fix] cute-dsl ops: keep optional emission tensors out of mutat…
siyidNV Jul 28, 2026
143453c
[None][fix] GVR top-k: rescue degenerate preIdx seeds instead of iden…
siyidNV Jul 28, 2026
5390d7f
[None][fix] DSA GVR ext wiring: live seq domains, block_max width, du…
siyidNV Jul 28, 2026
af9df05
[None][fix] GVR: xstate anchor visibility barrier + defensive contrac…
siyidNV Jul 28, 2026
58ebe34
[None][fix] test: seed-row col 6 is the adaptive-skip pass count, not…
siyidNV Jul 29, 2026
8af7c59
[None][perf] Fused indexer->top-k handshake: PDL early release + per-…
siyidNV Jul 29, 2026
bc2621b
Revert "[None][perf] Fused indexer->top-k handshake: PDL early releas…
siyidNV Jul 30, 2026
b3410e2
[None][perf] GVR top-k: take the PDL wait after the prologue, not at …
siyidNV Jul 29, 2026
6949d7f
[None][perf] GVR P4: filter the fine recursion by value range, not by…
siyidNV Jul 29, 2026
d66fa12
[None][perf] GVR P4: lane-parallel warp bin sums in both bin searches
siyidNV Jul 29, 2026
bdeabeb
[None][perf] GVR P4: classify the scatter by value range too
siyidNV Jul 29, 2026
9ede841
Merge branch 'main' into perf/gvr-emission-topk
siyidNV Jul 30, 2026
bba4dd4
[None][fix] disable the two P4 range-test fast paths (not fp32-exact)
siyidNV Jul 30, 2026
66cc032
[None][perf] GVR routing: stay on the stock kernel below 8k context
siyidNV Jul 30, 2026
d2ab3c1
[None][fix] GVR: keep the stock kernel path on upstream behaviour
siyidNV Jul 30, 2026
c65dc31
[None][chore] GVR: condense the narrative comment blocks
siyidNV Jul 30, 2026
d697e3d
[None][fix] GVR: drop the const-flag branch shells from the stock path
siyidNV Jul 30, 2026
e3db787
[None][chore] GVR: put the P1r degenerate rescue behind a const_expr
siyidNV Jul 30, 2026
1fe0fee
[None][fix] GVR: keep range-test and debug scalars off the stock path
siyidNV Jul 30, 2026
84f8a36
[None][fix] GVR: let the assist path keep the compacted tail repair
siyidNV Jul 30, 2026
cf87bf0
[None][perf] GVR routing: pick the tier from the shape-dependent tax
siyidNV Jul 30, 2026
52973e6
[None][fix] GVR routing: cap the list tier at batch 1, measured
siyidNV Jul 30, 2026
a048dfa
[None][perf] GVR routing: gate the list tier on emitted volume
siyidNV Jul 30, 2026
dc10a1e
[None][perf] GVR routing: widen the list tier, hand over to rungs sooner
siyidNV Jul 30, 2026
75f89e6
[None][perf] GVR: turn the block-skip prefix on one doubling later
siyidNV Jul 30, 2026
cda9b03
[None][chore] GVR: make the coarse histogram width a knob
siyidNV Jul 30, 2026
738ce68
[None][perf] GVR routing: fix the block-skip and cluster-split gates
siyidNV Jul 31, 2026
78227d5
[None][fix] GVR: keep the P1 refine rescue off the stock path
siyidNV Jul 31, 2026
1e70d74
[None][perf] GVR routing: hand the mid-row band back to the stock kernel
siyidNV Jul 31, 2026
0982dc1
Revert "[None][fix] GVR: keep the P1 refine rescue off the stock path"
siyidNV Aug 1, 2026
15f1c45
[None][perf] GVR routing: refit the tiers on kernel-only emission cost
siyidNV Aug 3, 2026
445bb7d
[None][perf] GVR: park the tight lines for the list tier
siyidNV Aug 3, 2026
f8e6871
[None][fix] GVR: stop dividing the routing length by the compress rat…
siyidNV Aug 3, 2026
428c25c
[None][perf] GVR: resolve P4's coarse bin search lane-parallel
siyidNV Aug 3, 2026
ed8ea66
[None][fix] GVR: pass the stream where the FP4 indexer's __call__ exp…
siyidNV Aug 3, 2026
8c147a7
[None][perf] GVR: give K=512 the compacted boundary-class repair
siyidNV Aug 4, 2026
d1c4cd5
[None][perf] GVR: one candidate pass and a lane-parallel fold in the …
siyidNV Aug 4, 2026
97639c2
[None][perf] GVR routing: narrow the weak band after the phase-4 work
siyidNV Aug 4, 2026
2eb7d92
[None][perf] GVR: drain the last two serial folds out of phase 4
siyidNV Aug 4, 2026
26f4bc6
[None][perf] GVR: stop paying the big-class tail machinery on a 2-way…
siyidNV Aug 4, 2026
f378bf8
[None][perf] GVR: let the scatter hand the tie class to the tail repair
siyidNV Aug 5, 2026
7f823de
[None][perf] GVR: rank the tail class with the block, not one warp
siyidNV Aug 6, 2026
d17a58a
[None][perf] GVR routing: the weak band no longer holds at n_comp 65536
siyidNV Aug 8, 2026
9466f59
[None][perf] GVR: drop phase 4's fine level, at compile time
siyidNV Aug 8, 2026
5383607
[None][perf] GVR: collect the dead weight the fine level left behind
siyidNV Aug 8, 2026
8d3cd3c
Merge branch 'main' into perf/gvr-emission-topk
siyidNV Aug 9, 2026
1bd86c0
[None][test] give the GVR decode unittest entry 120 minutes on B300
siyidNV Aug 9, 2026
fda4f39
[None][fix] correct GVR ext list-tier consumption and gate its state
siyidNV Aug 9, 2026
fb9643e
[None][chore] GVR routing exclusivity, entry guards, comment slimming
siyidNV Aug 9, 2026
1ec6f2f
[None][test] split the heavy CuTe DSL files out of the B200 attention…
siyidNV Aug 9, 2026
0620720
[None][fix] address GVR emission review findings
siyidNV Aug 11, 2026
c48cacf
[None][chore] commit the GVR per-step perf grid driver (provenance)
siyidNV Aug 11, 2026
c4de505
[None][chore] size list-tier candidate buffers by the routing cap
siyidNV Aug 11, 2026
d1c5cd7
Merge branch 'main' of https://github.com/NVIDIA/TensorRT-LLM into fi…
siyidNV Aug 13, 2026
412a657
Merge branch 'main' of https://github.com/NVIDIA/TensorRT-LLM into fi…
siyidNV Aug 13, 2026
ba1d698
[None][fix] restore the plateau terminal dropped in the main merge
siyidNV Aug 14, 2026
2a0c11f
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV Aug 20, 2026
c58812d
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV Aug 20, 2026
6c581a7
[None][fix] fp4 paged MQA logits: make_fragment -> make_rmem_tensor
siyidNV Aug 20, 2026
f24be00
[None][chore] restore the .claude subtree clobbered in the main merge
siyidNV Aug 21, 2026
8f16b4a
[None][fix] emission-tier op calls must bypass the tiered fast path
siyidNV Aug 21, 2026
f9a74a6
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV Aug 23, 2026
8efc753
[None][feat] house emission-assisted GVR inside the TopK module
siyidNV Aug 21, 2026
4091764
[None][test] cover emission-state parking for reused decode slots
siyidNV Aug 24, 2026
a29d92e
[None][chore] address perf-script and test-list review comments
siyidNV Aug 25, 2026
d9fc773
[None][chore] drop the perf experiment dir, move gvr_emission next to…
siyidNV Aug 27, 2026
be40f85
[None][test] cut the new GVR top-K tests' JIT cost by 29%
siyidNV Aug 27, 2026
c19e85a
Merge remote-tracking branch 'origin/main' into perf/gvr-emission-topk
siyidNV Aug 28, 2026
2a923f6
[None][fix] complete the #18094 port: phase 3's post-loop plateau ter…
siyidNV Aug 28, 2026
875ce23
[None][chore] ruff-format: drop the blank lines left by the phase-3 t…
siyidNV Aug 28, 2026
b9513c8
Merge branch 'main' into port
siyidNV Aug 31, 2026
c15f3d9
Merge remote-tracking branch 'origin/main' into port
siyidNV Aug 31, 2026
2bf2dae
[None][fix] Gate GVR emission on the FP4 indexer scoring path
siyidNV Sep 1, 2026
0d527a7
[None][fix] Do not arm GVR emission when Top-K row reordering is active
siyidNV Sep 1, 2026
eb81cc7
[None][fix] Decline the GVR emission tier under torch.compile tracing
siyidNV Sep 1, 2026
b73d4cd
Merge remote-tracking branch 'origin/main' into HEAD
siyidNV Sep 1, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
38 changes: 38 additions & 0 deletions tensorrt_llm/_torch/attention_backend/sparse/dsa/indexer.py
Original file line number Diff line number Diff line change
Expand Up @@ -720,6 +720,16 @@ def __init__(
decode_implementation=decode_top_k_implementation,
compress_ratio=self.compress_ratio,
)
# GVR emission-assisted decode (opt-in, experimental): the FP4/FP8
# indexer epilogue emits candidates the GVR Top-K consumes (see
# gvr_emission / gvr_routing; state lives on the TopK module)
# only the FP4 scoring op accepts emission kwargs
self.use_gvr_emission = (
os.environ.get("TRTLLM_GVR_EMISSION", "0") == "1"
and decode_top_k_implementation == TopKImplementation.CUTE_DSL_GVR
and self.use_cute_dsl_paged_mqa_logits
and self.use_fp4
)

# Fused wk + weights_proj weight for single FP32 cuBLAS GEMM
# (populated in cache_derived_state; maps to TF32 tensor cores on Ampere+)
Expand Down Expand Up @@ -1596,6 +1606,13 @@ def sparse_attn_indexer(
gvr_prior_indices,
request_offset=num_generations,
)
if self.use_gvr_emission:
# 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(
Comment thread
longcheng-nv marked this conversation as resolved.
slice(num_generations, num_generations + num_contexts)
)

reuse_topk = (
self.mtp_index_share
Expand Down Expand Up @@ -1688,6 +1705,25 @@ def sparse_attn_indexer(
metadata.dsl_expand_factor > 1
and next_n == metadata.dsl_expand_factor * metadata.dsl_atom
)
gvr_emit_kwargs: dict = {}
# emitting for a step the Top-K cannot consume only churns
# the closed-loop state, so gate on the consumable shape
if (
self.use_gvr_emission
and gvr_prior_indices is not None
and next_n == 1
Comment thread
longcheng-nv marked this conversation as resolved.
and not dsl_atom_split
and num_gen_tokens <= 256
# ext tiers are single-CTA/sort-path only; row reordering
# routes the Top-K through order_row, which excludes them
and metadata.kv_lens_row_reorder is None
):
gvr_emit_kwargs = self.top_k.prepare_gvr_emission(
num_generations,
indexer_max_seq_len,
torch.cuda.get_device_properties(q_decode.device).multi_processor_count,
gvr_prior_indices,
)
if self.use_fp4:
# FP4 DSL signature splits DG's (q, sf_q) tuple into two
# separate args and requires q.dtype == uint8 (q_decode
Expand Down Expand Up @@ -1720,6 +1756,7 @@ def sparse_attn_indexer(
dsl_block_table,
dsl_schedule_meta,
indexer_max_seq_len,
**gvr_emit_kwargs,
)
else:
# FP8 DSL kernel natively supports next_n ∈ {1, 2, 3, 4}.
Expand Down Expand Up @@ -1747,6 +1784,7 @@ def sparse_attn_indexer(
fp8_block_table,
fp8_schedule_meta,
indexer_max_seq_len,
**gvr_emit_kwargs,
Comment thread
longcheng-nv marked this conversation as resolved.
)
else:
decode_q_scale = (
Expand Down
500 changes: 486 additions & 14 deletions tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py

Large diffs are not rendered by default.

Loading
Loading