[None][feat] Multi-K (512/1024/2048) and Multi-dtype (fp32/bf16/fp16) GVR Top-K

[longcheng-nv]

Summary

Follow-up to #13477 (Scheme X dispatcher). Extends the GVR (Guess-Verify-Refine) Top-K micro-kernel to:

• Multi-K: 512 / 1024 / 2048 (was: 2048 only)
• Multi-dtype: fp32 / bf16 / fp16 logits (was: fp32 only)

via a GvrParams<T, K> trait class with 9 specializations. K=2048 fp32 production path is byte-identical SASS-wise (V2e LOCK preserved across all 3 production K=2048 kernel variants).

Commit Breakdown (16 commits)

1. feat: Multi-dtype GVR Top-K kernel: fp32/bf16/fp16
2. feat: Multi-dtype GVR Top-K dispatcher: bf16/fp16 path
3. feat: Multi-K (512/1024/2048) GVR Top-K via GvrParams trait class
4. fix: Forward TopK template arg into blockFusedSnapIter helpers
5. feat: Relax GVR Top-K caller gate to K ∈ {512, 1024, 2048}
6. feat: Per-(T, K) kNumBins tuning saves 5-9 % on K=512/1024 cells
7. feat: Defer dtype P3 conversion → −6 to −13 % on bf16/fp16 cells
8. feat: kFTarget retune for V4 M=K (K=512/1024 cells)
9. feat: Align standalone gvrTopKKernel launch_bounds with production
10. chore: Strip internal experiment-ID lineage from GVR Top-K comments
11. chore: Drop dead thresholdPos param + remove internal opt-registry tags
12. refactor: Hoist Scheme X bounds + collapse heuristicTopKDecode launcher
13. chore: Strip residual Opt-M / v0 references from selection-pass comment
14. feat: Symmetric Radix / Insertion fallback for indexer Top-K bf16/fp16 path
15. fix: GVR Top-K P4 snap-iter convergence + drop test tolerance widening (round-3 review fix)
16. chore: GVR Top-K drop dead blockFusedSnapIterDtype + clarify kFTarget docstring (round-3 review fix)

Key Files (6 files / +1621 / −355)

File	Change
cpp/tensorrt_llm/kernels/heuristic_topk.cuh	GvrParams<T,K> trait class; 9-way template specialization; per-(T,K) kNumBins / kFTarget tuning; deferred fp32 P3 conversion
cpp/tensorrt_llm/kernels/heuristicTopKDecode.cu	3 launcher overloads (fp32/bf16/fp16); GvrParams-driven kernel dispatch
cpp/tensorrt_llm/kernels/heuristicTopKDecode.h	Multi-dtype API surface
cpp/tensorrt_llm/kernels/indexerTopK.cu	Multi-dtype dispatcher; symmetric Insertion / Radix fallback for bf16/fp16
cpp/tensorrt_llm/kernels/IndexerTopK.h	dtype-templated invokeIndexerTopKDecode
cpp/tensorrt_llm/thop/IndexerTopKOp.cpp	Python op auto-routes by tensor.scalar_type()

API

No public API change. The Python op torch.ops.trtllm.indexer_topk_decode
now auto-routes by logits.scalar_type() to the matching dispatcher;
bf16/fp16 paths require pre_idx and heuristic_scratch (matching
dtype). K is selected via the index_topk argument as before; values
outside {512, 1024, 2048} fall back to Radix / Insertion.

Test plan

• Unit tests (pytest tests/unittest/_torch/thop/parallel/test_indexer_topk.py -k test_indexer_topk_decode):
  • test_indexer_topk_decode (5-arg / no pre_idx) — fp32 K∈{128, 2048} Radix/Insertion fallback covered (32 cases).
  • test_indexer_topk_decode_dist (7-arg / GVR Heuristic, parametrized in this PR over index_topk × dtype × num_tokens × batch_size × next_n × dist × success_ratio = 2268 cases) — all 9 new GVR (T, K) specializations covered. num_tokens=16384 cells hit GVR (1134); num_tokens=8192 cells route through Radix/Insertion fallback via Scheme X dispatcher (kSeqSmall=12288 gate). Post round-3 fix (commit 15): all 2268 cells pass under the strict default tolerance (1e-5) — flaky markers and the bin_tolerance widening are removed; the former 0.09 % K=512 / K=1024 intermittent at numRows ≈ 384 is fixed by the P4 snap-iter convergence guarantee.
