[None][perf] DSv4 GVR top-k (CuTe DSL): P4 histogram refinement + redundant-warp sync reduction by siyidNV · Pull Request #16424 · NVIDIA/TensorRT-LLM

siyidNV · 2026-07-15T10:57:38Z

Summary by CodeRabbit

Performance
- Optimized top‑K decoding on Blackwell GPUs with reduced synchronization and more efficient shared-memory access.
- Improved convergence and writeback efficiency for large-scale decoding workloads.
Reliability
- Fixed potential race conditions during clustered aggregation, improving result consistency across repeated operations.
Testing
- Added coverage for alternate execution configurations in BF16 and FP32 top‑K decoding paths.

Description

Kernel-level optimizations to the CuTe DSL GVR top-k decode kernel
(cute_dsl_gvr_topk_decode) for the DSv4 indexer path on Blackwell. No
algorithm changes: the secant search, candidate collection, and histogram
snap semantics (and their exact tie-aware outputs) are unchanged; every
change attacks synchronization overhead, leader-thread serialization, and
SMEM scalar round-trips inside the existing phases. Scope is
gvr_topk_decode.py only (per GVR owner guidance the load-balance
variant gvr_topk_decode_load_balance.py is a separate, pre-production
line — 0 lines changed there; it reuses run_one_row and all new
parameters are keyword-defaults, so it compiles unchanged and knob-off
restores the exact leader behavior).

Related PR: #14413 (DSv4 indexer top-k dispatcher tuning — this PR is the
CuTe DSL GVR kernel-side counterpart; measurement methodology and
synthetic-cell data are shared). Builds on top of #15835.

Why one PR and not two: the redundant-warp commit is measured AGAINST the
refinement chain as its baseline (the A/B table isolates it), the two
lines touch the same regions of Phase-4 code, and landing them separately
would force an interim retune of the barrier layout that the second
commit immediately deletes. Both are pure kernel-level changes with
identical (tie-aware exact) outputs; if preferred, the last commit
cherry-picks cleanly onto the first six as a follow-up PR.

Changes

P4 histogram refinement chain: the snap loop steps one distinct
value per iteration (~0.45us: full candidate re-scan + 2 barriers)
while real logits concentrate count mass at the k-th boundary, so the
selected 1/1024 bin often holds tens of values — multi-level histogram
refinement (rebuild over the selected bin, gated on bin occupancy)
leaves the snap loop 0-2 steps. Plus: SMEM window-base hoisting (kills
per-access cluster-window S2R recompute), single-pass converged
writeback (zero SMEM atomics), bracket-window level-1 histogram, an
fp32 range-overflow clamp (fix), and a parity double-buffer for the
cluster DSMEM count exchange (fix for a pre-existing PTX-model race).
Redundant-warp sync reduction: after per-warp partials land in
SMEM (one visibility barrier), EVERY warp redundantly performs the
block reduce + decision a single leader used to do. Identical inputs
reduced in an identical order are bit-identical across warps, so the
publish barrier, the leader serialization, and the SMEM scalar
re-reads all disappear:
- Phase 2 (p2_warp_redundant, cluster_size==1): one barrier per
  secant round (was 3); threshold/bracket/counts live in registers;
  staging parity-banked; canonical exit state published once for P3.
- Phase 4 (p4_warp_redundant): k-th bin search 3 barriers -> 1
  with the serial 16-slot prefix and 64-deep serial bin walk made
  lane-parallel (idx-shuffle inclusive scans + unique-crossing
  ballot); snap loop 2 barriers/iter -> 1 with parity double-buffered
  staging; threshold and convergence state in registers end to end.

New compile-time knobs (constructor): p4_warp_redundant: bool = True,
p2_warp_redundant: bool = True. Both off restores the leader cadence
byte-for-byte (used as the A/B baseline below; can be dropped in a
follow-up once soak is done).

Performance

Setup: B300 (148 SM, one GPU of an 8-GPU node, others idle — parallel
sharded measurement was found to cross-contaminate microsecond timings
via power/clock steering), V4-Pro shape (K=1024, compress_ratio=4, fp32
logits), realistic synthetic cells (per-layer empirical inverse-CDF + GPD
tail calibrated from production captures; temporally coherent preIdx),
cluster_size=1, 200-iteration CUDA-event wall means.

