Skip to content
Merged
Show file tree
Hide file tree
Changes from 107 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
149 changes: 142 additions & 7 deletions tensorrt_llm/_torch/attention_backend/sparse/dsa/indexer.py
Original file line number Diff line number Diff line change
Expand Up @@ -686,6 +686,16 @@ def __init__(
self.use_cute_dsl_paged_mqa_logits = (
sparse_params.use_cute_dsl_paged_mqa_logits and IS_CUTLASS_DSL_AVAILABLE
)
# GVR emission-assisted decode (opt-in, experimental): the FP4 indexer
# epilogue emits candidates for the top-k (see gvr_emission / gvr_routing).
self.use_gvr_emission = (
os.environ.get("TRTLLM_GVR_EMISSION", "0") == "1"
and self.use_cute_dsl_topk
and self.use_cute_dsl_paged_mqa_logits
and self.use_fp4
)
self._gvr_emission = None # lazy GvrEmissionState (first decode step)
self._gvr_route = None
self.weight_scale_factor = self.softmax_scale * self.n_heads**-0.5

self._enable_heuristic_topk = (
Expand Down Expand Up @@ -1365,6 +1375,24 @@ def _call_paged_mqa_logits(
max_seq_len,
)

def _ensure_gvr_emission(self, metadata: DSAtrtllmAttentionMetadata, device: torch.device):
"""Lazy per-layer GVR ext state (see gvr_emission.GvrEmissionState)."""
if self._gvr_emission is None:
from ..gvr_emission import LIST_EMIT_MIN_N, GvrEmissionState

# get_indexer_max_seq_len already returns the compressed length
n_comp = metadata.get_indexer_max_seq_len()
self._gvr_emission = GvrEmissionState(
max_rows=metadata.max_num_sequences,
top_k=self.index_topk,
device=device,
# the list tier is only reachable at n_comp >= LIST_EMIT_MIN_N;
# skip its large candidate buffers when the engine's static
# length cannot get there
enable_list_tier=n_comp >= LIST_EMIT_MIN_N,
)
return self._gvr_emission

def sparse_attn_indexer(
self,
metadata: DSAtrtllmAttentionMetadata,
Expand Down Expand Up @@ -1612,14 +1640,29 @@ def sparse_attn_indexer(
# bottom of the decode block): new gens from finishing prefill
# append after currently-active gens, i.e., slots
# [num_generations : num_generations + num_contexts].
if self._enable_heuristic_topk and has_prefill and not metadata.skip_indexer_for_ctx_reqs:
local_layer = metadata.kv_cache_manager.layer_offsets[self.layer_idx]
if (
has_prefill
and not metadata.skip_indexer_for_ctx_reqs
and (self._enable_heuristic_topk or self.use_gvr_emission)
):
ctx_seq_lens = metadata.seq_lens[:num_contexts]
# Per-sequence last context-token offset (exclusive cumsum minus 1).
last_ctx_idx = (torch.cumsum(ctx_seq_lens, dim=0) - 1).to(dtype=torch.long)
metadata.heuristic_prev_topk[
local_layer, num_generations : num_generations + num_contexts
].copy_(topk_indices_buffer[last_ctx_idx, :])
last_ctx_topk = topk_indices_buffer[last_ctx_idx, :]
if self._enable_heuristic_topk:
local_layer = metadata.kv_cache_manager.layer_offsets[self.layer_idx]
metadata.heuristic_prev_topk[
local_layer, num_generations : num_generations + num_contexts
].copy_(last_ctx_topk)
if self.use_gvr_emission:
# Ext state is positional like heuristic_prev_topk (no churn
# signal exists to key on): zeroed xstate cold-starts new gens
# onto the stock path; stale rows after churn only mis-place
# lines - counts are re-measured in-kernel, so exactness holds.
st = self._ensure_gvr_emission(metadata, topk_indices_buffer.device)
new_gens = slice(num_generations, num_generations + num_contexts)
st.prev_topk[new_gens].copy_(last_ctx_topk[:, : st.prev_topk.shape[1]])
st.xstate[new_gens].zero_()

reuse_topk = (
self.mtp_index_share
Expand All @@ -1633,6 +1676,7 @@ def sparse_attn_indexer(
metadata.shared_topk_indices[:num_generations, :]
)
elif has_decode and not metadata.skip_indexer_for_gen_reqs:
gvr_step_ok = False
# Get decode lengths per request (from seq_lens) for validation
gen_seq_lens = metadata.seq_lens[num_contexts : num_contexts + num_generations]
max_decode_len = gen_seq_lens.max().item()
Expand Down Expand Up @@ -1735,6 +1779,43 @@ def sparse_attn_indexer(
dsl_block_table = metadata.block_table_expanded[:exp_B]
dsl_schedule_meta = metadata.scheduler_metadata_buffer_expanded

gvr_emit_kwargs = {}
# one step gate shared with the consume branch below:
# emitting for a step the top-k cannot consume only
# churns the closed-loop state
gvr_step_ok = (
self.use_gvr_emission
and use_custom_topk
and next_n == 1
and not dsl_atom_split
and num_gen_tokens <= 256
)
if gvr_step_ok:
st = self._ensure_gvr_emission(metadata, q_fp8.device)
# indexer_max_seq_len is already the compressed length
# (get_indexer_max_seq_len divides); do not divide again.
# routing N is the ENGINE-STATIC max (graph capture
# bakes the tier in); short live rows in a long-max
# engine run assist machinery the planner would refuse
# at their true length - exactness holds via the
# in-kernel guards, the tax is routed pessimistically
emit_tier, self._gvr_route = st.plan(
Comment thread
longcheng-nv marked this conversation as resolved.
Outdated
batch_size,
indexer_max_seq_len,
torch.cuda.get_device_properties(q_fp8.device).multi_processor_count,
compress_ratio=max(self.compress_ratio, 1),
)
if emit_tier in ("counts", "list", "rungs"):
st.update_seed_rows(batch_size, emit_tier)
if emit_tier in ("counts", "list"):
gvr_emit_kwargs = st.indexer_emit_kwargs(emit_tier, batch_size)
if self._gvr_route.attach_block_max or emit_tier in (
"counts",
"list",
):
gvr_emit_kwargs["block_max_out"] = st.ensure_block_max(
indexer_max_seq_len
)[:batch_size]
logits_decode = torch.ops.trtllm.cute_dsl_fp4_paged_mqa_logits(
dsl_q,
decode_q_scale,
Expand All @@ -1744,6 +1825,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 @@ -1813,7 +1895,46 @@ def sparse_attn_indexer(
if not metadata.use_cute_dsl_topk:
heuristic_scratch = metadata.heuristic_scratch_values[:num_gen_tokens]

if self.use_cute_dsl_topk and self._enable_heuristic_topk:
# tier "none" (see gvr_routing) falls through to the stock
# branch; the untouched ext state reads as a cold start there.
if (
self.use_gvr_emission
and gvr_step_ok
and self._gvr_emission is not None
and self._gvr_route is not None
and self._gvr_route.tier != "none"
):
# emission-assisted GVR: consume what the indexer epilogue
# emitted this step (seed row / list / block_max per route)
st = self._gvr_emission
out_slice = topk_indices_buffer[token_offset : token_offset + num_gen_tokens, :]
# GVR op takes RAW-domain seq_lens (kernel ceil-divides by
# compress_ratio); on the DSL path the live compressed lens
# are gen_indexer_kv_lens_cuda_runtime, kv_lens_cuda_2d is 0.
if self.compress_ratio > 1:
gvr_lens = metadata.gen_indexer_kv_lens_cuda_runtime
assert gvr_lens is not None
seq_1d = gvr_lens.reshape(-1)[:num_gen_tokens] * self.compress_ratio
else:
seq_1d = gen_kv_lens_cuda.reshape(-1)[:num_gen_tokens]
ext_kw = st.topk_ext_kwargs(
self._gvr_route,
num_gen_tokens,
st.block_max[:num_gen_tokens] if st.block_max is not None else None,
)
torch.ops.trtllm.cute_dsl_gvr_topk_decode(
logits_decode,
st.prev_topk[:num_gen_tokens],
seq_1d.contiguous(),
out_slice,
self.index_topk,
next_n,
self.compress_ratio,
max_seq_len=indexer_max_seq_len,
**ext_kw,
)
st.prev_topk[:num_gen_tokens].copy_(out_slice)
Comment thread
brnguyen2 marked this conversation as resolved.
Outdated
Comment thread
longcheng-nv marked this conversation as resolved.
Outdated
elif self.use_cute_dsl_topk and self._enable_heuristic_topk:
torch.ops.trtllm.cute_dsl_gvr_topk_decode(
logits_decode,
pre_idx,
Expand All @@ -1826,9 +1947,23 @@ def sparse_attn_indexer(
order_row=metadata.kv_lens_row_reorder,
)
elif self.use_cute_dsl_topk and (self.compress_ratio == 1 or next_n == 1):
# latent-bug fix: the op takes 1-D request-level lens, but
# context_lens is the 2-D kv_lens_cuda_2d slice; on the
# FP4-DSL path the live compressed lens are
# gen_indexer_kv_lens_cuda_runtime
if self.compress_ratio > 1:
radix_lens = (
metadata.gen_indexer_kv_lens_cuda_runtime
Comment thread
brnguyen2 marked this conversation as resolved.
Outdated
if self.use_cute_dsl_paged_mqa_logits
and metadata.gen_indexer_kv_lens_cuda_runtime is not None
else context_lens
)
radix_lens = radix_lens.reshape(-1)
else:
radix_lens = gen_kv_lens_cuda
torch.ops.trtllm.cute_dsl_indexer_topk_decode(
logits_decode,
context_lens if self.compress_ratio > 1 else gen_kv_lens_cuda,
radix_lens,
topk_indices_buffer[token_offset : token_offset + num_gen_tokens, :],
self.index_topk,
next_n,
Expand Down
Loading
Loading