• Q9 ablation suite on B200 sm_100 — multi-K × multi-dtype × multi-BS perf characterization with bit-equivalence to torch.topk set (9 combos × 9 SWE-Bench-64K layers × 12 BS cells).
• V2e LOCK SASS byte-identity verified (production K=2048 path).
• CUDA Graph capture / replay safe via the warmup helper introduced in [None][perf] Scheme X L2-aware dispatcher and PDL launchers for sparse-attention GVR Top-K #13477 (no new framework hooks; static caches in dispatcher are dtype-independent).

V2e LOCK SASS byte-identity (production K=2048)

Production kernel	md5 baseline / md5 post-PR
heuristicTopKMultiRowKernel<2048> fp32	a9f5c6b8 / a9f5c6b8 ✅
heuristicTopKMultiRowKernelDtype<__half, 2048>	9caa920c / 9caa920c ✅
heuristicTopKMultiRowKernelDtype<__nv_bfloat16, 2048>	aede08cc / aede08cc ✅

K=2048 production SASS is byte-identical with pre-PR baseline for all
three dtypes. New K=512 / K=1024 specializations have stable SASS produced
fresh by the GvrParams<T,K> template.

Performance Results (B200 sm_100, SWE-Bench-64K decode logits)

BS=1 cross-layer pooled (Q9d 04b — 9 layers × 9 GVR combos vs Radix fp32 K=2048 baseline, production path)

Routed through torch.ops.trtllm.indexer_topk_decode (Scheme X v1.2 dispatcher
→ heuristicTopKMultiRowKernel{,Dtype}<TopK> with preIdxOffset=+1, V3.2
decode semantics). 1 827 timed rows / variant (9 layers × stride-10 over rows
1..2024), 5 reps + 3 warmup, 128 MiB L2 flush before every launch.

Variant	Latency (µs)	Speedup vs Radix
Radix fp32 K=2048 (baseline)	50.98	1.00×
K=512 fp32	20.16	2.53×
K=512 bf16	17.15	2.97×
K=512 fp16	18.02	2.83×
K=1024 fp32	22.85	2.23×
K=1024 bf16	19.26	2.65×
K=1024 fp16	18.82	2.71×
K=2048 fp32	26.27	1.94×
K=2048 bf16	21.15	2.41×
K=2048 fp16	21.25	2.40×

GVR wins all 9 combos × 9 layers. K↓ and dtype↓ are both monotonic
(smaller K = smaller smem residency / higher occupancy; smaller dtype =
smaller HBM-load + smem footprint). Cross-validates the BS=1 column of
Table 2 (Q9e 05) within ±10 % wall — both tables are now production path.

BS scaling — pooled across 9 layers, last row replicated (Q9e 05, production path)

Speedup vs Radix fp32 K=2048 baseline (50.1 µs at BS=1 → 93.5 µs at BS=400):

Variant	BS=1	BS=8	BS=32	BS=128	BS=200	BS=256	BS=400
K=2048 fp32	1.82×	1.85×	1.82×	1.71×	1.74×	1.46×	1.28×
K=2048 bf16	2.22×	2.28×	2.28×	2.20×	2.31×	2.08×	2.09×
K=2048 fp16	2.13×	2.21×	2.19×	2.13×	2.22×	2.07×	2.05×
K=1024 bf16	2.30×	2.37×	2.33×	2.30×	2.38×	2.24×	2.24×
K=512 bf16	2.62×	2.70×	2.67×	2.54×	2.55×	2.34×	2.29×

bf16 / fp16 maintain 2.0 – 2.3× speedup across the full BS range
(1 .. 400). fp32 fades from 1.85× to 1.28× as BS grows past 200 (smem
residency tightens). Half-precision dtypes scale gracefully because
their lower smem footprint preserves occupancy beyond 1 wave.

Reproducibility

• BS=1 (Q9d 04b, production path): single-GPU nsys, ~3 min wall, 91 350 NVTX-tagged launches.
• BS scaling (Q9e): single-GPU nsys, ~3 min wall, 5 400 NVTX-tagged
  launches (9 GVR combos × 12 BS × 9 layers × 5 reps).