A/B = both knobs off (leader paths, refinement chain included) vs both
knobs on — the table isolates the redundant-warp commit against the exact
base it merges onto (main @ f84b6da, includes #15835).

N \ BS	1	2	4	8	16	32	64	128	256	512	1024
4K	1.019	1.001	1.046	1.016	1.023	1.013	1.018	1.022	1.005	1.064	1.117
8K	1.029	1.022	1.007	1.010	1.027	1.016	1.035	0.997	1.034	1.088	1.100
16K	1.025	1.029	1.032	1.035	1.039	1.049	1.064	1.108	1.033	1.052	1.073
32K	1.020	1.118	1.119	1.119	1.115	1.117	1.122	1.094	1.041	1.055	1.044
64K	1.068	1.063	1.058	1.059	1.046	1.048	1.049	1.049	1.032	1.048	1.041
128K	1.056	1.043	1.035	1.037	1.036	1.037	1.036	1.037	1.052	1.031	1.029
256K	1.042	1.035	1.034	1.026	1.027	1.025	1.036	1.031	1.029	1.017	1.014

Speedup (off/on wall): geomean 1.045x, max 1.122x (32K/BS64; the
32K/BS2-64 band runs 1.10-1.12x, with 4K/BS1024, 8K/BS1024 and 16K/BS128
also >= 1.10x), min 0.997x (8K/BS128, within +/-0.3% run jitter).
4K/BS2 is shown as its 10-rep order-swapped median (1.001): with only 2
CTAs resident on a ~14us kernel, single measurements swing +/-12% in BOTH
directions from DVFS jitter. 0 correctness failures across all 77 cells
x 2 configs (tie-aware set equality vs torch.topk).

ncu (256K, BS=64): barrier-stall 14.1% -> 6.1% of warp-issue stalls;
registers 64 -> 64; zero local-memory traffic on both sides.

Notes for reviewers

The redundant-warp transformation is safe iff every warp's replay is
deterministic: the reduce reads a fixed slot range in a fixed order and
uses shuffle-based reductions only — no atomics, no arrival-order
dependence. The two parity double-buffers (P2 smem_wcnt 2x banks, P4
staging banks in smem_hist) exist because removing the publish
barrier legitimately lets a warp run one iteration ahead; the per-round
staging barrier bounds the drift to exactly one iteration.
_kth_bin_search_rw returns the selected (threshold, bin count) in
registers and deliberately has NO trailing barrier; both call sites
barrier before the next _hist_build or before the snap staging —
grep for "step-3 readers".
The P2 redundant cadence is gated to cluster_size == 1: at cs > 1 the
DSMEM all-reduce already dominates and the leader path is kept.
cute.arch.shuffle_sync_up was measured to ignore its offset argument
(a lane-id prefix probe returned 2^k*lane); the scans therefore use
shuffle_sync with a computed source lane. A broadcast-LDS prefix
fallback was also tried and rejected: +40% executed instructions — more
than the barriers it saved.

Test Coverage

tests/scripts/cute_dsl_kernels/top_k/run_gvr_topk.py pytest sweep
(dtype x K x N x seed x next_n x T x V x warp-parallel-reduce):
256/256 PASS with the default (redundant) cadences, plus a
dedicated 12-case test_gvr_topk_decode_leader_paths that parametrizes
the two knobs (off/off, off/on, on/off) x cluster_size {1, 4} so the
leader paths and mixed configurations stay compiling and exact. These
are developer-run sweeps (tests/scripts is not CI-collected).
CI coverage: tests/unittest/_torch/attention/sparse/ test_cute_dsl_gvr_topk_decode.py (added by [None][perf] Cute DSL GVR Top-K: short-row remove cluster sync in run_one_row #15835) exercises this
kernel through the production op wrapper in the CI image. (It imports
the full tensorrt_llm runtime and could not be executed in the local
dev container — binding skew against the available prebuilt wheel; its
kernel-content coverage is subsumed by the direct-import sweep above.)
77-cell BS x N performance sweep with per-cell tie-aware correctness
(table above), both knob settings: 0 failures.
Determinism: redundant-warp replay is bit-identical across warps by
construction (same SMEM inputs, same reduction order, no atomics on the
decision path); the output ordering contract is unchanged.

PR Checklist

Please review the following before submitting your PR:

PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
If PR introduces API changes, an appropriate PR label is added - either api-compatible or api-breaking. For api-breaking, include BREAKING in the PR title.
Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
Update tava architecture diagram if there is a significant design change in PR.
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.

GitHub Bot Help

To see a list of available CI bot commands, please comment /bot help.

Known bounds (bf16/fp16)

The redundant-warp cadence and refinement chain are tuned for the fp32 histogram configuration (kNumBins=1024, lane-parallel bin search). The bf16/fp16 kNumBins=512 configuration takes the non-lane-parallel bin-search fallback and does not collect the chain's benefit: measured bf16 geomean 0.987x vs pre-PR base (worst -3.2% at N=131K, B200; neutral at N<=32K). A dtype gate was implemented and measured: knobs-OFF is worse (0.917x — the chain is structural; the redundant cadence is what keeps its added decision points cheap), so knobs-ON defaults are kept for all dtypes. Tracked follow-up: lane-parallelize the kNumBins=512 bin search so bf16/fp16 also collect the chain's benefit. Details: #16424 (comment)

…n converged snap P3 stream-write (dense+compact), P4 keys/vals accesses, and the snap-loop convergence check now go through raw integer addressing off a per-method iterator.toint() base (_smem_ld/_smem_st helpers) instead of tensor indexing, which re-derived the cluster SMEM window (S2R SR_CgaCtaId + LEA<<24) around hot accesses. P4 writeback: when the snap loop converged, the final iteration's cgt (s_iscalars[3]) is the exact strictly-greater count at sel_thr, so gt entries pack via counter[4] from 0 while ties pack via counter[5] from offset cgt — ONE candidate sweep and one barrier instead of two-pass. Same [gt | eq | pad] output partition; non-converged snap falls back to the original two-pass (its cgt would be stale). Measured (B200, real V4-Pro/Flash production shapes): P4 7.80->7.58us, wall +1-2% on top of fastsync (my-pro fs2 geomean 1.032->1.043, my-flash fs0 1.063->1.075). Full sweep 450/450 PASS. ncu follow-up: the residual snap-top S2R stall is a post-barrier pipeline-refill accounting artifact plus the un-hoisted scalar scratch accesses — see analysis notes. Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

The P4 snap loop steps one distinct candidate value per iteration (~0.45us: full candidate re-scan + 2 barriers). Real logits concentrate count mass at the k-th boundary, so the selected 1/1024 bin often holds tens of values; snap stragglers of 10-20us set the wall clock at N<=32K (P4 max 15.8us at N=8K vs 7.6us mean on synth aggregate cells). When the selected level-1 bin holds >8 candidates, rebuild the histogram over just that bin (+2% slop; bin width shrinks 1024x) and re-run the 3-step k-th search before entering the snap loop. Costs one extra candidate scan (~1us), leaves snap 0-2 steps. The snap loop converges monotonically from any start, so exactness is untouched. Histogram build + 3-step search factored into _hist_build / _kth_bin_search (trace-identical to the previous inline code). Synth sweep (B300, bs64 K=1024 aggregate cells, us): 2K 12.24->10.48 4K 15.14->12.97 8K 20.32->13.60 16K 24.88->17.07 32K 32.99->21.57 64K 42.35->39.99 128K/256K flat (P2-dominated). Worst bound-gap 2.09x -> 1.48x. Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

Two follow-ups to the level-2 histogram refinement (eed5eb5931): 1. Refinement chain: level-2 gate lowered 8->2 (one rebuild breaks even vs ~2 snap steps) and a level-3 pass added (gate >8, fires only when level 2 fails to split the bin: heavy ties or sub-ulp window). Unrolled 2-iteration chain over (gate, binw_cur). 2. _kth_bin_search step 1: every lane redundantly walked a bins_per_warp-deep serial LDS+IADD chain (~7% of stall samples at N=8K). Now each lane sums bins_per_warp/32 bins + one REDUX warp reduce; const_expr fallback to the serial walk when bins_per_warp is not a multiple of 32. Synth sweep (B300, bs64 K=1024 aggregate cells, us, vs eed5eb5931): 2K 10.48->9.10 4K 12.97->10.71 8K 13.60->12.42 16K 17.07->15.80 32K 21.57->21.25 64K 39.99->38.14. Worst bound-gap 1.48x -> 1.43x (128K, P2-straggler dominated); N<=32K all within 1.08-1.27x. Cumulative vs pre-refinement baseline: 8K 20.3->12.4 (-39%), 16K 24.9->15.8 (-36%), 32K 33.0->21.3 (-36%). Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

…back 1. Histogram window from the P2 exit bracket: P4's candidate min/max scan is replaced (behind a runtime guard) by [vlo, vhi) from s_thr[1]/[2] — P3 collected v >= vlo and cnt(>=vhi) < kK puts the k-th inside the window. Out-of-window values clamp into edge bins (cumulative-from-top counts stay exact) and the window is P2's acceptance band, so level-1 resolution improves. Stale-bracket paths (degenerate fallback, probe variants) fail the guard and take the original scan; a plausible-but-wrong bracket costs extra snap steps only, never exactness. 2. Converged-path writeback drops its SMEM atomics: the converged snap iteration's per-warp packed(ge,gt) counts are still staged in smem_wcnt and cover exactly this scan's per-warp element sets, so warps derive deterministic [gt | eq] output bases from a register prefix instead of ~2*cand/32 serialized atomicAdd claims. Padding count becomes cge (= s_iscalars[2]) on this path. Synth sweep (B300, bs64 K=1024): 4K 10.71->10.25, 8K 12.42->12.14, 16K 15.80->15.66, 32K 21.25->21.15; large N unchanged. Both changes individually small (~1-2%) — the ncu stall attribution on the min/max and writeback-atomic regions was largely pipeline-accounting; kept for the latency-chain shortening and the documented bracket contract. Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

Pre-existing hole surfaced by adversarial review: a candidate span > FLT_MAX (|v| ~ 1.7e38, fuzz-only for real logits) overflows range1 to +inf -> inv1 = +0 -> all candidates in bin 0 -> thr = NaN. NaN breaks the snap walk's ordering (all comparisons false), so it no-ops for snap_limit iterations and the row is emitted as pure padding. Clamping range1 to FLT_MAX keeps the start threshold ordered; snap's monotone walk rescues any ordered (even +/-inf) start. No perf change (8K spot 12.17us, smoke PASS). Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

Adversarial review of the cs>1 path (now the recommended config for N>=64K) surfaced a pre-existing theoretical data race from the multi-CTA commit: block_count_ge uses a single s_cluster_partial slot and a single cluster rendezvous per call, so a straggler CTA's post-wait DSMEM read of a peer's slot is unfenced against that peer's NEXT-call overwrite. If it ever fired, per-CTA totals diverge -> thresholds diverge -> peer under-collection (wrong top-k) or a cluster-barrier phase mismatch (hang). The schedule window is astronomically small (peer must complete six CTA barriers + classify + a full slice re-scan while the straggler's tid0 sits between wait and its peer loads), but under the PTX memory model it is a race. Fix with zero added synchronization: call k writes/reads slot k&1 (s_cluster_partial grows [1] -> [3]; slot 2 = tid0-private call counter, zeroed in run_one_row; CTAs call in lockstep so parity stays cluster-aligned). The peer's call-(k+2) overwrite of the same slot is transitively ordered after my call-k reads by the call-(k+1) rendezvous. cs2 spot 64K/128K/256K + smoke all PASS, perf unchanged. Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

After per-warp partials land in SMEM (one visibility barrier), EVERY warp redundantly performs the block reduce + decision a single leader used to do, in a fixed order — identical inputs and identical fp32 operation order make the result bit-identical across warps, so the publish barrier, the leader serialization, and the SMEM scalar re-reads (each costing a cluster-window S2R recompute in SASS) all disappear. Phase 4 (p4_warp_redundant, default on): - k-th bin search: 3 barriers -> 1; the serial 16-slot prefix and the 64-deep serial bin walk become lane-parallel (idx-shuffle inclusive scans + unique-crossing ballot); the selected threshold/count return in registers through level-2 refinement. - snap loop: 2 barriers/iter -> 1; threshold + convergence state in registers; warp partial staging parity double-buffered in smem_hist (a warp one iteration ahead cannot clobber a slot a slower warp still reads); the converged writeback reads the winning parity bank. Phase 2 (p2_warp_redundant, default on, cluster_size == 1): - one barrier per secant round (the staging barrier inside block_count_ge); classify + secant update replayed redundantly in registers by every warp; smem_wcnt staging parity-banked (2x num_warps); canonical exit state published once for Phase 3. Implementation notes from the falsification loop (B300, N=256K bs=64): - cute.arch.shuffle_sync_up IGNORES its offset (probed: a lane-id prefix scan returned 2^k*lane) — scans use shuffle_sync with a computed source lane instead. The interim broadcast-LDS prefix (64 predicated reads per thread per call) cost +40% instructions and erased the barrier win (ncu: barrier stall 14.1%->6.1% but 16.2M vs 11.5M inst). ncu (256K, bs64): barrier-stall 14.1% -> 6.1%; regs 64 -> 64, zero local-memory traffic both sides. A/B wall: 4K +2.6%, 32K +12%, 256K +3.3% (knobs off = leader paths, byte-identical). Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

coderabbitai · 2026-07-15T11:05:14Z

📝 Walkthrough

Walkthrough

The GVR Top-K decode kernel adds configurable redundant-warp execution for Phases 2 and 4, parity-buffered reduction scratch, hoisted shared-memory access helpers, optimized histogram/writeback paths, and driver coverage across redundant and leader configurations.

Changes

GVR Top-K redundant-warp execution

Layer / File(s)	Summary
Phase 2 redundant execution and scratch handling `tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py`	The kernel adds redundant-warp secant search, direct aggregate candidate counts, and parity-buffered shared-memory and cluster scratch handling.
Phase 3 and Phase 4 memory and snap paths `tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py`	Shared-memory bases are hoisted for candidate access; histogram helpers, redundant snap convergence, and deterministic writeback are added.
Runtime flags and path validation `tests/scripts/cute_dsl_kernels/top_k/run_gvr_topk.py`	Compilation and runtime dispatch expose both redundant-warp flags, with tests covering leader and redundant combinations for BF16 and FP32.
Estimated code review effort: 5 (Critical)	~90+ minutes

Suggested reviewers: kris1025, limin2021

Sequence Diagram(s)

sequenceDiagram
  participant gvr_topk_decode
  participant GvrTopKKernel
  participant phase2_secant_search
  participant phase4_histogram_snap
  participant output
  gvr_topk_decode->>GvrTopKKernel: compile with redundant-warp flags
  GvrTopKKernel->>phase2_secant_search: execute configured Phase 2 path
  phase2_secant_search->>phase4_histogram_snap: pass candidate state
  phase4_histogram_snap->>output: write top-K values and indices

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Title check	✅ Passed	The title is specific, concise, and matches the main PR theme: DSv4 GVR top-k kernel performance optimizations with redundant-warp sync reduction.
Description check	✅ Passed	The description includes the required sections and covers the change, tests, performance, and checklist items in sufficient detail.

✨ Finishing Touches

🧪 Generate unit tests (beta)

Create PR with unit tests

_{Comment @coderabbitai help to get the list of available commands.}

coderabbitai

Actionable comments posted: 2

🧹 Nitpick comments (1)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py (1)
529-570: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Annotate the new CuTe-JIT helpers
_smem_*, _hist_build, _kth_bin_search, and _kth_bin_search_rw still omit parameter and return annotations. Add the supported CuTe-JIT types used elsewhere in this file, or add a brief DSL exemption if these signatures are intentionally left dynamic.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py`
around lines 529 - 570, Add parameter and return annotations to the CuTe-JIT
helpers _smem_ld_f32, _smem_ld_i32, _smem_st_f32, _smem_st_i32, _hist_build,
_kth_bin_search, and _kth_bin_search_rw, using the supported CuTe-JIT types
already used elsewhere in the file. If these signatures must remain dynamic,
document a concise DSL exemption instead.
Source: Coding guidelines

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py`:
- Around line 1699-1705: In the bin-index calculation within the visible top-k
decode path, clamp the floating-point bin value bin_f to the valid [0, kBins -
1] range before converting it with cutlass.Int32. Preserve the existing integer
bounds checks as needed, and add a regression covering a narrow histogram window
with a very large positive candidate.

In `@tests/scripts/cute_dsl_kernels/top_k/run_gvr_topk.py`:
- Around line 179-180: Update the gvr_topk_decode() docstring’s Args section to
document p4_warp_redundant and p2_warp_redundant, including their True defaults
and how each controls redundant warps in Phase 4 and Phase 2 respectively.

---

Nitpick comments:
In `@tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py`:
- Around line 529-570: Add parameter and return annotations to the CuTe-JIT
helpers _smem_ld_f32, _smem_ld_i32, _smem_st_f32, _smem_st_i32, _hist_build,
_kth_bin_search, and _kth_bin_search_rw, using the supported CuTe-JIT types
already used elsewhere in the file. If these signatures must remain dynamic,
document a concise DSL exemption instead.

🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

Push a commit to this branch (recommended)
Create a new PR with the fixes

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 483e6126-b8e7-4281-a2a2-ac1ebe87f611

📥 Commits

Reviewing files that changed from the base of the PR and between e54fe55 and b8b616b.

📒 Files selected for processing (2)

tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py
tests/scripts/cute_dsl_kernels/top_k/run_gvr_topk.py

- _hist_build: clamp the bin index in the float domain (fmax/fmin) BEFORE the int cast — fptosi is undefined for out-of-range/NaN at the IR level even though PTX cvt.rzi saturates; fmax first canonicalizes NaN to bin 0. Edge-bin semantics for in-range values are unchanged. - run_gvr_topk.gvr_topk_decode: document p4_warp_redundant / p2_warp_redundant in the Args section. Made-with-Claude-Code: fable-5 Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com>

siyidNV · 2026-07-15T12:02:54Z

/bot run

siyidNV · 2026-07-16T04:52:12Z

/bot run

siyidNV · 2026-07-16T06:25:46Z

/bot run

longcheng-nv · 2026-07-16T06:28:52Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-16T06:34:57Z

PR_Github #59641 [ run ] triggered by Bot. Commit: a2ebe3a Link to invocation

tensorrt-cicd · 2026-07-16T11:34:48Z

PR_Github #59641 [ run ] completed with state SUCCESS. Commit: a2ebe3a
/LLM/main/L0_MergeRequest_PR pipeline #48079 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

siyidNV · 2026-07-16T12:39:42Z

/bot run --disable-fail-fast

longcheng-nv · 2026-07-16T13:00:08Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-16T13:14:20Z

PR_Github #59714 [ run ] triggered by Bot. Commit: c6ff467 Link to invocation

tensorrt-cicd · 2026-07-21T01:42:27Z

PR_Github #60471 [ run ] completed with state SUCCESS. Commit: 9b7f0bf
/LLM/main/L0_MergeRequest_PR pipeline #48803 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

siyidNV · 2026-07-21T03:07:43Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-21T03:20:43Z

PR_Github #60552 [ run ] triggered by Bot. Commit: 9b7f0bf Link to invocation

tensorrt-cicd · 2026-07-21T05:31:30Z

PR_Github #60552 [ run ] completed with state SUCCESS. Commit: 9b7f0bf
/LLM/main/L0_MergeRequest_PR pipeline #48869 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

siyidNV · 2026-07-21T05:43:26Z

/bot run --disable-fail-fast

longcheng-nv · 2026-07-21T06:07:45Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-21T06:13:37Z

PR_Github #60606 [ run ] triggered by Bot. Commit: 9269219 Link to invocation

…stogram-ladder Stack this PR on top of NVIDIA#16424 (semantic dependency): brings in the P4 histogram level-2 refinement (_hist_build/_kth_bin_search/_kth_bin_search_rw) and the p4/p2_warp_redundant sync reductions for the classic path. Conflict resolution (gvr_topk_decode.py, 3 hunks): - __init__ signature: union of both sides' new kwargs. - Insertion-point overlap between block reduce and phase4_histogram_snap: kept BOTH NVIDIA#16424's P4 helper methods and this branch's phase4_rank_scatter. - phase4_histogram_snap/phase2_secant_search bodies: taken from NVIDIA#16424 (this branch never modified them; R0 path uses phase4_rank_scatter). Validated standalone on B200: 24/24 exact vs torch.topk across {r0_default, classic, warp_redundant on/off} x {fp32,bf16} x {K512/cr4, K1024/cr4, K2048/cr1} x N in {8k,64k,128k}. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: longcheng-nv <243710427+longcheng-nv@users.noreply.github.com>

…Check) Formatting-only (AST-verified identical): 4 long-line rewraps inside phase4_rank_scatter that predate the NVIDIA#16424 merge, plus one duplicate blank line introduced by the merge commit. Root cause of the L0_MergeRequest_PR #48924 (and likely #48707) Release-Check failure; the two deleted-file paths in the CI ruff-format log were misleading output from the performance-summary printer. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: longcheng-nv <243710427+longcheng-nv@users.noreply.github.com>

