Repository navigation
[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 116 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
272 changes: 272 additions & 0 deletions
272
tensorrt_llm/_torch/attention_backend/sparse/gvr_emission.py
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,272 @@ | ||
| # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & | ||
| # AFFILIATES. All rights reserved. | ||
| # SPDX-License-Identifier: Apache-2.0 | ||
| # | ||
| # Licensed under the Apache License, Version 2.0 (the "License"); | ||
| # you may not use this file except in compliance with the License. | ||
| # You may obtain a copy of the License at | ||
| # | ||
| # http://www.apache.org/licenses/LICENSE-2.0 | ||
| # | ||
| # Unless required by applicable law or agreed to in writing, software | ||
| # distributed under the License is distributed on an "AS IS" BASIS, | ||
| # WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. | ||
| # See the License for the specific language governing permissions and | ||
| # limitations under the License. | ||
| """Emission-assisted GVR top-k state for the DSA decode path. | ||
|
|
||
| Owns the persistent (graph-address-stable) buffers the emission tiers | ||
| ride on, the device-side closed-loop seed-row update (pure tensor ops, | ||
| CUDA-graph capturable) and the per-step routing decision. Opt-in: | ||
| without the flag the DSA decode path is unchanged. | ||
|
|
||
| Tier semantics (see gvr_routing): | ||
| * this step's TOP-K consumes what the PREVIOUS step's indexer | ||
| epilogue emitted; | ||
| * this step's INDEXER emits what the routing planned for the NEXT | ||
| step. | ||
| """ | ||
|
|
||
| import math | ||
| from typing import Optional | ||
|
|
||
| import torch | ||
|
|
||
| from ...cute_dsl_kernels.blackwell.top_k.gvr_routing import ( | ||
| LIST_EMIT_MAX_B, | ||
| LIST_EMIT_MIN_N, | ||
| TopkRoute, | ||
| pick_config, | ||
| plan_emission, | ||
| ) | ||
|
|
||
| # Bucketed candidate-list geometry: two tight segments of LIST_SEG_A | ||
| # entries plus a LIST_CAP_C-entry loose segment. | ||
| LIST_SEG_A = 8192 | ||
| LIST_CAP_C = 24576 | ||
| LIST_WIDTH = 2 * LIST_SEG_A + LIST_CAP_C | ||
|
|
||
| __all__ = ["GvrEmissionState", "LIST_EMIT_MIN_N", "LIST_PARK_LINE"] | ||
|
|
||
| # Closed-loop line placement: fit the slope of log2(count) vs threshold | ||
| # from the previous step's (lines, counts) and place the new lines at | ||
| # these K-relative target counts (t0 loosest .. t2 tightest). | ||
| LINE_TARGETS = (8.0, 5.0, 2.0) | ||
| LIST_T0_TARGET = 2.5 # list tier: single collect-line target (xK) | ||
| LIST_T0_COUNT_MAX = 6144.0 # keep n0 inside the [K+64, segA] admission band | ||
| SLOPE_MIN = 0.05 | ||
| SLOPE_MAX = 64.0 | ||
|
|
||
| # No-fit fallback: multiplicative guards around the published k-th value | ||
| # (t1 hugs it from below; t0/t2 guard by GUARD_LO/GUARD_HI spans). | ||
| FALLBACK_REL = 2.0**0.125 - 1.0 | ||
| FALLBACK_ABS = 1e-3 | ||
| GUARD_LO = 2.0 | ||
| GUARD_HI = 0.5 | ||
|
|
||
| # List tier only: park the two tight lines above any score so every | ||
| # admitted entry lands in the loosest segment. Any finite value above | ||
| # the score range works; the kernel's eligibility check only needs the | ||
| # three lines increasing and the loosest one finite. | ||
| LIST_PARK_LINE = 1.0e30 | ||
|
|
||
|
|
||
| class GvrEmissionState: | ||
| """Per-attention-backend emission state (persistent buffers).""" | ||
|
|
||
| def __init__( | ||
| self, | ||
| max_rows: int, | ||
| top_k: int, | ||
| device: torch.device, | ||
| enable_list_tier: bool = True, | ||
| own_prior: bool = True, | ||
| ): | ||
| self.max_rows = max_rows | ||
| self.top_k = top_k | ||
| # packed seed row: lines at cols 0..2, counts (emission-filled) | ||
| # at 3..5, adaptive-skip pass count at 6 | ||
| self.seed_row = torch.zeros((max_rows, 8), dtype=torch.float32, device=device) | ||
| # contiguous alias of the three lines for the rungs tier (a | ||
| # [rows, 3] column view of the packed row is non-contiguous) | ||
| self.seed_rungs = torch.zeros((max_rows, 3), dtype=torch.float32, device=device) | ||
| self.xstate = torch.zeros((max_rows, 8), dtype=torch.float32, device=device) | ||
| self.cand_vals: Optional[torch.Tensor] = None | ||
| self.cand_idx: Optional[torch.Tensor] = None | ||
| self.cand_ctl: Optional[torch.Tensor] = None | ||
| self.cand_cur: Optional[torch.Tensor] = None | ||
| if enable_list_tier: | ||
| # the routing only ever picks the list tier at | ||
| # batch <= LIST_EMIT_MAX_B, so the wide candidate buffers | ||
| # need that many rows, not max_rows (~0.33 MB/row/layer) | ||
| cand_rows = min(max_rows, LIST_EMIT_MAX_B) | ||
| self.cand_vals = torch.zeros( | ||
| (cand_rows, LIST_WIDTH), dtype=torch.float32, device=device | ||
| ) | ||
| self.cand_idx = torch.zeros((cand_rows, LIST_WIDTH), dtype=torch.int32, device=device) | ||
| self.cand_ctl = torch.zeros((cand_rows, 4), dtype=torch.int32, device=device) | ||
| self.cand_cur = torch.zeros((cand_rows, 4), dtype=torch.int32, device=device) | ||
| # previous-step top-k feedback (address-stable; zero-init -> | ||
| # first step's pre_idx points at index 0, a benign candidate). | ||
| # own_prior=False when the caller already keeps this state (the | ||
| # TopK module rides metadata's gvr_prior_indices). | ||
| self.prev_topk = ( | ||
| torch.zeros((max_rows, top_k), dtype=torch.int32, device=device) if own_prior else None | ||
| ) | ||
| # block_max prefix ([rows, nb_pad*4] fp32 warp-partials), | ||
| # allocated lazily once max_seq_len is known | ||
| self.block_max: Optional[torch.Tensor] = None | ||
|
|
||
| def ensure_block_max(self, max_seq_len: int) -> torch.Tensor: | ||
| nb4 = ((max_seq_len + 255) // 256 * 256) // 128 * 4 | ||
| # exact width: the runner asserts shape == (rows, nrec), so a | ||
| # wider reused buffer would trip it | ||
| if self.block_max is None or self.block_max.shape[1] != nb4: | ||
| # allocating during CUDA graph capture would bake a dangling | ||
| # address into the graph, so fail loudly instead | ||
| if torch.cuda.is_current_stream_capturing(): | ||
| raise RuntimeError( | ||
| "GvrEmissionState.ensure_block_max: (re)allocation requested " | ||
| "during CUDA graph capture; the block_max buffer must be " | ||
| "created by a warmup step before capture" | ||
| ) | ||
| self.block_max = torch.zeros( | ||
| (self.max_rows, nb4), dtype=torch.float32, device=self.seed_row.device | ||
| ) | ||
| return self.block_max | ||
|
|
||
| def plan( | ||
| self, batch: int, n_comp: int, num_sms: int, compress_ratio: int = 4 | ||
| ) -> tuple[str, TopkRoute]: | ||
| """Route this step: (tier the epilogue emits, launch knobs the | ||
| top-k consumes it with).""" | ||
| emit_tier = plan_emission( | ||
| batch, n_comp, self.top_k, have_epilogue=True, compress_ratio=compress_ratio | ||
| ) | ||
| if emit_tier == "list" and self.cand_vals is None: | ||
| # constructed with enable_list_tier=False: no candidate | ||
| # buffers to emit into, demote to the counts tier | ||
| emit_tier = "counts" | ||
| # emission and consumption happen inside the SAME forward (zero, | ||
| # emit, consume), so the consumer routes on this step's tier | ||
| route = pick_config(emit_tier, batch, n_comp, self.top_k, num_sms) | ||
| return emit_tier, route | ||
|
|
||
| def update_seed_rows(self, num_rows: int, emit_tier: str = "counts") -> None: | ||
| """Device-side closed-loop line update from the last publish. | ||
|
|
||
| Slope-fits log2(count) vs threshold from the previous step's | ||
| (lines, counts) and places the new lines at K-relative target | ||
| counts. Counts come from the packed row (counts/list emission) | ||
| or from the kernel's rung-count publish in xstate cols 4..6 | ||
| (rungs tier). Rows without a usable fit get multiplicative | ||
| guards around the published k-th value; rows with invalid | ||
| xstate (col 0 == 0, e.g. cold start) get non-finite lines, | ||
| which the kernel's validity guard routes to the stock path. | ||
| Pure tensor ops (graph-capturable). | ||
| """ | ||
| s = self.seed_row[:num_rows] | ||
| x = self.xstate[:num_rows] | ||
| valid = x[:, 0] > 0 | ||
| kth = x[:, 1] | ||
| anchor = x[:, 2] | ||
| t_prev0 = s[:, 0] | ||
| t_prev2 = s[:, 2] | ||
| cnts = x[:, 4:7] if emit_tier == "rungs" else s[:, 3:6] | ||
| k = float(self.top_k) | ||
| inf = torch.full_like(kth, float("inf")) | ||
| d_fb = kth.abs() * FALLBACK_REL + FALLBACK_ABS | ||
| if emit_tier == "list": | ||
| # two-point fit (t0_prev, n0) / (kth, K): kth is the exact | ||
| # k-th boundary on list rows | ||
| n0 = cnts[:, 0].clamp_min(1.0) | ||
| dthr = (kth - t_prev0).clamp_min(1e-3) | ||
| slope = ((torch.log2(n0) - math.log2(k)) / dthr).clamp(SLOPE_MIN, SLOPE_MAX) | ||
| tgt0 = min(LIST_T0_TARGET * k, LIST_T0_COUNT_MAX) | ||
| t0 = kth - math.log2(tgt0 / k) / slope | ||
| fit_ok = torch.isfinite(t_prev0) & (n0 > k) | ||
| t0 = torch.where(fit_ok, t0, kth - GUARD_LO * d_fb) | ||
| park = torch.full_like(kth, LIST_PARK_LINE) | ||
| new0 = torch.where(valid, t0, inf) | ||
| new1 = torch.where(valid, park, inf) | ||
| new2 = torch.where(valid, park + park, inf) | ||
| else: | ||
| c0 = cnts[:, 0].clamp_min(1.0) | ||
| c2 = cnts[:, 2].clamp_min(1.0) | ||
| dthr = (t_prev2 - t_prev0).clamp_min(1e-3) | ||
| slope = ((torch.log2(c0) - torch.log2(c2)) / dthr).clamp(SLOPE_MIN, SLOPE_MAX) | ||
| # anchor count estimate: slide the anchor onto the prev line fit | ||
| anch_c = (c2 * torch.exp2(-(anchor - t_prev2) * slope)).clamp(1.0, 1e6) | ||
| t0 = anchor + torch.log2(anch_c / (LINE_TARGETS[0] * k)) / slope | ||
| t1 = anchor + torch.log2(anch_c / (LINE_TARGETS[1] * k)) / slope | ||
| t2 = anchor + torch.log2(anch_c / (LINE_TARGETS[2] * k)) / slope | ||
| # t_prev2 < 1e29 also rejects a parked line left by a tier flip | ||
| fit_ok = torch.isfinite(t_prev0) & (t_prev2 < 1e29) & (c0 > c2) | ||
| t0 = torch.where(fit_ok, t0, kth - GUARD_LO * d_fb) | ||
| t1 = torch.where(fit_ok, t1, kth - 1e-6) | ||
| t2 = torch.where(fit_ok, t2, kth + GUARD_HI * d_fb) | ||
| # strictly ascending (kernel line-validity contract) | ||
| t1 = torch.maximum(t1, t0 + 1e-4) | ||
| t2 = torch.maximum(t2, t1 + 1e-4) | ||
| new0 = torch.where(valid, t0, inf) | ||
| new1 = torch.where(valid, t1, inf) | ||
| new2 = torch.where(valid, t2, inf) | ||
| s[:, 0] = new0 | ||
| s[:, 1] = new1 | ||
| s[:, 2] = new2 | ||
| s[:, 3:8] = 0.0 | ||
| rungs = self.seed_rungs[:num_rows] | ||
| rungs[:, 0] = new0 | ||
| rungs[:, 1] = new1 | ||
| rungs[:, 2] = new2 | ||
| if self.cand_ctl is not None: | ||
| nc = min(num_rows, self.cand_ctl.shape[0]) | ||
| self.cand_ctl[:nc].zero_() | ||
| self.cand_cur[:nc].zero_() | ||
|
|
||
| def indexer_emit_kwargs(self, emit_tier: str, num_rows: int) -> dict: | ||
| """kwargs for CuteDSLFP4PagedMQALogitsRunner.forward covering the | ||
| planned emission tier (caller merges into its call).""" | ||
| kw: dict = {} | ||
| if emit_tier in ("counts", "list"): | ||
| kw["seed_thr"] = self.seed_row[:num_rows] | ||
| if emit_tier == "list": | ||
| kw.update( | ||
| accept_cap=LIST_SEG_A, | ||
| cand_out=self.cand_vals[:num_rows], | ||
| cand_idx_out=self.cand_idx[:num_rows], | ||
| cand_ctl_out=self.cand_ctl[:num_rows], | ||
| cand_cur_out=self.cand_cur[:num_rows], | ||
| ) | ||
| return kw | ||
|
|
||
| def topk_ext_kwargs( | ||
| self, route: TopkRoute, num_rows: int, block_max: Optional[torch.Tensor] | ||
| ) -> dict: | ||
| """kwargs for trtllm::cute_dsl_gvr_topk_decode consuming this | ||
| step's emission per the picked route.""" | ||
| kw: dict = { | ||
| "xstate": self.xstate[:num_rows], | ||
| "cluster_size": route.cluster_size, | ||
| } | ||
| if route.num_threads is not None: | ||
| kw["num_threads"] = route.num_threads | ||
| if route.tier in ("counts", "list"): | ||
| kw["seed_thr"] = self.seed_row[:num_rows] | ||
| elif route.tier == "rungs": | ||
| # [rows, 3] seed selects the op's ext_rungs variant | ||
| kw["seed_thr"] = self.seed_rungs[:num_rows] | ||
| if route.tier == "list": | ||
| # accept_cap must match the emitter's segment geometry: the | ||
| # buffers are laid out at bases 0 / LIST_SEG_A / 2*LIST_SEG_A, | ||
| # and the consumer derives the C capacity from the tensor | ||
| # width minus 2*accept_cap. | ||
| kw.update( | ||
| cand_vals=self.cand_vals[:num_rows], | ||
| cand_idx=self.cand_idx[:num_rows], | ||
| cand_ctl=self.cand_ctl[:num_rows], | ||
| accept_cap=LIST_SEG_A, | ||
| ) | ||
| if route.attach_block_max and block_max is not None: | ||
| kw["block_max"] = block_max | ||
| return kw |
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.