• All variants use 128 MiB L2 flush before every launch (cold-cache).

Round-3 review fixes (commit 15)

Addresses lfr-0531 and mingyangHao review comments on the 0.09 %
intermittent at K=512 / K=1024:

• Kernel (cpp/tensorrt_llm/kernels/heuristic_topk.cuh): P4 snap
  loop's snap_limit = cand_count/4 was insufficient for cells where
  the initial bin-edge threshold sits far from the true K-th value. Bug:
  on cgt >= kK non-convergence, Pass 1 emitted K elements in scan
  order from the cgt > kK strictly-greater set, missing some true
  top-K members. Fix: tighten snap_limit = cand_count, which is the
  provably-sufficient bound (each snap iter strictly decreases cgt or
  strictly increases cge by ≥ 1, so worst-case convergence takes
  cand_count − kK + 1 iters). Common path still breaks early at iter
  1–3, so steady-state perf is unchanged.
• Test (tests/unittest/_torch/thop/parallel/test_indexer_topk.py):
  remove pytest.mark.flaky(reruns=2) markers on K=512 / K=1024; drop
  the full_range/256 bin-quantized tolerance and the unsound
  |mean| * eps * 4 bf16/fp16 branch (bf16 → fp32 promotion inside the
  kernel is bit-exact and order-preserving, so no eps slack is
  warranted); restore compare_top_k_results default
  tolerance = 1e-5.

Additional cleanup in commit 4880bb6 (chore, no live SASS impact):

• Removed dead blockFusedSnapIterDtype helper (-64 lines) and stale doc paragraphs — the dtype P4 snap path uses the fp32 blockFusedSnapIter<TopK> directly after commit 7's deferred-conversion optimization stored smem keys[] as fp32.
• Expanded the kFTarget docstring to call out that the value is the secant solver's soft steering target (not the convergence condition), and that kFTarget < kK is intentional for small K — addresses the line-215 review question.

Local validation on B200 sm_100 with the fix + strict default
tolerance:

Scope	Cells	Result	Wall
K=512 + K=1024 (all dist / dtype / BS / next_n / ratio)	1512	1512/1512	137 s
K=2048 (same axes)	756	756/756	70 s
Stress: 5× focused corner (logistic m0.47, BS=128,	5×36	180/180	~5 min
next_n=3 → numRows=384)

0 failures across 2484 invocations under default strict tolerance.
The former 0.09 % intermittent corner is now stable.

Risks

• Smem ceiling at K=2048 + fp32 + BS=200..256: the K=512 fp32 path
  shows an anomalous BS=200→256 latency spike (29 → 77 µs).
  Production deployment is bounded by Scheme X dispatcher (kBsLarge
  derived from L2 capacity), so the cliff is not reached in
  serving traffic — but flagged for follow-up perf work.
• bf16 / fp16 path is decode-only: bf16 / fp16 prefill has not been
  added in this PR. Prefill remains fp32-only.

🤖 Generated with Claude Code

Summary by CodeRabbit

Release Notes

• New Features
  • Added support for bfloat16 and float16 data types in TopK decoding operations (previously fp32-only)
  • Extended topK parameter support to values of 512, 1024, and 2048
  • Introduced utility function to determine GVR heuristic optimization eligibility

[longcheng-nv]

/bot run --disable-fail-fast

[coderabbitai]

📝 Walkthrough

Walkthrough

This PR extends TensorRT-LLM's top-K decoding kernels to support bfloat16 and float16 inputs alongside existing float32 support. The changes introduce compile-time dtype specialization via traits, refactor the GVR algorithm to accept compile-time TopK values (512, 1024, 2048), and extend public APIs and dispatchers across IndexerTopK and HeuristicTopK subsystems.

Changes

Multi-dtype Top-K Kernel Extension