tensorrt-cicd · 2026-07-21T09:41:32Z

PR_Github #60606 [ run ] completed with state SUCCESS. Commit: 9269219
/LLM/main/L0_MergeRequest_PR pipeline #48915 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

siyidNV · 2026-07-21T10:39:37Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-21T10:45:44Z

PR_Github #60677 [ run ] triggered by Bot. Commit: 694cd0e Link to invocation

longcheng-nv · 2026-07-21T12:08:04Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-21T12:13:51Z

PR_Github #60686 [ run ] triggered by Bot. Commit: 694cd0e Link to invocation

tensorrt-cicd · 2026-07-21T12:15:36Z

PR_Github #60677 [ run ] completed with state ABORTED. Commit: 694cd0e

Link to invocation

tensorrt-cicd · 2026-07-21T15:32:38Z

PR_Github #60686 [ run ] completed with state FAILURE. Commit: 694cd0e
/LLM/main/L0_MergeRequest_PR pipeline #48980 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

…id kernel Fixes the B300 CI IMA in test_lb_main_branches (CUDA illegal memory access, present since the R0-default commit; not introduced by the NVIDIA#16424 merge — reproduced on the pre-merge head). Root cause: GvrTopKLBKernel inlines a cluster_size>1 and a cluster_size=1 GvrTopKKernel instance into ONE launch, but the CuTe DSL sizes the launch's dynamic SMEM from the LAST-traced SmemAllocator only (CuTeDSL.track_smem_allocator holds a single slot). Making enable_r0 the default diverged the two instances' SMEM layouts — p1b_cache defaults to True at cs>1 but False for fp32 at cs=1, so the long branch (traced first, sized by the smaller single instance traced second) overflowed by the top_k*4B smem_gath buffer. compute-sanitizer: invalid __shared__ write at the reserved-size boundary; repro = fp32 K2048 all-long LB. Fix: - gvr_topk_decode.py: allocate s_cluster_partial / s_cluster_partial_m unconditionally (usage stays cs>1-gated; +16-32B at cs=1) so layouts no longer depend on cluster_size; new kc_diet ctor knob (None keeps the tuned cs=1 K512 diet) so LB can disable it symmetrically. - gvr_topk_decode_load_balance.py: pin p1b_cache=True and kc_diet=False for BOTH member instances + a constructor drift-guard asserting the smem-layout attributes match. Validation on B200: 11/11 LB configs exact (fp32/bf16 x K512/1024/2048 x all_long/all_short/mixed x next_n 1/2, incl. the crashing config), and the 24-case single-CTA smoke matrix stays exact. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: longcheng-nv <243710427+longcheng-nv@users.noreply.github.com>

longcheng-nv · 2026-07-22T01:08:17Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-22T01:27:21Z

PR_Github #60839 [ run ] triggered by Bot. Commit: 694cd0e Link to invocation

tensorrt-cicd · 2026-07-22T03:21:24Z

PR_Github #60839 [ run ] completed with state SUCCESS. Commit: 694cd0e
/LLM/main/L0_MergeRequest_PR pipeline #49113 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

siyidNV · 2026-07-22T06:39:15Z

/bot run --disable-fail-fast

siyidNV · 2026-07-22T07:31:37Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-07-22T07:39:00Z

PR_Github #60920 [ run ] triggered by Bot. Commit: 694cd0e Link to invocation

tensorrt-cicd · 2026-07-22T08:19:45Z

PR_Github #60920 [ run ] completed with state SUCCESS. Commit: 694cd0e
/LLM/main/L0_MergeRequest_PR pipeline #49187 completed with status: 'SUCCESS'

CI Report

Link to invocation

…undant-warp sync reduction (NVIDIA#16424) Signed-off-by: siyidNV <297196620+siyidNV@users.noreply.github.com> Co-authored-by: siyidNV <297196620+siyidNV@users.noreply.github.com> Co-authored-by: Claude <noreply@anthropic.com>

siyidNV and others added 7 commits July 15, 2026 02:00

siyidNV requested a review from a team as a code owner July 15, 2026 10:57

siyidNV requested a review from kris1025 July 15, 2026 10:57

github-actions Bot assigned siyidNV Jul 15, 2026

coderabbitai Bot reviewed Jul 15, 2026

View reviewed changes

Comment thread tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode.py Outdated

Comment thread tests/scripts/cute_dsl_kernels/top_k/run_gvr_topk.py

siyidNV mentioned this pull request Jul 15, 2026

[None][infra] Add blossom-ci authorized users #16425

Closed

1 task

coderabbitai Bot mentioned this pull request Jul 15, 2026

[None][perf] GVR top-K decode: enable R0 histogram-ladder admission by default #16457

Open

Merge branch 'main' into perf/gvr-kernel-opts

a2ebe3a

longcheng-nv requested review from lfr-0531, limin2021 and longcheng-nv and removed request for kris1025 July 16, 2026 05:59

Merge branch 'main' into perf/gvr-kernel-opts

c6ff467

siyidNV added 2 commits July 21, 2026 13:43

Merge branch 'main' into perf/gvr-kernel-opts

f485738

Merge branch 'main' into perf/gvr-kernel-opts

9269219

Merge branch 'main' into perf/gvr-kernel-opts

694cd0e

mingyangHao approved these changes Jul 22, 2026

View reviewed changes

lfr-0531 merged commit 9095cc1 into NVIDIA:main Jul 22, 2026
7 checks passed

Uh oh!

Conversation

siyidNV commented Jul 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Description

Changes

Performance

Notes for reviewers

Test Coverage

PR Checklist

GitHub Bot Help

Known bounds (bf16/fp16)

Uh oh!

coderabbitai Bot commented Jul 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Walkthrough

Changes

Sequence Diagram(s)

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

siyidNV commented Jul 15, 2026

Uh oh!

siyidNV commented Jul 16, 2026

Uh oh!

siyidNV commented Jul 16, 2026

Uh oh!

longcheng-nv commented Jul 16, 2026

Uh oh!

tensorrt-cicd commented Jul 16, 2026

Uh oh!

tensorrt-cicd commented Jul 16, 2026

Uh oh!

siyidNV commented Jul 16, 2026

Uh oh!

longcheng-nv commented Jul 16, 2026

Uh oh!

tensorrt-cicd commented Jul 16, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

siyidNV commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

siyidNV commented Jul 21, 2026

Uh oh!

longcheng-nv commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

siyidNV commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

longcheng-nv commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

tensorrt-cicd commented Jul 21, 2026

Uh oh!

longcheng-nv commented Jul 22, 2026

Uh oh!

tensorrt-cicd commented Jul 22, 2026

Uh oh!

tensorrt-cicd commented Jul 22, 2026

Uh oh!

siyidNV commented Jul 22, 2026

Uh oh!

siyidNV commented Jul 22, 2026

siyidNV commented Jul 15, 2026 •

edited

Loading

coderabbitai Bot commented Jul 15, 2026 •

edited

Loading