Skip to content
Open
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(

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

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

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

this gate can arm an emission tier while metadata.kv_lens_row_reorder is non-None. MTP computes row reordering from the configured draft width, but later draft iterations run with next_n == 1, satisfying this emission gate. The Top-K call then receives both emission tensors and order_row, which violates the explicit order_row is None assertion in the custom op. On a 148-SM B200 with three draft tokens, for example, this is reachable at batch sizes ≥74 and ≤256. Please disable emission when row ordering is active, or select only one of these two paths.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Confirmed and fixed in 0d527a7 — arming now also requires metadata.kv_lens_row_reorder is None, so the later MTP draft iterations that come through with next_n == 1 take the order_row dispatch without emission tensors. The op-side order_row is None assertion stays as the loud safety net rather than a silent drop.

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,

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

gvr_emit_kwargs can be non-empty on the FP8 path because self.use_gvr_emission does not require self.use_fp4. For list/counts routes this passes seed_thr, block_max_out, and candidate buffers to cute_dsl_fp8_paged_mqa_logits, whose registered schema does not accept these arguments, causing an unknown-keyword error before kernel launch. This is production-reachable because indexer_k_dtype defaults to FP8. Please either gate emission on self.use_fp4 or implement the corresponding FP8 emission interface.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Confirmed and fixed in 2bf2dae — the gate now also requires self.use_fp4, so gvr_emit_kwargs stays empty on the FP8 path (the FP8 op's registered schema indeed takes no emission arguments).

One nuance on reachability: emission sits behind TRTLLM_GVR_EMISSION=1 (default off), so the default production configuration never entered this path; with the flag on and the default FP8 indexer k dtype it failed exactly as you describe, before launch.

FP8 emission is not abandoned — the kernel-side port exists on a separate branch (validated on B200, FP8 400/0 plus FP4 cross-check 256/0) and will come as a follow-up PR that relaxes this gate together with the op schema.

)
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