Layer / File(s)	Summary
Public API cpp/tensorrt_llm/kernels/IndexerTopK.h, cpp/tensorrt_llm/kernels/heuristicTopKDecode.h	Adds bf16/fp16 overload declarations for invokeIndexerTopKDecode and launchHeuristicTopKDecode, plus canIndexerTopKDecodeUseGvr(...) predicate.
Dtype Traits & Shared Memory cpp/tensorrt_llm/kernels/heuristic_topk.cuh	Introduces GvrDtypeTraits<T> specializations (float, bf16, fp16), GvrParams<T, TopK> compile-time constants per dtype/TopK pair, and templated KernelSmemTplK layout matching selected parameters.
Block Helpers cpp/tensorrt_llm/kernels/heuristic_topk.cuh	Templatizes blockCountGE, blockFusedSnapIter, and dtype-specific variants (blockCountGEDtype, blockFusedSnapIterDtype) to use K-dependent traits and shared-memory layouts.
GVR Algorithm cpp/tensorrt_llm/kernels/heuristic_topk.cuh	Refactors gvrTopKJob<TopK> and adds gvrTopKJobDtype<InputT, TopK> implementations; updates Phase 2–4 logic to use compile-time kCC/kBins/kK instead of fixed constants; deterministic partition with two-pass emission.
Global Kernels & Instantiations cpp/tensorrt_llm/kernels/heuristic_topk.cuh, cpp/tensorrt_llm/kernels/heuristicTopKDecode.cu	Defines gvrTopKKernel<TopK> and gvrTopKKernelDtype<InputT, TopK> with extern shared memory; explicit instantiations for 6 (dtype × TopK) combinations.
HeuristicTopK Dispatcher cpp/tensorrt_llm/kernels/heuristicTopKDecode.cu, cpp/tensorrt_llm/kernels/heuristic_topk.cuh	Implements launchHeuristicTopK and launchHeuristicTopKDecodeImpl<InputT> runtime dispatchers; validates topK against {512, 1024, 2048}; computes dtype-specific shared-memory size and selects kernel via switch(topK).
IndexerTopK Kernel Templating cpp/tensorrt_llm/kernels/indexerTopK.cu	Extends processHistogramStep, processBins, topKPerRowJob, topKPerRowDecode, and topKPerRowPrefill to accept InputT template parameter; converts bf16/fp16 to float at load sites.
IndexerTopK Dispatcher cpp/tensorrt_llm/kernels/indexerTopK.cu	Refactors Scheme X logic via new SchemeXBounds/getSchemeXBounds(...) helper; updates fp32 dispatch to use cached bounds; adds bf16/fp16 overloads with split-work restriction; implements canIndexerTopKDecodeUseGvr(...) predicate.
Torch Integration cpp/tensorrt_llm/thop/IndexerTopKOp.cpp	Discovers logits dtype, validates heuristic_scratch dtype matching, dispatches to appropriate invokeIndexerTopKDecode overload with reinterpret_casted bf16/half pointers and dtype-matched scratch buffers.

Estimated code review effort

🎯 4 (Complex) | ⏱️ ~60 minutes

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 12.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly and specifically describes the main changes: extending GVR Top-K to support multiple K values (512/1024/2048) and multiple data types (fp32/bf16/fp16).
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.
Description check	✅ Passed	The PR description is comprehensive and well-structured, covering summary, commit breakdown, key files, API changes, test plan with detailed test coverage, performance results with tables, round-3 review fixes, and identified risks.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

Tip

💬 Introducing Slack Agent: The best way for teams to turn conversations into code.

Slack Agent is built on CodeRabbit's deep understanding of your code, so your team can collaborate across the entire SDLC without losing context.

• Generate code and open pull requests
• Plan features and break down work
• Investigate incidents and troubleshoot customer tickets together
• Automate recurring tasks and respond to alerts with triggers
• Summarize progress and report instantly

Built for teams:

• Shared memory across your entire org—no repeating context
• Per-thread sandboxes to safely plan and execute work
• Governance built-in—scoped access, auditability, and budget controls

One agent for your entire SDLC. Right inside Slack.

👉 Get started

---

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

[coderabbitai]

🧹 Nitpick comments (2)

cpp/tensorrt_llm/kernels/heuristic_topk.cuh (2)

1720-1726: ⚡ Quick win

Use getEnvEnablePDL() for consistency.

heuristicTopKDecode.cu (line 222) uses tensorrt_llm::common::getEnvEnablePDL() while this launcher manually parses TRTLLM_ENABLE_PDL. Using the common utility ensures consistent behavior and avoids duplicating env-parsing logic.

♻️ Suggested fix

Add the include at the top of the file:

`#include` "tensorrt_llm/common/envUtils.h"

Then replace the manual parsing:

-    // Honor the standard TRTLLM_ENABLE_PDL env var (default on; set "0" to
-    // disable).
-    bool enablePDL = true;
-    if (char const* env = std::getenv("TRTLLM_ENABLE_PDL"))
-    {
-        if (env[0] == '0' && env[1] == '\0')
-            enablePDL = false;
-    }
+    bool const enablePDL = tensorrt_llm::common::getEnvEnablePDL();

🤖 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 `@cpp/tensorrt_llm/kernels/heuristic_topk.cuh` around lines 1720 - 1726,
Replace the manual getenv parsing that sets the local bool enablePDL with a call
to the shared utility tensorrt_llm::common::getEnvEnablePDL(); add the include
"tensorrt_llm/common/envUtils.h" to the top of the file, remove the block that
directly reads TRTLLM_ENABLE_PDL and instead initialize enablePDL by calling
tensorrt_llm::common::getEnvEnablePDL() to keep behavior consistent with
heuristicTopKDecode.cu.

---

1684-1686: 💤 Low value

Remove unused SmemKey type alias.

SmemKey is defined from the trait but never used — SmemT correctly uses float for keys as documented in the optimization note at lines 1173-1176. The unused alias adds confusion.

♻️ Suggested fix

 template <typename InputT, int TopK = TOP_K>
 __global__ void __launch_bounds__(BLOCK_SIZE)
     gvrTopKKernelDtype(InputT const* __restrict__ input, int const N, int const* __restrict__ preIdx, int const M,
         int const topK, InputT* __restrict__ outputValues, int* __restrict__ outputIndices)
 {
-    using SmemKey = typename GvrDtypeTraits<InputT>::SmemKey;
-    using SmemT
-        = KernelSmemTplK<float, GvrParams<InputT, TopK>::kC, GvrParams<InputT, TopK>::kNumBins>; // dtype keys fp32
+    // dtype path uses fp32 smem keys (deferred convert to output).
+    using SmemT = KernelSmemTplK<float, GvrParams<InputT, TopK>::kC, GvrParams<InputT, TopK>::kNumBins>;
     extern __shared__ unsigned char smem_raw[];

🤖 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 `@cpp/tensorrt_llm/kernels/heuristic_topk.cuh` around lines 1684 - 1686, Remove
the unused type alias SmemKey introduced from GvrDtypeTraits<InputT>; it is
never referenced and causes confusion—delete the line "using SmemKey = typename
GvrDtypeTraits<InputT>::SmemKey;" and keep the existing SmemT definition that
intentionally uses float keys (KernelSmemTplK<float, GvrParams<InputT,
TopK>::kC, GvrParams<InputT, TopK>::kNumBins>), leaving GvrDtypeTraits, InputT
and TopK untouched.

🤖 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.

Nitpick comments:
In `@cpp/tensorrt_llm/kernels/heuristic_topk.cuh`:
- Around line 1720-1726: Replace the manual getenv parsing that sets the local
bool enablePDL with a call to the shared utility
tensorrt_llm::common::getEnvEnablePDL(); add the include
"tensorrt_llm/common/envUtils.h" to the top of the file, remove the block that
directly reads TRTLLM_ENABLE_PDL and instead initialize enablePDL by calling
tensorrt_llm::common::getEnvEnablePDL() to keep behavior consistent with
heuristicTopKDecode.cu.
- Around line 1684-1686: Remove the unused type alias SmemKey introduced from
GvrDtypeTraits<InputT>; it is never referenced and causes confusion—delete the
line "using SmemKey = typename GvrDtypeTraits<InputT>::SmemKey;" and keep the
existing SmemT definition that intentionally uses float keys
(KernelSmemTplK<float, GvrParams<InputT, TopK>::kC, GvrParams<InputT,
TopK>::kNumBins>), leaving GvrDtypeTraits, InputT and TopK untouched.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 71e22bbe-2801-4613-8946-0f31f85d685a

📥 Commits

Reviewing files that changed from the base of the PR and between b26c000 and 17b5227.

📒 Files selected for processing (6)

• cpp/tensorrt_llm/kernels/IndexerTopK.h
• cpp/tensorrt_llm/kernels/heuristicTopKDecode.cu
• cpp/tensorrt_llm/kernels/heuristicTopKDecode.h
• cpp/tensorrt_llm/kernels/heuristic_topk.cuh
• cpp/tensorrt_llm/kernels/indexerTopK.cu
• cpp/tensorrt_llm/thop/IndexerTopKOp.cpp

[tensorrt-cicd]

PR_Github #47533 [ run ] triggered by Bot. Commit: 17b5227 Link to invocation

[tensorrt-cicd]

PR_Github #47533 [ run ] completed with state FAILURE. Commit: 17b5227
/LLM/main/L0_MergeRequest_PR pipeline #37450 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

[longcheng-nv]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #47544 [ run ] triggered by Bot. Commit: 17b5227 Link to invocation

[tensorrt-cicd]

PR_Github #47544 [ run ] completed with state FAILURE. Commit: 17b5227
/LLM/main/L0_MergeRequest_PR pipeline #37460 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

Link to invocation

[longcheng-nv]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #47548 [ run ] triggered by Bot. Commit: e5ad15e Link to invocation

[tensorrt-cicd]

PR_Github #47548 [ run ] completed with state SUCCESS. Commit: e5ad15e
/LLM/main/L0_MergeRequest_PR pipeline #37463 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

[longcheng-nv]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #47580 [ run ] triggered by Bot. Commit: e5ad15e Link to invocation

[tensorrt-cicd]

PR_Github #47652 [ run ] triggered by Bot. Commit: 607349e Link to invocation

[longcheng-nv]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #47653 [ run ] triggered by Bot. Commit: 0c9f705 Link to invocation

[tensorrt-cicd]

PR_Github #47653 [ run ] completed with state FAILURE. Commit: 0c9f705
/LLM/main/L0_MergeRequest_PR pipeline #37557 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

[longcheng-nv]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #47764 [ run ] triggered by Bot. Commit: 4880bb6 Link to invocation

[tensorrt-cicd]

PR_Github #47764 [ run ] completed with state SUCCESS. Commit: 4880bb6
/LLM/main/L0_MergeRequest_PR pipeline #37656 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

[longcheng-nv]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #47798 [ run ] triggered by Bot. Commit: 4880bb6 Link to invocation

[tensorrt-cicd]

PR_Github #47798 [ run ] completed with state SUCCESS. Commit: 4880bb6
/LLM/main/L0_MergeRequest_PR pipeline #37689 completed with status: 'SUCCESS'

CI Report

Link to invocation

[longcheng-nv]

Review Status Summary

All reviewer feedback has been addressed in commit 4880bb6. CI is fully green on this HEAD — two consecutive /bot run cycles (#47764 and #47798) both completed with SUCCESS, all 7 GitHub check-runs ✅, blossom-ci ✅.

Addressed threads

Reviewer	Thread	Resolution
@mingyangHao	heuristic_topk.cuh:227 — "Why for 512 topk, target is 384?"	Expanded kFTarget docstring with derivation rationale (4880bb6)
@mingyangHao	heuristic_topk.cuh — "function defined but never referenced"	Removed dead code: −64 lines body, −3 lines companion doc (4880bb6)
@mingyangHao	heuristic_topk.cuh:553 — unroll pragma on 16-iter loops	Verified -O3 nvcc already unrolls; SASS byte-identical with/without pragma — no change needed
@mingyangHao	test_indexer_topk.py — "will be zero when mean=0, dedicated design?"	Fixed unsound bf16/fp16 quantization-noise formula — now scales with value magnitude near the K-th boundary (4880bb6)
@lfr-0531	test_indexer_topk.py — "still accuracy issues at K=512/1024?"	Confirmed a real correctness bug (not tie-noise); GVR Top-K is now exact-match against same-dtype logits.topk
@lfr-0531	(same thread) — flaky(reruns=2) masking	Removed both flaky markers in d9be9d6; tests now pass deterministically

Status

• Mergeable: true (state: blocked — awaiting review approval)
• All 6 inline review threads have a response from the author; none re-opened
• Latest CI run on HEAD 4880bb600e: ✅ SUCCESS (2026-05-12 00:23 UTC)

@juney-nvidia could you take a look when you have a moment? @mingyangHao @lfr-0531 if my responses satisfactorily resolve your threads, please mark them resolved or leave any follow-up. Happy to address further comments. Thanks!

[yuxianq]

Approve the warmup_heuristic_topk_decode change in cpp_custom_ops.py.

[mingyangHao]

LGTM