diff --git a/docs/source/blogs/media/gvr_v2/README.md b/docs/source/blogs/media/gvr_v2/README.md
new file mode 100644
index 000000000000..c3c658cf6a7a
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/README.md
@@ -0,0 +1,188 @@
+---
+orphan: true
+---
+
+
+
+# GVR V2: Benchmark Methodology and Figure Reproduction
+
+This companion to [GVR V2: Self-Sampling and Multi-Thresholding for Faster Exact Top-K](../../tech_blog/blog29_GVR_V2_Self_Sampling_Exact_TopK_for_Sparse_Attention.md) contains the observations and definitions needed to reproduce the article's performance figures, plus the scope of its supplied temporal-overlap illustration and statistics. The article focuses on the algorithm and its performance; this document records the measurement boundaries.
+
+## Regenerate the Figures
+
+With NumPy and Matplotlib installed, run from the repository root:
+
+```bash
+python docs/source/blogs/media/gvr_v2/plot_results.py
+```
+
+The script regenerates `summary.json` and ten SVGs: `speedup.svg`, `evolution.svg`, `candidate_work.svg`, `algorithm.svg`, `gpu_sampling.svg`, `radix_cuda_map.svg`, `gvr_v1_map.svg`, `latency.svg`, `roofline.svg`, and `integration.svg`. It requires no GPU. The algorithm diagrams are schematic; every performance panel uses the bundled timing observations.
+
+The remaining figure, [temporal_overlap.svg](temporal_overlap.svg), is an author-supplied illustration converted directly from its standalone PDF. Its labels, axes, and plotted contents are preserved. It is not generated by `plot_results.py` or derived from the kernel-timing CSVs.
+
+## GPU Sampling Implementation
+
+`gpu_sampling.svg` illustrates the public device implementation at [the measured PR #19076 revision](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode_self_sampling.py). It is schematic and adds no measured performance claim.
+
+Both streaming families use packed/vectorized current-row windows. Let `d = SS2` be the integer stride in window units and `SMP` the number of windows. The score offsets below are relative to the aligned sampling base, with `0 <= j < SMP`:
+
+| Family | Window start | Loads and logical ownership |
+| :--- | :--- | :--- |
+| `main` | `8*j*d` | Work item `j` loads two adjacent `float4` vectors, covering eight FP32 scores |
+| `clus` | `16*j*d` | Work item `j` loads the lower eight scores and `j + SMP` loads the upper eight, using two `float4` vectors each |
+
+The clustered work items need not be adjacent lanes or distinct physical threads when the logical sample exceeds the CTA size. In both families, physical threads stride through additional work items in CTA-sized increments. Vector loads use `ld.global.nc.v4.f32`; the stride and count guards keep them inside the permitted interval. Prefill substitutes a valid score for leading alignment lanes before histogramming. The initial two vectors remain available for the sample histogram; extra work items are reloaded after sample extrema establish the bin scale. Eight retained sample scores are not the kernel's total register count.
+
+The [CUDA Best Practices Guide](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#coalesced-access-to-global-memory) documents 32-byte access granularity and the cost of sparse or misaligned accesses. A 32-byte-aligned eight-score window occupies one sector and a similarly aligned sixteen-score window occupies two. The supported API guarantees vector alignment, not universal 32/64-byte row alignment; a window offset by 16 bytes touches two or three sectors respectively. Cache reuse, warp issue patterns, and replay can further affect traffic. A 64-byte window is a software grouping, not a claim about cache-line size or a single memory instruction.
+
+The locality/coverage and register-lifetime explanation is an engineering interpretation of this implementation, not a measured 8-versus-16 ablation or evidence that either width is universally optimal. At fixed total sample count, the larger window has fewer distinct locations. Correlated values within a window do not supply independent random observations.
+
+The runtime geometry starts with a route-dependent budget proportional to `N / aim`, clamps the requested score count between 256 scores and half the row, then derives the integer window stride and count. Targets use the actual `S = SMP*8` or `S = SMP*16` population after rounding. For desired candidate population `A = aim`, the descending target ranks are approximately `A*S/N`, `K*S/N`, and `2*A*S/N`; the third target is implemented as twice the first integer target. The 256-score budget floor and 256 histogram bins are separate tuning parameters.
+
+Sample extrema reduce within each warp, pass through shared memory, and combine across warps after a CTA barrier to establish the 256 sample bins. Workers increment shared-memory counters, then synchronize before warp 0 calls `scan_cross0`. Each lane scans eight counters using two vector shared-memory loads and warp shuffles. This single scan extracts up to three rank crossings and clears the counters; a caller-side barrier publishes its results.
+
+Sample-bin lower edges define the primary threshold `T` near rank `A*S/N`, the upper anchor `T_K` near rank `K*S/N`, and the lower anchor `T_3` near rank `2*A*S/N`. For sample-bin width `w`, sampled `main` sets `HIC = max(T + 4*max(T_K - T, 0), T + 8*w)`: fourfold extrapolation of the nonnegative upper gap with at least eight bin widths of headroom. In `clus`, when a valid lower anchor satisfies `T_3 < T`, the upper gap is first capped at `2*(T - T_3)` before applying that same extrapolation and headroom. `HIC` tightens the classification bound; scores above it saturate into the top bin and remain candidates.
+
+Where enabled and valid, `T_3` supplies the lower admission floor `TSH`. Non-split streaming can retry with this floor after an aggressive primary threshold; the gated split-row route can stage down to the floor during its scan so the merged histogram already includes that wider admitted set. Sample misses, incomplete staging, and degenerate brackets still require exact verification or recovery. Non-sampled short-row and register routes retain their separate initialization policies.
+
+In variable-length `main`, warp 0 derives sampling geometry and publishes it through shared memory while the other warps issue register-free L2 prefetch hints for their CTA's upcoming row slice. Later register preloads and L2 hints are placed after the sample-reduction publication barrier, limiting overlap between live sample-reduction state and prefetched values. Cluster ranks repeat an identical full-row sample and prime their row slices after the corresponding barrier. Their sample histograms are CTA-local; subsequent verification histograms may be merged across the cluster. The prefetch placement is an implementation detail, not an independently measured contribution to speedup.
+
+## Candidate-Work Illustration
+
+`candidate_work.svg` is a schematic, independent of the timing observations. For a finite-score row with K-th-largest boundary $\tau$ and $q\le\tau$, the admitted population obeys $C_p=C(q)=K+E+D(q,\tau)$. Here $E=C(\tau)-K$ counts excess boundary ties and $D$ counts entries with $q\le x_i\lt\tau$. The simplified curve labels the three population segments directly, and the matching bar explains their roles without numerical ratios. Their proportions are illustrative, not benchmark measurements.
+
+Both panels are computed from the same finite synthetic row in `plot_results.py`. The tail-count axis is linear, and the stacked bar is proportional to population size. At the exact boundary, the filled marker includes ties and the open marker excludes them; their difference is the entire tied population, of which only the excess contributes to E.
+
+The admission band requires $K\le C(q)\le B_r$. A large tie plateau can leave no threshold in this band; the exact recovery path still applies. In V2, the admitted count describes candidate handling, while only the crossing bin requires the remaining exact selection. The illustration does not imply that every execution family materializes the same buffer or that candidate count alone predicts latency. The tail-count curve explains threshold admission; it does not prescribe a scalar search or a particular number of verification thresholds.
+
+## Temporal-Overlap Illustration
+
+Figure 2 shows SWE-bench-64K decode traces for DeepSeek-V3.2 and DeepSeek-V4 Pro. The temporal-hint index mapping shifts each prior index by +1 for V3.2 and leaves each compressed-bin index unchanged for V4 Pro. Overlap is the fraction of current Top-K indices covered by that mapped prior. Blue points are current selections matched by the mapped hint; orange points are current selections it misses. For V3.2, +1 predicts a one-position shift as decoding advances; this evaluates a prediction rule, not retention at identical token indices. The top panels show selected 1,024-position crops for layer 60. The bottom panels use the full index domain over 298 transitions per trace, without smoothing, for V3.2 layers 0/20/60 and Pro layers 2/22/60. Parenthesized legend values are mean overlaps.
+
+The figure illustrates temporal-hint variability, not kernel speedup. Its source Top-K streams are separate from the timing observations bundled below. `provenance.json` records the supplied SVG's checksum without private source paths or submission metadata.
+
+### Temporal-Overlap Statistics
+
+The table following Figure 2 uses author-supplied near-64K overlap summaries. Layer IDs are matched between SWE-bench and random-token inputs within each indexer. First average the adjacent-step hit ratio after temporal-hint index mapping over time for each layer; then compute statistics across those layer means, with equal weight per layer despite unequal decode durations. P10 and P90 are linearly interpolated empirical quantiles. V3.2 maps prior indices by +1; V4 retains the same compressed-bin indices.
+
+| Indexer | Input | Layers | Mean | Median | P10 | P90 | Min–max |
+| :--- | :--- | ---: | ---: | ---: | ---: | ---: | ---: |
+| V4 Pro | SWE-bench | 30 | 71.5% | 70.5% | 60.3% | 82.6% | 57.1–85.0% |
+| V4 Pro | Random tokens | 30 | 57.9% | 58.1% | 33.6% | 80.4% | 28.4–86.7% |
+| V4 Flash | SWE-bench | 21 | 62.8% | 61.5% | 53.4% | 73.3% | 53.2–83.7% |
+| V4 Flash | Random tokens | 21 | 52.8% | 49.2% | 30.7% | 74.7% | 25.9–81.3% |
+| V3.2 | SWE-bench | 61 | 47.4% | 47.8% | 37.9% | 59.7% | 5.7–70.7% |
+| V3.2 | Random tokens | 61 | 46.0% | 46.5% | 33.1% | 62.2% | 6.0–67.9% |
+
+These percentages preserve the supplied summaries' precision. They are independent of the timing CSVs and are not regenerated by `plot_results.py`. They summarize per-layer means, not individual-step extremes or the selected traces in Figure 2. Different K values and index-mapping rules make this evidence of hint variability, not a controlled ranking of models or a universal causal effect of prompt type. Low overlap concerns prediction quality and execution cost; exact verification and recovery preserve selection correctness.
+
+## Published Data
+
+| File | Contents |
+| :--- | :--- |
+| [flash_timings.csv.gz](flash_timings.csv.gz) | 2,079 DeepSeek-V4 Flash cases |
+| [pro_timings.csv.gz](pro_timings.csv.gz) | 2,970 DeepSeek-V4 Pro cases |
+| [v32_timings.csv.gz](v32_timings.csv.gz) | 4,697 DeepSeek-V3.2 cases |
+| [temporal_comparison.csv.gz](temporal_comparison.csv.gz) | One GVR V1 observation for each of the same 9,746 cases |
+| [provenance.json](provenance.json) | Published-file checksums, implementation labels, pairing, and roofline constants |
+| [summary.json](summary.json) | Recomputed statistics |
+| [plot_results.py](plot_results.py) | Figure generation |
+
+The CSVs begin with a copyright comment. `cell`, `model`, `isl_bucket`, and `layer` identify a workload. `batch`, `n`, and `k` specify its dimensions. Columns ending in `_us` contain case-level mean kernel durations in microseconds. Join the temporal supplement by `(cell, batch)`. Blank entries indicate missing or unsupported measurements, never zero latency. The loader maps blanks to `None` in both timing sources; statistics exclude missing observations, and complete-case comparisons exclude rows with any missing comparison timing. The full figure-regeneration command requires the bundled complete grid for all three implementations and rejects missing timings before writing outputs, preserving the article's common-coverage claims.
+
+The files contain kernel timings and workload dimensions. They do not contain input scores, prompt contents, or individual timing repetitions. They reproduce the published statistics and charts; repeating the GPU experiment requires suitable score inputs and a benchmark harness.
+
+## Measurement and Comparison Scope
+
+Measurements use NVIDIA B200, FP32 indexer scores, and batch sizes from 1 to 1,024. The GVR V2 reference covers 886 row geometries and 9,746 workload/batch cases, all passing tie-aware exactness checks. Each published GVR time is the arithmetic mean of 10 cold-L2 repetitions, rounded to 0.001 µs; five warm-L2 repetitions are excluded from these figures. A 512 MiB cache eviction runs outside the timed region. Compilation, input preparation, allocation during setup, and Python launch overhead are excluded; required device kernels remain timed.
+
+Each workload/batch case repeats one captured layer/step score row into distinct batch rows. This controls the input distribution and valid width while measuring batch scaling; it is not a heterogeneous batch of independent serving requests. GVR uses `next_n=1` and sets `max_seq_len` to that case's valid row length times its compression ratio. A serving graph may use a larger stable envelope and choose a different execution plan. The bundled grid does not independently benchmark ragged mixed-length batches, MTP, or prefill.
+
+The primary reference is the hint-free GVR V2 `run_varlen` implementation from [PR #19076](https://github.com/NVIDIA/TensorRT-LLM/pull/19076), measured at its [final public revision](https://github.com/NVIDIA/TensorRT-LLM/commit/be1b9885e8df9bf070e8cb68459e24a7119afaa9). Its FP32 comparisons with GVR V1 and TensorRT LLM radix CUDA match observations from separate runs by workload identity and batch size, with shape metadata checked where available. All three implementations cover the same 9,746 cases. Baseline times are retained as measured; no aggregate correction factor is applied. These comparisons describe the bundled kernel implementations and workloads, rather than the performance of entire serving frameworks.
+
+The device kernel and host dispatcher at the [main revision audited on September 17, 2026](https://github.com/NVIDIA/TensorRT-LLM/commit/73c70633b2547eedf0c91f85e760aec534f6af84) are byte-identical to those measured for the reference. This establishes source continuity for those files, not a new whole-framework performance measurement.
+
+Figures 1, 3, and 7–10 and their numerical summaries all use this PR #19076 reference. B300 measurements are outside the B200 comparison. Historical serving experiments retain their original implementation scope, as described below.
+
+| Implementation | Relevant comparison contract |
+| :--- | :--- |
+| GVR V2 | FP32 scores, valid row lengths, unordered INT32 indices |
+| GVR V1 | Temporal-prior pivot/rescue admission; complete implementation paired by workload and batch |
+| TensorRT LLM radix CUDA | Production dispatcher, including short-row insertion and long-row split-work paths |
+
+The GVR V1 implementation reference appears below. Complete build revisions for the historical radix observations are unavailable in the timing export.
+
+## GVR V1 and Algorithm Evolution
+
+The GVR V1 baseline is the tiered temporal implementation from public [PR #16877](https://github.com/NVIDIA/TensorRT-LLM/pull/16877), recorded in the `temporal_tiered_us` column. It already uses multi-threshold admission: sampled ladder counts choose a pivot and rescue rung, then a fused count/collect pass verifies both exactly. Its short-row and register routes have different execution strategies; Figure 3 sketches the streaming comparison, while its bars measure complete implementations.
+
+V2 changes calibration to packed/vectorized current-row windows and couples it to dense verification-bin counts and crossing-bin refinement; multi-thresholding itself is not a V2-only contribution. The GVR V1 comparison spans full implementations, including scheduling and integration; it is not an isolated self-sampling ablation. The public source at the article's implementation reference is [the GVR V1 kernel](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode_tp.py).
+
+The prior contract depends on the selected engine. At the measured revision, [`TopK.needs_gvr_prior`](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/modules/top_k.py#L69-L75) is true only for the temporal GVR path. V2 decode calls the self-sampling [`run_varlen`](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode_self_sampling_host.py#L1435-L1445), whose public signature has no `pre_idx`. The V1 [`tiered_topk`](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode_dispatch.py#L334-L357) dispatcher retains its required prior input. The self-sampling launcher fills an internal kernel argument slot named `pre_idx` with the output-index tensor; its [hint-free kernels do not read that slot](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode_self_sampling_host.py#L1598-L1612). This internal ABI slot is not a previous-step input to the V2 public entry point.
+
+Figure 3 computes each implementation's speedup directly from its paired radix times. The geometric means are 3.471858× for GVR V1 and 5.051864× for V2. The direct GVR V1/V2 time ratio is 1.455089×. `summary.json` records these under `evolution_vs_radix` and `temporal_vs_v2`.
+
+The direct V2 comparison also exposes the lower end of the measured speedup distribution:
+
+| Baseline | V2 speedup, geometric mean | V2 speedup, P5 | V2 wins / cases | V2 win rate |
+| :--- | ---: | ---: | ---: | ---: |
+| GVR V1 | 1.455089× | 1.095922× | 9,704 / 9,746 | 99.569054% |
+
+The P5 value is the fifth percentile across per-case `temporal_tiered_us / gvr_v2_us` mean-time ratios, not a runtime tail latency or a percentile of individual timing repetitions. It describes the bundled case distribution and does not establish a worst-case latency guarantee.
+
+## Aggregation and Coverage
+
+Speedup is the geometric mean of per-case `baseline_us / gvr_v2_us` ratios. Every case has equal weight. A win is a ratio strictly above one; minima and percentiles also use individual ratios of case-level mean durations. `_stats` in `plot_results.py` calls `numpy.percentile` without a method override, using its default linear interpolation between adjacent sorted ratios at fractional index `(cases - 1)*p/100` for percentile `p`. No slower case is discarded.
+
+Figure 1 uses the same cases for all three implementations within each model: 2,079 Flash, 2,970 Pro, and 4,697 V3.2 cases, totaling 9,746. V2 is fixed at 1.00; shorter bars mean less time. These values are recorded under `comparison_common_cases`. The article's overall and per-model tables use this same full paired coverage.
+
+The latency and roofline curves use arithmetic-mean durations over all captured layers at each row-length/batch point: 21 layers for Flash, 30 for Pro, and 61 for V3.2. All layers at each plotted point have the same valid width. The radix CUDA comparison heatmap (Figure 7) instead geometrically averages per-layer `radix_cuda_us / gvr_v2_us` ratios at each shape, using the same layers. It contains 275 cells: 99 each for Flash and Pro, and 77 for V3.2, covering all 11 batch sizes. The panels share a 1–21× color scale with parity at 1.0, and cell labels round to one decimal place. The cell values range from 1.522571× to 20.182551×, with no clipping by the color scale. A shape average can hide variation among individual cases.
+
+The GVR V1 comparison heatmap (Figure 8) uses the same 275 shapes and layer weights, with per-layer `temporal_tiered_us / gvr_v2_us` ratios. Its shared 1–3× scale is separate from the radix heatmap's scale; cell labels round to two decimal places. Shape averages range from 1.053828× to 2.865232× without clipping. Per-model minima/maxima are 1.058428×/2.865232× for Flash, 1.053828×/2.828985× for Pro, and 1.100878×/1.862122× for V3.2.
+
+All shape averages exceed one, but 15 shapes contain at least one slower constituent layer: one for Flash, six for Pro, and eight for V3.2. An orange corner identifies these cells using the unrounded per-layer ratios and a strict below-one comparison. There are 41 individual regressing cases overall; the markers indicate their presence, not their count or magnitude, and do not claim statistical significance. At Pro N=131,075 and B=512, the 30-layer geometric mean is 1.053828× while the minimum layer ratio is 0.688535×. These are the same observations used in the article's overall comparison.
+
+Correctness compares selected value multisets with `torch.topk`, allowing tied indices to differ. The capture-grid checks do not establish NaN ordering parity or a universal tie order. The linked implementation PRs additionally cover padding, variable lengths, exceptional values, and graph replay.
+
+## Additional Numerical Views
+
+The article uses Figure 1 for the model-level comparison. The following table uses the same full paired coverage and reports geometric-mean speedups of V2 over each baseline.
+
+| Baseline | V4 Flash, $K=512$ | V4 Pro, $K=1024$ | V3.2, $K=2048$ |
+| :--- | ---: | ---: | ---: |
+| GVR V1 | 1.53× | 1.54× | 1.37× |
+| TensorRT LLM radix CUDA | 4.88× | 4.87× | 5.25× |
+
+*Each model column uses the same workloads for both baselines.*
+
+For a concrete large-batch slice, the following times are at $B=1024$ and $N\approx131{,}072$, averaged over the same layers as Figure 9:
+
+| Model | GVR V2 | GVR V1 | Radix CUDA |
+| :--- | ---: | ---: | ---: |
+| V4 Flash | **113.1 µs** | 138.4 µs | 461.3 µs |
+| V4 Pro | **132.8 µs** | 146.1 µs | 477.7 µs |
+| V3.2 | **123.9 µs** | 166.0 µs | 496.7 µs |
+
+
+## Roofline Definitions
+
+The logical work is `W = B*N` abstract comparisons and the minimum traffic is `Q_min = 4*B*(N+K)` bytes, accounting for FP32 score reads and INT32 index writes. Operational intensity is `I = W/Q_min`; measured useful throughput is `P = B*N/(time_us*1e6)` Tcompare/s. Every kernel uses this same normalization. Temporal-prior accesses, padding, staging, repeated scans, and synchronization remain in measured time but do not enlarge the ideal work or minimum-traffic terms. The plotted throughput is not a hardware instruction rate or measured DRAM bandwidth.
+
+The full model is `min(R, BW*I)`. The theoretical parameters are `R=37.224960 Tcompare/s` and `BW=8 TB/s`; calibrated parameters are `R=37.047490 Tcompare/s` and `BW=6.912116 TB/s`. Their knees are 4.65312 and approximately 5.35979 compare/byte, above Top-K's ideal `[0.125, 0.25)` intensity range. The semantic comparison convention counts two binary comparisons per FMNMX3 result; it is not FP32 FLOPS.
+
+Figure 10B uses linear axes at B=1024. In this article, a Pareto curve denotes each operator's measured intensity–throughput trace across row lengths at that fixed batch. The points connect in intensity order; the curve is not a computed nondominated frontier or a search over configurations. Figure 9 also shows B=1, and the heatmaps in Figures 7–8 cover all 11 batches.
+
+Reachable rate is `100 * P / min(R, BW*I)` percent, using the calibrated roof. Compute each point from the arithmetic-mean duration across the same 21/30/61 layers used in Figure 10B for Flash/Pro/V3.2. The average reachable rate is the unweighted arithmetic mean of these point-level percentages, and the peak is their maximum. Flash and Pro each contribute nine intensity points; V3.2 contributes seven. The average is not weighted by row length or serving frequency. `summary.json` records the point counts and average/peak percentages under `roofline_reachable_rate`.
+
+Every plotted point lies on the bandwidth branch of the roof. With elapsed time `t` and consistent units, the fractional reachable rate simplifies to `(W/t)/(BW*W/Q_min) = Q_min/(BW*t)`; multiply by 100 for percent. This is the ideal minimum-traffic time `Q_min/BW` divided by measured time. The cancellation of `W` explains why the comparison convention does not change the reachable rate on this branch. Additional traffic and kernel work remain in `t`, so the rate does not measure actual DRAM bytes transferred or bandwidth utilization.
+
+The read-dominated roof is optimistic; mixed read/write behavior and additional kernel work can lower achievable throughput. Reachable rate describes useful selection work relative to this model, not measured DRAM bandwidth utilization.
+
+The input-read term applies to nontrivial selection with more valid scores than output slots. Short-row identity/padding paths can skip score reads; they are outside the measured grid and should not be evaluated against this traffic bound.
+
+## Serving Results
+
+Decode TPOT results come from public [PR #18410](https://github.com/NVIDIA/TensorRT-LLM/pull/18410); prefill results come from public [PR #18702](https://github.com/NVIDIA/TensorRT-LLM/pull/18702). These are separate historical serving experiments, not transformations of the operator speedups or measurements of the current complete framework. The 1.84–2.61× prefill-kernel range compares aggregate rank-0 Top-K kernel durations within measured Flash/Pro prefill windows; it is not the latency of a single Top-K invocation. The 2.9–4.5% throughput increase isolates adding V2 prefill to a deployment already using V2 decode on the tested long-input, batched configurations. It is not the throughput gain of replacing radix in both phases.
diff --git a/docs/source/blogs/media/gvr_v2/algorithm.svg b/docs/source/blogs/media/gvr_v2/algorithm.svg
new file mode 100644
index 000000000000..b5d0a5e14063
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/algorithm.svg
@@ -0,0 +1,762 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/candidate_work.svg b/docs/source/blogs/media/gvr_v2/candidate_work.svg
new file mode 100644
index 000000000000..49296172dedd
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/candidate_work.svg
@@ -0,0 +1,690 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/evolution.svg b/docs/source/blogs/media/gvr_v2/evolution.svg
new file mode 100644
index 000000000000..b25d7afddeb7
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/evolution.svg
@@ -0,0 +1,374 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/flash_timings.csv.gz b/docs/source/blogs/media/gvr_v2/flash_timings.csv.gz
new file mode 100644
index 000000000000..33c0b9e954f1
Binary files /dev/null and b/docs/source/blogs/media/gvr_v2/flash_timings.csv.gz differ
diff --git a/docs/source/blogs/media/gvr_v2/gpu_sampling.svg b/docs/source/blogs/media/gvr_v2/gpu_sampling.svg
new file mode 100644
index 000000000000..878e72799d96
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/gpu_sampling.svg
@@ -0,0 +1,469 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/gvr_v1_map.svg b/docs/source/blogs/media/gvr_v2/gvr_v1_map.svg
new file mode 100644
index 000000000000..3ad9fd7430c6
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/gvr_v1_map.svg
@@ -0,0 +1,1770 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/integration.svg b/docs/source/blogs/media/gvr_v2/integration.svg
new file mode 100644
index 000000000000..9df3baa85d07
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/integration.svg
@@ -0,0 +1,292 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/latency.svg b/docs/source/blogs/media/gvr_v2/latency.svg
new file mode 100644
index 000000000000..a16acb90cdb8
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/latency.svg
@@ -0,0 +1,1877 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/plot_results.py b/docs/source/blogs/media/gvr_v2/plot_results.py
new file mode 100644
index 000000000000..5833527d5541
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/plot_results.py
@@ -0,0 +1,1258 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+"""Regenerate the blog's figures and statistics from the bundled timing data.
+
+Requires matplotlib and numpy. Run from any directory; outputs stay beside this file.
+"""
+
+import csv
+import gzip
+import json
+from pathlib import Path
+from statistics import geometric_mean, mean, median
+
+import matplotlib
+
+matplotlib.use("Agg")
+import matplotlib.pyplot as plt
+import numpy as np
+from matplotlib.colors import Normalize
+from matplotlib.lines import Line2D
+from matplotlib.patches import FancyBboxPatch, Polygon, Rectangle
+from matplotlib.ticker import FuncFormatter
+
+ROOT = Path(__file__).resolve().parent
+COPYRIGHT = (
+ "Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. SPDX-License-Identifier: Apache-2.0"
+)
+ARMS = ["radix_cuda"]
+MODELS = {
+ "flash": "DeepSeek-V4 Flash · K=512",
+ "pro": "DeepSeek-V4 Pro · K=1024",
+ "v32": "DeepSeek-V3.2 · K=2048",
+}
+LABELS = {
+ "gvr_v2": "GVR V2",
+ "temporal_tiered": "GVR V1",
+ "radix_cuda": "TensorRT LLM radix CUDA",
+}
+COLORS = {
+ "gvr_v2": "#579600",
+ "temporal_tiered": "#386781",
+ "radix_cuda": "#64748b",
+}
+REACHABLE_BW = 6.912116
+TEMPORAL = ["temporal_tiered"]
+CROSS_CAMPAIGN = set(TEMPORAL + ARMS)
+COMPARISON_ARMS = ["gvr_v2", *TEMPORAL, *ARMS]
+
+
+def _load() -> list[dict]:
+ rows = []
+ for path in sorted(ROOT.glob("*_timings.csv.gz")):
+ with gzip.open(path, "rt") as handle:
+ for row in csv.DictReader(line for line in handle if not line.startswith("#")):
+ for name in ("batch", "n", "k", "layer"):
+ row[name] = int(row[name])
+ for name in list(row):
+ if name.endswith("_us"):
+ row[name] = float(row[name]) if row[name] else None
+ rows.append(row)
+ if len(rows) != 9746 or len({(r["cell"], r["batch"]) for r in rows}) != 9746:
+ raise ValueError("Expected exactly 9,746 unique cases")
+ with gzip.open(ROOT / "temporal_comparison.csv.gz", "rt") as handle:
+ historical = list(csv.DictReader(line for line in handle if not line.startswith("#")))
+ lookup = {(r["cell"], int(r["batch"])): r for r in historical}
+ if len(lookup) != len(rows) or len(historical) != len(rows):
+ raise ValueError("Temporal observations must cover the same 9,746 unique cases")
+ for row in rows:
+ old = lookup[(row["cell"], row["batch"])]
+ for arm in TEMPORAL:
+ row[arm + "_us"] = float(old[arm + "_us"]) if old[arm + "_us"] else None
+ return rows
+
+
+def _stats(rows: list[dict], arm: str, reference: str = "gvr_v2") -> dict:
+ valid = [r for r in rows if r[arm + "_us"] is not None]
+ ratios = [r[arm + "_us"] / r[reference + "_us"] for r in valid]
+ if not valid:
+ return {"cases": 0}
+ return {
+ "cases": len(valid),
+ "geomean": geometric_mean(ratios),
+ "minimum": min(ratios),
+ "p5": float(np.percentile(ratios, 5)),
+ "p95": float(np.percentile(ratios, 95)),
+ "wins": sum(x > 1 for x in ratios),
+ "win_percent": 100 * mean(x > 1 for x in ratios),
+ "baseline_median_us": median(r[arm + "_us"] for r in valid),
+ "gvr_median_us": median(r[reference + "_us"] for r in valid),
+ }
+
+
+def _save(fig: plt.Figure, name: str) -> None:
+ path = ROOT / (name + ".svg")
+ fig.savefig(
+ path,
+ bbox_inches="tight",
+ metadata={"Date": None, "Description": COPYRIGHT},
+ )
+ path.write_text("\n".join(line.rstrip() for line in path.read_text().splitlines()) + "\n")
+ plt.close(fig)
+
+
+def _comparison(rows: list[dict]) -> dict:
+ """Normalize every bar to V2 over one common case set per model."""
+ result = {}
+ for model in MODELS:
+ matched = _matching(rows, model)
+ result[model] = {
+ "cases": len(matched),
+ "layers": len({r["layer"] for r in matched}),
+ "latency_relative_to_v2": {
+ arm: _stats(matched, arm)["geomean"] for arm in COMPARISON_ARMS
+ },
+ }
+ return result
+
+
+def _overview(rows: list[dict]) -> None:
+ data = _comparison(rows)
+ fig, axes = plt.subplots(1, 3, figsize=(14, 4.7), sharey=True)
+ fig.subplots_adjust(left=0.185, right=0.97, bottom=0.25, top=0.70, wspace=0.14)
+ labels = [
+ "GVR V2",
+ "GVR V1",
+ "TRT-LLM radix CUDA",
+ ]
+ positions = [2.8, 1.6, 0.4]
+ for ax, (model, title) in zip(axes, MODELS.items()):
+ panel = data[model]
+ ax.axhspan(2.3, 3.3, color="#edf5df", zorder=0)
+ ax.axvline(1, color="#579600", alpha=0.55, linewidth=1, linestyle=(0, (2, 3)))
+ for y, arm in zip(positions, COMPARISON_ARMS):
+ value = panel["latency_relative_to_v2"][arm]
+ ax.barh(y, value, height=0.63, color=COLORS[arm], zorder=3)
+ ax.text(
+ value + 0.10,
+ y,
+ f"{value:.2f}×",
+ va="center",
+ fontsize=10.5,
+ weight="bold" if arm == "gvr_v2" else "normal",
+ color="#447a00" if arm == "gvr_v2" else "#334155",
+ )
+ name, kval = title.split(" · ")
+ ax.set_title(name, loc="left", fontsize=12.5, weight="bold", pad=28)
+ ax.text(
+ 0,
+ 1.025,
+ kval,
+ transform=ax.transAxes,
+ fontsize=9.5,
+ color="#52616f",
+ )
+ ax.set(xlim=(0, 5.95), ylim=(-0.15, 3.35), xticks=[0, 1, 2, 3, 4, 5])
+ ax.xaxis.set_major_formatter(FuncFormatter(lambda value, _: f"{value:g}×"))
+ ax.set_yticks(positions, labels, fontsize=10.5)
+ ax.tick_params(axis="both", length=0, pad=8)
+ ax.spines["left"].set_visible(False)
+ ax.spines["bottom"].set_color("#d5dce3")
+ ax.set_axisbelow(True)
+ ax.grid(axis="x", color="#e9edf1", linewidth=0.7)
+ axes[0].get_yticklabels()[0].set(color="#447a00", weight="bold")
+ fig.text(
+ 0.035,
+ 0.935,
+ "The GVR evolution: V1, V2, and radix CUDA",
+ fontsize=20,
+ weight="bold",
+ color="#17202b",
+ )
+ fig.text(
+ 0.035,
+ 0.875,
+ "Geometric-mean kernel time relative to GVR V2 · shorter is faster · B200 / FP32",
+ fontsize=11,
+ color="#52616f",
+ )
+ fig.text(
+ 0.185,
+ 0.13,
+ "Same workloads within each panel. GVR V2 (PR #19076) = 1.00×.",
+ fontsize=10,
+ color="#334155",
+ )
+ fig.text(
+ 0.035,
+ 0.067,
+ "GVR V1: temporal hint calibration. GVR V2: current-row self-sampling.",
+ fontsize=9,
+ color="#52616f",
+ )
+ _save(fig, "speedup")
+
+
+def _box(
+ ax: plt.Axes,
+ xy: tuple[float, float],
+ size: tuple[float, float],
+ text: str,
+ color: str = "#edf5df",
+ fontsize: float = 11,
+) -> None:
+ ax.add_patch(
+ FancyBboxPatch(
+ xy,
+ *size,
+ boxstyle="round,pad=0.08,rounding_size=0.08",
+ facecolor=color,
+ edgecolor="#ccd5dd",
+ linewidth=0.8,
+ )
+ )
+ ax.text(
+ xy[0] + size[0] / 2,
+ xy[1] + size[1] / 2,
+ text,
+ ha="center",
+ va="center",
+ fontsize=fontsize,
+ color="#17202b",
+ linespacing=1.5,
+ )
+
+
+def _evolution(rows: list[dict]) -> None:
+ fig = plt.figure(figsize=(14, 5.6))
+ ax = fig.add_axes((0.02, 0.16, 0.61, 0.69))
+ ax.set(xlim=(0, 10), ylim=(0, 6))
+ ax.axis("off")
+ ax.text(
+ 0.2,
+ 5.55,
+ "GVR V1 (temporal hint): calibration + multi-thresholding",
+ fontsize=11,
+ weight="bold",
+ color=COLORS["temporal_tiered"],
+ )
+ for x, label in [
+ (0.2, "Previous indices\n→ current-score gather"),
+ (3.6, "Hint-derived thresholds\n→ multiple exact counts"),
+ (7.0, "Collect candidates\n→ exact refinement"),
+ ]:
+ _box(ax, (x, 3.8), (2.8, 1.2), label, "#eef3f7", 10)
+ ax.text(
+ 0.2,
+ 3.25,
+ "Biased toward previous winners; overlap varies across layers and steps.\n"
+ "Weak hints add verification and recovery work.",
+ fontsize=9.5,
+ color="#475569",
+ )
+ ax.text(
+ 0.2,
+ 2.55,
+ "GVR V2 streaming: self-sampling + multi-thresholding",
+ fontsize=11,
+ weight="bold",
+ color="#447a00",
+ )
+ for x, label in [
+ (0.2, "Coalesced current-row\nsample → tail bracket"),
+ (3.6, "Full-row classification\n→ many exact counts"),
+ (7.0, "Emit certain winners\n→ refine crossing bin"),
+ ]:
+ _box(ax, (x, 0.8), (2.8, 1.2), label, fontsize=10)
+ for y in (4.4, 1.4):
+ for x in (3.02, 6.42):
+ ax.annotate(
+ "",
+ (x + 0.5, y),
+ (x, y),
+ arrowprops={"arrowstyle": "->", "color": "#64748b", "lw": 1.5},
+ )
+ ax.text(
+ 0.2,
+ 0.23,
+ "No temporal-overlap dependency; one calibration rule for both phases.",
+ fontsize=9.5,
+ color="#447a00",
+ )
+ bars = fig.add_axes((0.76, 0.30, 0.21, 0.4))
+ arms = [*TEMPORAL, "gvr_v2"]
+ for i, arm in enumerate(arms):
+ ratios = [r["radix_cuda_us"] / r[arm + "_us"] for r in rows]
+ value = geometric_mean(ratios)
+ bars.barh(i, value, height=0.52, color=COLORS[arm])
+ bars.text(value + 0.1, i, f"{value:.2f}×", va="center", weight="bold", fontsize=12)
+ bars.set_yticks(
+ range(len(arms)),
+ [LABELS[arm] for arm in arms],
+ fontsize=10,
+ )
+ bars.invert_yaxis()
+ bars.set_xlim(0, 6)
+ bars.set_xticks([0, 1, 2, 3, 4, 5])
+ bars.set_xlabel("Speedup over radix CUDA", fontsize=10)
+ bars.set_title("Measured evolution", fontsize=12, loc="left", weight="bold", pad=18)
+ bars.grid(axis="x", alpha=0.12)
+ bars.set_axisbelow(True)
+ fig.suptitle(
+ "From temporal prediction to current-row calibration",
+ x=0.03,
+ ha="left",
+ fontsize=20,
+ weight="bold",
+ y=0.98,
+ )
+ fig.text(
+ 0.035,
+ 0.06,
+ "Design goals: improve the practical performance floor and average latency; "
+ "remove the temporal prior's framework lifecycle.",
+ fontsize=10,
+ color="#475569",
+ )
+ fig.text(
+ 0.035,
+ 0.005,
+ "Flows show streaming paths; bars compare complete GVR V1 and GVR V2 implementations.",
+ fontsize=9,
+ color="#475569",
+ )
+ _save(fig, "evolution")
+
+
+def _candidate_work() -> None:
+ """Illustrate tail counts and candidate amplification without measured data."""
+ fig = plt.figure(figsize=(16.7, 7.5), facecolor="white")
+ ink, muted = "#17202b", "#52616f"
+ green, orange, rose = "#447a00", "#b56b0b", "#ad4b70"
+ canvas = fig.add_axes((0, 0, 1, 1))
+ canvas.set(xlim=(0, 1), ylim=(0, 1))
+ canvas.axis("off")
+ for x, width in ((0.025, 0.47), (0.52, 0.455)):
+ canvas.add_patch(
+ FancyBboxPatch(
+ (x, 0.205),
+ width,
+ 0.65,
+ boxstyle="round,pad=0.008,rounding_size=0.014",
+ facecolor="#f8fafc",
+ edgecolor="#dce3e9",
+ linewidth=0.9,
+ )
+ )
+ fig.text(
+ 0.04, 0.947, "Threshold quality sets candidate work", fontsize=23, weight="bold", color=ink
+ )
+ fig.text(
+ 0.04,
+ 0.90,
+ "A tighter threshold reduces extra candidates; exact counts keep admission safe.",
+ fontsize=12.5,
+ color=muted,
+ )
+ fig.text(
+ 0.044, 0.80, "A Choose an admission threshold", fontsize=15.5, weight="bold", color=ink
+ )
+ fig.text(
+ 0.54, 0.80, "B Explain the admitted population", fontsize=15.5, weight="bold", color=ink
+ )
+
+ # One finite row defines both panels, including the left-continuous tie jump.
+ scores = np.array([1.4, 3.0, 5.8, 7.5, 9.0])
+ multiplicities = np.array([70, 80, 70, 70, 60])
+ row = np.repeat(scores, multiplicities)
+ k, q, tau = 100, 4.4, 7.5
+ admitted = int(np.count_nonzero(row >= q))
+ at_boundary = int(np.count_nonzero(row >= tau))
+ above_boundary = int(np.count_nonzero(row > tau))
+ excess = at_boundary - k
+ shell = admitted - at_boundary
+
+ ax = fig.add_axes((0.089, 0.32, 0.335, 0.425), facecolor="#f8fafc")
+ ax.set(xlim=(0, 10), ylim=(0, 3.9))
+ ax.axhspan(1, 2.8, color="#edf4e5", zorder=0)
+ for level in (1, 2.8):
+ ax.axhline(level, color="#adc096", linewidth=1, linestyle=(0, (4, 4)))
+ thresholds = np.r_[0, scores, 10]
+ counts = np.array([np.count_nonzero(row >= t) / k for t in thresholds])
+ ax.step(thresholds, counts, where="pre", color=ink, linewidth=2.3, zorder=3)
+ ax.text(9.7, 3.28, "Over capacity", color=muted, fontsize=11.5, ha="right")
+ ax.text(0.25, 0.25, "Too few candidates", color=muted, fontsize=11.5)
+ ax.vlines(q, 0, admitted / k, color=green, linewidth=1.2, linestyles=(0, (3, 3)))
+ ax.scatter([q], [admitted / k], s=110, color=green, edgecolor="white", linewidth=1.4, zorder=5)
+ ax.text(
+ 4.6,
+ 2.31,
+ r"$C_p=C(q)$",
+ color=green,
+ fontsize=15,
+ ha="center",
+ )
+ ax.vlines(tau, 0, at_boundary / k, color=rose, linewidth=1.1, linestyles=(0, (3, 3)))
+ ax.scatter([tau], [at_boundary / k], s=42, color=rose, zorder=6)
+ ax.scatter([tau], [above_boundary / k], s=35, facecolor="white", edgecolor=rose, zorder=6)
+ ax.text(
+ 7.6,
+ 1.56,
+ r"$C(\tau)$",
+ ha="center",
+ fontsize=12.5,
+ color=rose,
+ )
+ # Brackets split C(q), not the entire tie jump: K cuts through that jump.
+ for start, stop, color, symbol in (
+ (0, 1, green, r"$K$"),
+ (1, at_boundary / k, rose, r"$E$"),
+ (at_boundary / k, admitted / k, orange, r"$D$"),
+ ):
+ ax.plot(
+ [10.14, 10.38, 10.38, 10.14],
+ [start, start, stop, stop],
+ color=color,
+ linewidth=1.8,
+ clip_on=False,
+ )
+ ax.text(
+ 10.72,
+ (start + stop) / 2,
+ symbol,
+ color=color,
+ fontsize=16,
+ va="center",
+ clip_on=False,
+ )
+ ax.hlines(
+ [at_boundary / k, admitted / k],
+ [tau, q],
+ [10.14, 10.14],
+ colors=[rose, green],
+ linewidth=0.9,
+ linestyles=(0, (3, 3)),
+ clip_on=False,
+ )
+ ax.set_yticks([0, 1, 2.8], ["0", r"$K$", r"$B_r$"])
+ ax.set_xticks([q, tau], [r"$q$", r"$\tau$"])
+ for tick, color in zip(ax.get_xticklabels(), (green, rose)):
+ tick.set_color(color)
+ ax.set_ylabel(r"Candidates $C(t)$", fontsize=12.5, labelpad=10)
+ ax.set_xlabel("Higher threshold →", fontsize=12, labelpad=7)
+ ax.tick_params(length=0, pad=7, labelsize=13)
+ ax.spines["left"].set_color("#cbd5e1")
+ ax.spines["bottom"].set_color("#cbd5e1")
+ fig.text(
+ 0.089,
+ 0.223,
+ r"Safe admission: $K\leq C(q)\leq B_r$",
+ fontsize=12,
+ color=green,
+ )
+
+ fig.text(0.55, 0.732, "CANDIDATE AMPLIFICATION", fontsize=10.5, weight="bold", color=muted)
+ fig.text(0.55, 0.666, r"$C_p\,/\,K$", fontsize=26, color=green)
+ fig.text(0.665, 0.678, "Candidates per required winner", fontsize=12.5, color=ink)
+ right = fig.add_axes((0.55, 0.564, 0.392, 0.072))
+ right.set(xlim=(0, admitted), ylim=(0, 1))
+ right.axis("off")
+ segments = [
+ (k, "#deedc8", green, r"$K$", "Required output"),
+ (excess, "#f3dce5", rose, r"$E$", r"Excess boundary ties: $C(\tau)-K$"),
+ (shell, "#fae6c7", orange, r"$D$", r"Boundary shell: $q\leq x<\tau$"),
+ ]
+ left = 0
+ for i, (width, fill, color, symbol, label) in enumerate(segments):
+ right.add_patch(
+ Rectangle((left, 0), width, 1, facecolor=fill, edgecolor="white", linewidth=2)
+ )
+ right.text(
+ left + width / 2, 0.5, symbol, fontsize=21, ha="center", va="center", color=color
+ )
+ y = 0.499 - i * 0.079
+ fig.text(0.555, y, symbol, fontsize=18, color=color, va="center")
+ fig.text(0.589, y, label, fontsize=13, color=ink, va="center")
+ canvas.plot([0.55, 0.943], [y - 0.037, y - 0.037], color="#e1e7ec", lw=0.8)
+ left += width
+ fig.text(0.746, 0.24, r"$C_p=C(q)=K+E+D(q,\tau)$", fontsize=19, ha="center", color=ink)
+
+ canvas.add_patch(
+ FancyBboxPatch(
+ (0.027, 0.07),
+ 0.941,
+ 0.096,
+ boxstyle="round,pad=0.008,rounding_size=0.013",
+ facecolor="#edf4e5",
+ edgecolor="none",
+ )
+ )
+ fig.text(0.045, 0.113, "V2", fontsize=17, weight="bold", color=green)
+ fig.text(0.11, 0.112, r"$C(q)$ Candidate handling", fontsize=14, color=ink)
+ canvas.annotate(
+ "",
+ (0.624, 0.108),
+ (0.413, 0.108),
+ arrowprops={"arrowstyle": "->", "color": green, "lw": 1.4},
+ )
+ fig.text(0.515, 0.132, "Exact bin counts", fontsize=10.5, color=green, ha="center")
+ fig.text(0.65, 0.112, r"$m$ Crossing-bin refinement", fontsize=14, color=green)
+ fig.text(
+ 0.04,
+ 0.025,
+ "Schematic finite-score example · population counts only",
+ fontsize=10.5,
+ color=muted,
+ )
+ _save(fig, "candidate_work")
+
+
+def _algorithm() -> None:
+ fig, ax = plt.subplots(figsize=(14, 9))
+ ax.set(xlim=(0, 14), ylim=(0, 7))
+ ax.axis("off")
+ headings = [
+ (0.2, "1 SELF-SAMPLE", "Place the bracket near the current tail"),
+ (4.9, "2 MULTI-THRESHOLD", "Get many exact counts from one classification"),
+ (9.65, "3 REFINE", "Finish only the uncertain boundary"),
+ ]
+ for x, heading, subtitle in headings:
+ ax.text(x, 6.5, heading, fontsize=14, weight="bold", color="#447a00")
+ ax.text(x, 6.04, subtitle, fontsize=9.3, color="#475569")
+ ax.text(0.2, 5.5, "Current row: sparse, regularly spaced vector loads", fontsize=9)
+ for i in range(32):
+ sampled = i % 8 < 2
+ ax.add_patch(
+ Rectangle(
+ (0.2 + i * 0.126, 4.85), 0.112, 0.36, facecolor="#76b900" if sampled else "#dfe5eb"
+ )
+ )
+ ax.text(0.2, 4.46, "Sample histogram ≈ current score distribution", fontsize=9)
+ sample = [1, 3, 5, 8, 11, 14, 12, 9, 6, 3, 2, 1]
+ for i, h in enumerate(sample):
+ ax.add_patch(Rectangle((0.25 + i * 0.32, 2.7), 0.28, h * 0.085, facecolor="#a7c976"))
+ for x, label, y in [(2.37, "T_floor", 2.16), (2.69, "T", 2.45), (3.33, "T_K", 2.16)]:
+ ax.plot([x, x], [2.66, 3.9], color="#7557a6", linestyle="--", linewidth=1)
+ ax.text(x, y, label, ha="center", fontsize=10, color="#7557a6")
+ ax.text(
+ 0.2,
+ 1.45,
+ "Ranks ≈ 2AS/N, AS/N, KS/N\n→ safety floor, admission threshold, upper anchor",
+ fontsize=9.3,
+ linespacing=1.5,
+ )
+ ax.annotate(
+ "", (4.7, 3.6), (4.2, 3.6), arrowprops={"arrowstyle": "->", "lw": 2, "color": "#64748b"}
+ )
+ _box(ax, (4.95, 4.8), (4.05, 0.65), "Every valid score is examined", "#e9eff5", 11)
+ ax.text(4.9, 4.25, "Exact histogram → suffix counts at all bin boundaries", fontsize=9)
+ counts = [1000, 700, 400, 250, 150, 100, 80, 73, 380, 330, 270]
+ for i, count in enumerate(counts):
+ color = "#f5b642" if i == 7 else ("#579600" if i > 7 else "#cbd5e1")
+ ax.add_patch(Rectangle((5.0 + i * 0.36, 2.7), 0.31, 0.18 + count / 800, facecolor=color))
+ ax.text(5.0, 2.35, "T", fontsize=10)
+ ax.text(8.9, 2.35, "H", fontsize=10)
+ ax.annotate(
+ "73 in crossing bin",
+ (7.7, 3.02),
+ (6.0, 3.65),
+ fontsize=9,
+ arrowprops={"arrowstyle": "->", "color": "#8c661c"},
+ color="#8c661c",
+ )
+ ax.text(8.43, 3.82, "980\nabove", fontsize=10, ha="center", color="#447a00")
+ ax.text(
+ 4.9,
+ 1.45,
+ "256 verification bins (shown schematically)\nOne bin assignment per survivor, then an on-chip scan",
+ fontsize=9.3,
+ linespacing=1.5,
+ )
+ ax.annotate(
+ "", (9.5, 3.6), (9.05, 3.6), arrowprops={"arrowstyle": "->", "lw": 2, "color": "#64748b"}
+ )
+ _box(ax, (9.75, 4.45), (3.75, 0.85), "980 certain winners", fontsize=13)
+ _box(ax, (9.75, 3.0), (3.75, 0.85), "Select 44 of 73 boundary candidates", "#fff3d9", 10.5)
+ _box(ax, (9.75, 1.55), (3.75, 0.85), "1,024 exact output indices", fontsize=12)
+ for y in (4.05, 2.6):
+ ax.annotate(
+ "",
+ (11.6, y - 0.1),
+ (11.6, y + 0.25),
+ arrowprops={"arrowstyle": "->", "color": "#64748b"},
+ )
+ fig.subplots_adjust(left=0.01, right=0.99, bottom=0.27, top=0.98)
+ guard = fig.add_axes((0.025, 0.03, 0.95, 0.26))
+ guard.set(xlim=(0, 14), ylim=(0, 3))
+ guard.axis("off")
+ guard.text(
+ 0.05, 2.8, "EXACTNESS CHECKS BEFORE OUTPUT", fontsize=12, weight="bold", color="#334155"
+ )
+ _box(
+ guard,
+ (3.2, 1.96),
+ (7.6, 0.44),
+ "Full-row coverage · valid bracket · complete candidates",
+ "#f1f4f7",
+ 11,
+ )
+ paths = [
+ (0.15, "Enough survivors\nRefine the crossing → exact Top-K", "#edf5df"),
+ (4.85, "Too few survivors\nLower admission and verify again", "#fff3d9"),
+ (9.55, "Overflow or unusable bracket\nExact complete-set / whole-row recovery", "#e9eff5"),
+ ]
+ for x, label, color in paths:
+ _box(guard, (x, 0.24), (4.15, 0.98), label, color, 10)
+ guard.annotate(
+ "",
+ (x + 2.075, 1.34),
+ (7, 1.86),
+ arrowprops={"arrowstyle": "->", "color": "#64748b", "lw": 1.4},
+ )
+ _save(fig, "algorithm")
+
+
+def _gpu_sampling() -> None:
+ """Show sample-window ownership and the on-chip calibration pipeline."""
+ fig, ax = plt.subplots(figsize=(16.7, 7.0))
+ fig.subplots_adjust(left=0.025, right=0.985, bottom=0.03, top=0.97)
+ ax.set(xlim=(0, 16), ylim=(0, 7))
+ ax.axis("off")
+ ink, muted = "#17202b", "#52616f"
+ blue, purple = "#386781", "#7557a6"
+ ax.text(
+ 0.1, 6.68, "Spread short sample windows across the current row", fontsize=21, weight="bold"
+ )
+ windows = (0.7, 5.5, 10.3)
+ for j, x in enumerate(windows):
+ ax.add_patch(Rectangle((x, 5.65), 2.4, 0.4, facecolor="#e3efcd", edgecolor="#a7c976"))
+ ax.text(
+ x + 1.2, 5.85, f"Window {j}", ha="center", va="center", fontsize=15, color="#447a00"
+ )
+ if j < 2:
+ ax.plot([x + 2.55, x + 4.65], [5.85, 5.85], color="#b6c0c9", linestyle=(0, (2, 3)))
+ ax.text(13.05, 5.83, "…", fontsize=22, color=muted)
+ ax.text(
+ 0.7,
+ 5.2,
+ "Window starts use a regular stride; no previous-index lookup or random-number generation.",
+ fontsize=15,
+ color=muted,
+ )
+
+ panels = [
+ (0.7, "MAIN · 8 scores / 32 bytes", 8, 0.56),
+ (8.3, "CLUS · 16 scores / 64 bytes", 16, 0.38),
+ ]
+ for x, heading, scores, width in panels:
+ ax.text(x, 4.57, heading, fontsize=19, weight="bold", color=ink)
+ for i in range(scores):
+ fill = "#dcebf3" if i < 8 else "#e9dff4"
+ ax.add_patch(
+ Rectangle(
+ (x + i * width, 3.38),
+ width,
+ 0.56,
+ facecolor=fill,
+ edgecolor="white",
+ linewidth=1,
+ )
+ )
+ for vector in range(scores // 4):
+ left = x + vector * 4 * width
+ ax.add_patch(
+ Rectangle(
+ (left, 3.38), 4 * width, 0.56, fill=False, edgecolor="#8e9ba7", linewidth=1
+ )
+ )
+ ax.text(left + 2 * width, 4.1, "float4 · 16 B", ha="center", fontsize=13, color=muted)
+ for worker in range(scores // 8):
+ start = x + worker * 8 * width
+ color = blue if worker == 0 else purple
+ ax.plot(
+ [start, start, start + 8 * width, start + 8 * width],
+ [3.21, 3.08, 3.08, 3.21],
+ color=color,
+ linewidth=1.2,
+ )
+ label = "Work item j" if worker == 0 else "Work item j + P"
+ ax.text(start + 4 * width, 2.78, label, fontsize=15, ha="center", color=color)
+ ax.text(
+ x,
+ 2.3,
+ "Two vector loads per work item; eight retained sample scores.",
+ fontsize=14,
+ color=muted,
+ )
+
+ stages = [
+ (0.2, "1 REGISTER VALUES", "Warp min/max reductions"),
+ (5.6, "2 SHARED HISTOGRAM", "256 bins · atomic increments"),
+ (11.0, "3 WARP-0 SCAN", "Rank crossings → anchors"),
+ ]
+ for x, heading, description in stages:
+ _box(ax, (x, 0.63), (4.75, 1.05), heading + "\n" + description, "#f1f5f8", 15)
+ if x < 11:
+ ax.annotate(
+ "",
+ (x + 5.25, 1.15),
+ (x + 4.85, 1.15),
+ arrowprops={"arrowstyle": "->", "color": muted, "lw": 1.5},
+ )
+ ax.text(
+ 0.2,
+ 0.11,
+ "P = number of sample windows. CTA threads process work items in strides; a window is not a CUDA block.",
+ fontsize=13,
+ color=muted,
+ )
+ _save(fig, "gpu_sampling")
+
+
+def _integration() -> None:
+ fig, ax = plt.subplots(figsize=(14, 8.4))
+ ax.set(xlim=(0, 14), ylim=(0, 9))
+ ax.axis("off")
+ ax.text(
+ 0.3,
+ 8.65,
+ "One selection core, two row interfaces",
+ fontsize=21,
+ weight="bold",
+ color="#17202b",
+ )
+ _box(
+ ax,
+ (0.75, 7.35),
+ (12.5, 0.82),
+ "Sparse-attention indexer → TopK dispatcher\nOne self-sampling configuration for both phases",
+ "#e9eff5",
+ 12,
+ )
+ adapters = [
+ (
+ 0.75,
+ "DECODE · run_varlen\nKV lengths + MTP offset + compression → valid prefix\n"
+ "Routing: streaming, register, or cluster families",
+ "#edf5df",
+ ),
+ (
+ 7.45,
+ "PREFILL · run_prefill\nCompressed-column window [start, end)\n"
+ "Local indices · one thread block per row",
+ "#e9eff5",
+ ),
+ ]
+ for x, label, color in adapters:
+ _box(ax, (x, 5.1), (5.8, 1.38), label, color, 11.5)
+ ax.annotate(
+ "",
+ (x + 2.9, 6.6),
+ (7, 7.23),
+ arrowprops={"arrowstyle": "->", "color": "#64748b", "lw": 1.5},
+ )
+ ax.annotate(
+ "",
+ (x + 2.9, 4.34),
+ (x + 2.9, 4.98),
+ arrowprops={"arrowstyle": "->", "color": "#64748b", "lw": 1.5},
+ )
+ ax.text(
+ 3.65,
+ 4.62,
+ "streaming route",
+ ha="center",
+ fontsize=9,
+ color="#447a00",
+ backgroundcolor="white",
+ )
+ ax.text(
+ 10.35,
+ 4.62,
+ "compile-time window mode",
+ ha="center",
+ fontsize=9,
+ color="#52616f",
+ backgroundcolor="white",
+ )
+ ax.add_patch(
+ FancyBboxPatch(
+ (0.75, 2.03),
+ 12.5,
+ 2.15,
+ boxstyle="round,pad=0.08,rounding_size=0.08",
+ facecolor="#f5f9ee",
+ edgecolor="#99bb6c",
+ linewidth=1.3,
+ )
+ )
+ ax.text(
+ 7,
+ 3.73,
+ "Shared streaming implementation · GvrMainKernel",
+ ha="center",
+ fontsize=14,
+ weight="bold",
+ color="#447a00",
+ )
+ stages = ["Self-sample", "Multi-threshold\nexact counts", "Collect + refine", "Exact indices"]
+ for i, label in enumerate(stages):
+ x = 1.03 + i * 3.08
+ _box(ax, (x, 2.49), (2.55, 0.7), label, "#ffffff", 11)
+ if i < 3:
+ ax.annotate(
+ "",
+ (x + 2.95, 2.84),
+ (x + 2.68, 2.84),
+ arrowprops={"arrowstyle": "->", "color": "#64748b", "lw": 1.4},
+ )
+ ax.text(
+ 7,
+ 2.14,
+ "Prefill specializes addressing, masks, and index origin; selection logic is shared.",
+ ha="center",
+ fontsize=10,
+ color="#52616f",
+ )
+ _box(
+ ax,
+ (0.75, 0.72),
+ (12.5, 0.66),
+ "Caller-owned INT32 output · no temporal-prior seed, handoff, or write-back",
+ "#edf5df",
+ 12,
+ )
+ ax.annotate(
+ "", (7, 1.5), (7, 1.91), arrowprops={"arrowstyle": "->", "color": "#64748b", "lw": 1.5}
+ )
+ ax.text(
+ 0.75,
+ 0.1,
+ "Runtime support: layout gates · precompiled launchers · exact native fallback",
+ fontsize=10.5,
+ color="#52616f",
+ )
+ fig.subplots_adjust(left=0.015, right=0.985, bottom=0.025, top=0.99)
+ _save(fig, "integration")
+
+
+def _matching(rows: list[dict], model: str) -> list[dict]:
+ required = COMPARISON_ARMS
+ return [
+ r for r in rows if r["model"] == model and all(r[a + "_us"] is not None for a in required)
+ ]
+
+
+def _line_data(rows: list[dict], arm: str, batch: int) -> tuple[list[float], list[float]]:
+ selected = [r for r in rows if r["batch"] == batch]
+ buckets = sorted({r["isl_bucket"] for r in selected}, key=lambda x: int(x[:-1]))
+ groups = [[r for r in selected if r["isl_bucket"] == bucket] for bucket in buckets]
+ return (
+ [median(r["n"] for r in group) for group in groups],
+ [mean(r[arm + "_us"] for r in group) for group in groups],
+ )
+
+
+def _legend(fig: plt.Figure) -> None:
+ handles = [
+ Line2D(
+ [0],
+ [0],
+ color=COLORS[a],
+ linewidth=2.5,
+ linestyle="--" if a in CROSS_CAMPAIGN else "-",
+ label=LABELS[a],
+ )
+ for a in COMPARISON_ARMS
+ ]
+ fig.legend(
+ handles=handles,
+ loc="lower center",
+ ncol=3,
+ frameon=False,
+ bbox_to_anchor=(0.5, 0.01),
+ fontsize=10,
+ )
+
+
+def _latency(rows: list[dict]) -> None:
+ fig, axes = plt.subplots(3, 2, figsize=(13, 11))
+ for i, (model, title) in enumerate(MODELS.items()):
+ matched = _matching(rows, model)
+ for j, batch in enumerate((1, 1024)):
+ ax = axes[i, j]
+ for arm in COMPARISON_ARMS:
+ x, y = _line_data(matched, arm, batch)
+ ax.plot(
+ x,
+ y,
+ marker="o",
+ markersize=3,
+ color=COLORS[arm],
+ linewidth=2.3 if arm == "gvr_v2" else 1.6,
+ linestyle="--" if arm in CROSS_CAMPAIGN else "-",
+ )
+ ax.set_xscale("log", base=2)
+ ax.set_yscale("log")
+ ax.set_xticks(
+ [2**n for n in ((10, 12, 14, 16, 18) if model != "v32" else (12, 14, 16, 18))]
+ )
+ ax.xaxis.set_major_formatter(FuncFormatter(lambda value, _: f"{value / 1024:g}K"))
+ ax.yaxis.set_major_formatter(FuncFormatter(lambda value, _: f"{value:g}"))
+ ax.set_title(f"{title} | B={batch}", fontsize=12, loc="left")
+ ax.set_ylabel("Mean kernel time (µs)")
+ ax.set_xlabel("Valid indexer row length N")
+ ax.grid(which="major", alpha=0.18)
+ fig.suptitle("Latency across row lengths and batch sizes", fontsize=18, weight="bold", y=0.995)
+ fig.subplots_adjust(hspace=0.52, wspace=0.2, bottom=0.14, top=0.94)
+ _legend(fig)
+ _save(fig, "latency")
+
+
+def _roofline_reachable_rates(rows: list[dict]) -> dict:
+ calibration = json.loads((ROOT / "provenance.json").read_text())["roofline_model"]
+ by_model = {}
+ for model in MODELS:
+ matched = _matching(rows, model)
+ k = matched[0]["k"]
+ by_model[model] = {}
+ for arm in COMPARISON_ARMS:
+ widths, times = _line_data(matched, arm, 1024)
+ rates = []
+ for n, us in zip(widths, times):
+ intensity = n / (4 * (n + k))
+ roof = min(
+ calibration["measured_compare_t_s"],
+ calibration["measured_bandwidth_tb_s"] * intensity,
+ )
+ rates.append(100 * (1024 * n / (us * 1e6)) / roof)
+ by_model[model][arm] = {
+ "points": len(rates),
+ "average_percent": mean(rates),
+ "peak_percent": max(rates),
+ }
+ return {
+ "batch": 1024,
+ "reference": "Calibrated roof at each plotted intensity",
+ "aggregation": "Arithmetic mean and maximum of plotted-point reachable rates",
+ "by_model": by_model,
+ }
+
+
+def _roofline(rows: list[dict]) -> None:
+ model_data = json.loads((ROOT / "provenance.json").read_text())["roofline_model"]
+ bw = model_data["measured_bandwidth_tb_s"]
+ compare = model_data["measured_compare_t_s"]
+ fig = plt.figure(figsize=(14, 9.4), facecolor="white")
+ top = fig.add_axes((0.075, 0.61, 0.875, 0.245))
+ intensity = np.geomspace(0.05, 24, 400)
+ top.axvspan(0.125, 0.25, color="#edf5df", zorder=0)
+ top.plot(
+ intensity,
+ np.minimum(
+ model_data["theoretical_bandwidth_tb_s"] * intensity,
+ model_data["theoretical_compare_t_s"],
+ ),
+ color="#94a3b8",
+ linestyle=(0, (5, 3)),
+ linewidth=1.8,
+ )
+ top.plot(intensity, np.minimum(bw * intensity, compare), color="#273746", linewidth=2.4)
+ top.plot([0.125, 0.25], [bw * 0.125, bw * 0.25], color=COLORS["gvr_v2"], linewidth=5)
+ top.annotate(
+ "Ideal Top-K band\n0.125–0.25 compare/byte",
+ xy=(0.18, bw * 0.18),
+ xytext=(0.055, 12),
+ fontsize=10.5,
+ color="#447a00",
+ arrowprops={"arrowstyle": "->", "color": "#579600", "connectionstyle": "arc3,rad=.15"},
+ )
+ top.text(0.85, 2.0, "Bandwidth slope\n6.912 TB/s × intensity", fontsize=10.5, color="#334155")
+ top.annotate(
+ "Compare ceiling\n37.047 Tcompare/s",
+ xy=(12, compare),
+ xytext=(7, 4.5),
+ fontsize=10.5,
+ color="#334155",
+ arrowprops={"arrowstyle": "->", "color": "#64748b"},
+ )
+ top.plot(compare / bw, compare, "o", color="#273746", markersize=5)
+ top.text(5.0, 63, "Knee: 5.36", fontsize=10, ha="center", color="#52616f")
+ top.set(xscale="log", yscale="log", xlim=(0.05, 24), ylim=(0.2, 100))
+ top.set_xticks([0.125, 0.25, 1, 4, 16], ["0.125", "0.25", "1", "4", "16"])
+ top.set_yticks([1, 10, 100], ["1", "10", "100"])
+ top.minorticks_off()
+ top.set_xlabel("Operational intensity · compare/byte", fontsize=10.5, labelpad=8)
+ top.set_ylabel("Tcompare/s", fontsize=10.5)
+ top.grid(axis="y", color="#e9edf1", linewidth=0.7)
+ top.tick_params(length=0, pad=7)
+ for spine in ("left", "bottom"):
+ top.spines[spine].set_color("#d5dce3")
+ top.legend(
+ handles=[
+ Line2D([0], [0], color="#273746", lw=2.4, label="Measured calibration"),
+ Line2D(
+ [0],
+ [0],
+ color="#94a3b8",
+ lw=1.8,
+ linestyle=(0, (5, 3)),
+ label="Theoretical reference",
+ ),
+ ],
+ loc="upper left",
+ bbox_to_anchor=(0.015, 1.21),
+ ncol=2,
+ fontsize=9.5,
+ frameon=False,
+ )
+ for i, (model, title) in enumerate(MODELS.items()):
+ ax = fig.add_axes((0.075 + i * 0.305, 0.205, 0.26, 0.23))
+ matched = _matching(rows, model)
+ k = matched[0]["k"]
+ xroof = np.linspace(0.125, 0.25, 100)
+ ax.fill_between(xroof, xroof * bw, 1.95, color="#f2f5f7", zorder=0)
+ ax.plot(xroof, xroof * bw, color="#273746", linestyle=(0, (2, 2)), linewidth=1.5)
+ for arm in [*TEMPORAL, *ARMS, "gvr_v2"]:
+ widths, times = _line_data(matched, arm, 1024)
+ x = [n / (4 * (n + k)) for n in widths]
+ y = [1024 * n / (us * 1e6) for n, us in zip(widths, times)]
+ ax.plot(
+ x,
+ y,
+ color=COLORS[arm],
+ linewidth=2.8 if arm == "gvr_v2" else 1.5,
+ linestyle="--" if arm in CROSS_CAMPAIGN else "-",
+ marker="o" if arm == "gvr_v2" else ".",
+ markersize=4.5,
+ alpha=1 if arm == "gvr_v2" else 0.8,
+ zorder=5 if arm == "gvr_v2" else 3,
+ )
+ ax.set(xlim=(0.123, 0.253), ylim=(0, 1.95))
+ ax.set_xticks([0.125, 0.175, 0.225, 0.25], [".125", ".175", ".225", ".250"])
+ ax.set_yticks([0, 0.5, 1, 1.5], ["0", "0.5", "1.0", "1.5"])
+ ax.set_xlabel("Intensity (compare/byte)", fontsize=10, labelpad=8)
+ ax.set_title(title, fontsize=11, loc="left", pad=12, weight="bold")
+ if i == 0:
+ ax.set_ylabel("Work throughput (Tcompare/s)", fontsize=10)
+ ax.text(0.133, 1.73, "Calibrated roof", fontsize=9, color="#52616f")
+ ax.grid(axis="y", color="#e9edf1", linewidth=0.7)
+ ax.tick_params(length=0, labelsize=9, pad=7)
+ for spine in ("left", "bottom"):
+ ax.spines[spine].set_color("#d5dce3")
+ fig.text(
+ 0.035,
+ 0.96,
+ "Top-K lives on the bandwidth slope",
+ fontsize=21,
+ weight="bold",
+ color="#17202b",
+ )
+ fig.text(
+ 0.035,
+ 0.92,
+ "A. The full B200 roofline · Top-K intensity stays far below the compute knee",
+ fontsize=11.5,
+ color="#52616f",
+ )
+ fig.text(
+ 0.075,
+ 0.515,
+ "B. Pareto curves across intensities",
+ fontsize=13,
+ weight="bold",
+ color="#17202b",
+ )
+ fig.text(
+ 0.075,
+ 0.482,
+ "B = 1024 · identical layers per model · higher is faster",
+ fontsize=10.5,
+ color="#52616f",
+ )
+ handles = [
+ Line2D(
+ [0],
+ [0],
+ color=COLORS[a],
+ linewidth=2.5,
+ linestyle="--" if a in CROSS_CAMPAIGN else "-",
+ label=LABELS[a],
+ )
+ for a in COMPARISON_ARMS
+ ]
+ fig.legend(
+ handles=handles,
+ loc="lower center",
+ ncol=3,
+ frameon=False,
+ bbox_to_anchor=(0.53, 0.077),
+ fontsize=10,
+ )
+ fig.text(
+ 0.075,
+ 0.046,
+ "Shared ideal work: BN comparisons. Minimum traffic: 4B(N + K) bytes. "
+ "Extra passes and output work remain in measured time.",
+ fontsize=9,
+ color="#52616f",
+ )
+ fig.text(
+ 0.075,
+ 0.018,
+ "Line styles distinguish benchmark runs. Work throughput uses the same logical task for every kernel.",
+ fontsize=9,
+ color="#52616f",
+ )
+ _save(fig, "roofline")
+
+
+def _speedup_map(rows: list[dict], arm: str, label: str, scope: str) -> None:
+ fig, axes = plt.subplots(1, 3, figsize=(14, 5.8))
+ batches = sorted({r["batch"] for r in rows})
+ is_v1 = arm == "temporal_tiered"
+ norm = Normalize(vmin=1, vmax=3 if is_v1 else 21)
+ cmap = plt.get_cmap("YlGnBu")
+ for ax, (model, title) in zip(axes, MODELS.items()):
+ selected = [r for r in rows if r["model"] == model and r[arm + "_us"] is not None]
+ buckets = sorted({r["isl_bucket"] for r in selected}, key=lambda s: int(s[:-1]))
+ values = []
+ widths = []
+ for bucket in buckets:
+ group = [r for r in selected if r["isl_bucket"] == bucket]
+ widths.append(median(r["n"] for r in group) / 1024)
+ values.append(
+ [
+ geometric_mean(
+ r[arm + "_us"] / r["gvr_v2_us"] for r in group if r["batch"] == b
+ )
+ for b in batches
+ ]
+ )
+ data = np.asarray(values)
+ graphic = ax.imshow(data, aspect="auto", cmap=cmap, norm=norm)
+ for y in range(len(buckets)):
+ for x in range(len(batches)):
+ if is_v1 and any(
+ r[arm + "_us"] < r["gvr_v2_us"]
+ for r in selected
+ if r["isl_bucket"] == buckets[y] and r["batch"] == batches[x]
+ ):
+ ax.add_patch(
+ Polygon(
+ [(x + 0.21, y - 0.5), (x + 0.5, y - 0.5), (x + 0.5, y - 0.21)],
+ facecolor="#d46d24",
+ edgecolor="white",
+ linewidth=0.3,
+ )
+ )
+ red, green, blue, _ = cmap(norm(data[y, x]))
+ brightness = 0.299 * red + 0.587 * green + 0.114 * blue
+ ax.text(
+ x,
+ y,
+ f"{data[y, x]:.2f}" if is_v1 else f"{data[y, x]:.1f}",
+ ha="center",
+ va="center",
+ fontsize=6.9 if is_v1 else 7.1,
+ color="white" if brightness < 0.5 else "#17202b",
+ )
+ ax.set_xticks(range(len(batches)), batches, rotation=60, fontsize=9)
+ ax.set_yticks(range(len(widths)), [f"{n:.0f}K" for n in widths], fontsize=9)
+ ax.set_xlabel("Batch size B")
+ ax.set_ylabel("Valid row length N (rounded)")
+ ax.set_title(title, loc="left", fontsize=11, pad=12)
+ fig.suptitle(
+ f"GVR V2 vs {label}: gains across the full length–batch grid",
+ fontsize=18,
+ x=0.035,
+ ha="left",
+ weight="bold",
+ y=1.02,
+ )
+ fig.subplots_adjust(left=0.06, right=0.99, top=0.87, bottom=0.34, wspace=0.27)
+ cax = fig.add_axes((0.34, 0.055 if is_v1 else 0.09, 0.32, 0.026))
+ ticks = [1, 1.5, 2, 2.5, 3] if is_v1 else [1, 5, 10, 15, 21]
+ bar = fig.colorbar(graphic, cax=cax, orientation="horizontal", ticks=ticks)
+ if is_v1:
+ bar.ax.set_xticklabels(["1× · parity", "1.5×", "2×", "2.5×", "3×"])
+ bar.set_label(f"{label} time / GVR V2 time · geometric mean across layers", fontsize=9)
+ fig.text(
+ 0.06,
+ 0.19 if is_v1 else 0.17,
+ f"{scope} · 1.0× is parity · shared color scale across all three models.",
+ fontsize=9,
+ )
+ if is_v1:
+ fig.add_artist(
+ Polygon(
+ [(0.06, 0.155), (0.07, 0.155), (0.07, 0.133)],
+ transform=fig.transFigure,
+ facecolor="#d46d24",
+ edgecolor="none",
+ )
+ )
+ fig.text(
+ 0.08,
+ 0.138,
+ "Orange corner: at least one layer is slower in V2. Layer averages do not show every case.",
+ fontsize=9,
+ color="#52616f",
+ )
+ _save(fig, "gvr_v1_map" if is_v1 else arm + "_map")
+
+
+def main() -> None:
+ """Validate the frozen dataset, then regenerate statistics and ten figures."""
+ plt.rcParams.update(
+ {
+ "font.family": "DejaVu Sans",
+ "font.size": 11,
+ "svg.fonttype": "none",
+ "svg.hashsalt": "gvr-v2-blog",
+ "axes.spines.top": False,
+ "axes.spines.right": False,
+ }
+ )
+ rows = _load()
+ if any(row[arm + "_us"] is None for row in rows for arm in COMPARISON_ARMS):
+ raise ValueError(
+ "Figure regeneration requires complete paired timings for all comparison arms"
+ )
+ summary = {
+ "copyright": COPYRIGHT,
+ "reference": json.loads((ROOT / "provenance.json").read_text())["reference"],
+ "overall": {a: _stats(rows, a) for a in ARMS},
+ "comparison_common_cases": _comparison(rows),
+ "by_model": {
+ m: {a: _stats([r for r in rows if r["model"] == m], a) for a in ARMS} for m in MODELS
+ },
+ "temporal_vs_v2": {a: _stats(rows, a) for a in TEMPORAL},
+ "roofline_reachable_rate": _roofline_reachable_rates(rows),
+ "evolution_vs_radix": {
+ a: {
+ "geomean": geometric_mean(r["radix_cuda_us"] / r[a + "_us"] for r in rows),
+ "wins_percent": 100 * mean(r["radix_cuda_us"] > r[a + "_us"] for r in rows),
+ }
+ for a in [*TEMPORAL, "gvr_v2"]
+ },
+ }
+ (ROOT / "summary.json").write_text(json.dumps(summary, indent=2) + "\n")
+ for name, result in summary["overall"].items():
+ print(
+ name, f"{result['geomean']:.6f}×", result["cases"], f"wins {result['win_percent']:.3f}%"
+ )
+ _overview(rows)
+ _evolution(rows)
+ _candidate_work()
+ _algorithm()
+ _gpu_sampling()
+ _speedup_map(rows, "radix_cuda", "radix CUDA", "TensorRT LLM production dispatcher")
+ _speedup_map(rows, "temporal_tiered", "GVR V1", "GVR V1 (temporal hint)")
+ _latency(rows)
+ _roofline(rows)
+ _integration()
+
+
+if __name__ == "__main__":
+ main()
diff --git a/docs/source/blogs/media/gvr_v2/pro_timings.csv.gz b/docs/source/blogs/media/gvr_v2/pro_timings.csv.gz
new file mode 100644
index 000000000000..7daaa1773cf1
Binary files /dev/null and b/docs/source/blogs/media/gvr_v2/pro_timings.csv.gz differ
diff --git a/docs/source/blogs/media/gvr_v2/provenance.json b/docs/source/blogs/media/gvr_v2/provenance.json
new file mode 100644
index 000000000000..6317dd6488e1
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/provenance.json
@@ -0,0 +1,46 @@
+{
+ "copyright": "Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. SPDX-License-Identifier: Apache-2.0",
+ "reference": {
+ "repository": "NVIDIA/TensorRT-LLM",
+ "implementation": "GVR V2 self-sampling",
+ "entry": "run_varlen",
+ "mode": "hint-free",
+ "pull_request": "https://github.com/NVIDIA/TensorRT-LLM/pull/19076",
+ "commit": "be1b9885e8df9bf070e8cb68459e24a7119afaa9"
+ },
+ "hardware": "NVIDIA B200 (SM100)",
+ "baseline_versions": {
+ "radix_cuda": "Production insertion/radix/split-work dispatcher",
+ "temporal_tiered": "GVR V1 (temporal hint), tiered implementation from public PR #16877"
+ },
+ "pairing": {
+ "radix_cuda": "cross-campaign",
+ "temporal_tiered": "cross-campaign, same (cell, batch), not an isolated algorithm ablation"
+ },
+ "published_files": {
+ "flash_timings.csv.gz": "0f1c4799296bea1db9a012fe9ddf90521918af3ce73e6b39ad25a6c05115dcc6",
+ "pro_timings.csv.gz": "bfd713272ccd50cab57d8f436d1d788e29d3ef6cc2014818104b2872f4476cbd",
+ "v32_timings.csv.gz": "59bccd05ba685d01ed1ebb34c9a0b15df02d1879f8c16bbaffa2a264f512fd8d",
+ "temporal_comparison.csv.gz": "2abdeb8eafa1366be5673c79dfdb938a02716431cf8745f18d3a2fc240a778ca"
+ },
+ "scope": "Per-case FP32 cold kernel means with PR #19076 as the GVR V2 reference. GVR V1 (temporal hint, PR #16877) and TensorRT LLM radix CUDA observations are paired across benchmark runs by workload and batch. All three implementations cover the same 9,746 cases.",
+ "roofline_model": {
+ "theoretical_bandwidth_tb_s": 8.0,
+ "measured_bandwidth_tb_s": 6.912116,
+ "theoretical_compare_t_s": 37.22496,
+ "measured_compare_t_s": 37.04749,
+ "work_definition": "One abstract comparison per input score; W=B*N, Qmin=4*B*(N+K).",
+ "compare_calibration": "Study counts two semantic binary comparisons per FMNMX3 result.",
+ "scope": "Single-GPU read-dominated ceiling; actual Top-K instruction count and extra traffic are excluded."
+ },
+ "illustrative_figures": {
+ "temporal_overlap.svg": {
+ "sha256": "4b9081122fb8f4fa208f4e3c66acc76bfd0ee131fab14cd9ef424db5b28be63c",
+ "source": "Temporal Top-K overlap figure supplied by the article authors",
+ "conversion": "Standalone figure converted from PDF to SVG; plotted values, labels, vector axes, and embedded point clouds preserved.",
+ "scope": "SWE-bench-64K traces for DeepSeek-V3.2 and DeepSeek-V4 Pro; temporal-hint hit rates after index mapping (+1 for V3.2, unchanged compressed-bin indices for V4 Pro), distinct from kernel timing comparisons.",
+ "reproduction": "Imported figure asset; not generated from the bundled kernel-timing CSVs.",
+ "display": "Intrinsic width is 1200 CSS pixels for responsive full-column display; the original viewBox, aspect ratio, and plotted content are preserved."
+ }
+ }
+}
diff --git a/docs/source/blogs/media/gvr_v2/radix_cuda_map.svg b/docs/source/blogs/media/gvr_v2/radix_cuda_map.svg
new file mode 100644
index 000000000000..b5ffeccfcfa9
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/radix_cuda_map.svg
@@ -0,0 +1,1655 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/roofline.svg b/docs/source/blogs/media/gvr_v2/roofline.svg
new file mode 100644
index 000000000000..dcd12b9ce644
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/roofline.svg
@@ -0,0 +1,1761 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/speedup.svg b/docs/source/blogs/media/gvr_v2/speedup.svg
new file mode 100644
index 000000000000..47b16b23ed45
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/speedup.svg
@@ -0,0 +1,502 @@
+
+
+
diff --git a/docs/source/blogs/media/gvr_v2/summary.json b/docs/source/blogs/media/gvr_v2/summary.json
new file mode 100644
index 000000000000..eedc74aa7a54
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/summary.json
@@ -0,0 +1,175 @@
+{
+ "copyright": "Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. SPDX-License-Identifier: Apache-2.0",
+ "reference": {
+ "repository": "NVIDIA/TensorRT-LLM",
+ "implementation": "GVR V2 self-sampling",
+ "entry": "run_varlen",
+ "mode": "hint-free",
+ "pull_request": "https://github.com/NVIDIA/TensorRT-LLM/pull/19076",
+ "commit": "be1b9885e8df9bf070e8cb68459e24a7119afaa9"
+ },
+ "overall": {
+ "radix_cuda": {
+ "cases": 9746,
+ "geomean": 5.051863720455904,
+ "minimum": 1.3352283494409702,
+ "p5": 1.800987459271426,
+ "p95": 10.216837270210021,
+ "wins": 9746,
+ "win_percent": 100,
+ "baseline_median_us": 46.3535,
+ "gvr_median_us": 9.6435
+ }
+ },
+ "comparison_common_cases": {
+ "flash": {
+ "cases": 2079,
+ "layers": 21,
+ "latency_relative_to_v2": {
+ "gvr_v2": 1.0,
+ "temporal_tiered": 1.5270656764856136,
+ "radix_cuda": 4.875738163985888
+ }
+ },
+ "pro": {
+ "cases": 2970,
+ "layers": 30,
+ "latency_relative_to_v2": {
+ "gvr_v2": 1.0,
+ "temporal_tiered": 1.54320332519381,
+ "radix_cuda": 4.871976103268161
+ }
+ },
+ "v32": {
+ "cases": 4697,
+ "layers": 61,
+ "latency_relative_to_v2": {
+ "gvr_v2": 1.0,
+ "temporal_tiered": 1.3723454906900026,
+ "radix_cuda": 5.250852136354158
+ }
+ }
+ },
+ "by_model": {
+ "flash": {
+ "radix_cuda": {
+ "cases": 2079,
+ "geomean": 4.875738163985888,
+ "minimum": 1.4130089899524063,
+ "p5": 1.6569445902710334,
+ "p95": 10.993023430337592,
+ "wins": 2079,
+ "win_percent": 100,
+ "baseline_median_us": 39.606,
+ "gvr_median_us": 6.608
+ }
+ },
+ "pro": {
+ "radix_cuda": {
+ "cases": 2970,
+ "geomean": 4.871976103268161,
+ "minimum": 1.3352283494409702,
+ "p5": 1.6883335124343106,
+ "p95": 10.58141910734875,
+ "wins": 2970,
+ "win_percent": 100,
+ "baseline_median_us": 42.3375,
+ "gvr_median_us": 6.979
+ }
+ },
+ "v32": {
+ "radix_cuda": {
+ "cases": 4697,
+ "geomean": 5.250852136354158,
+ "minimum": 2.1828381113051605,
+ "p5": 3.648166973469981,
+ "p95": 8.976731655272205,
+ "wins": 4697,
+ "win_percent": 100,
+ "baseline_median_us": 49.974,
+ "gvr_median_us": 10.768
+ }
+ }
+ },
+ "temporal_vs_v2": {
+ "temporal_tiered": {
+ "cases": 9746,
+ "geomean": 1.4550894088383663,
+ "minimum": 0.6885347106288259,
+ "p5": 1.095922446440455,
+ "p95": 2.102507274523004,
+ "wins": 9704,
+ "win_percent": 99.56905397085984,
+ "baseline_median_us": 12.232,
+ "gvr_median_us": 9.6435
+ }
+ },
+ "roofline_reachable_rate": {
+ "batch": 1024,
+ "reference": "Calibrated roof at each plotted intensity",
+ "aggregation": "Arithmetic mean and maximum of plotted-point reachable rates",
+ "by_model": {
+ "flash": {
+ "gvr_v2": {
+ "points": 9,
+ "average_percent": 41.56601571959186,
+ "peak_percent": 77.81571160254971
+ },
+ "temporal_tiered": {
+ "points": 9,
+ "average_percent": 26.123944119758303,
+ "peak_percent": 63.276309173607764
+ },
+ "radix_cuda": {
+ "points": 9,
+ "average_percent": 7.73037106360671,
+ "peak_percent": 16.904373664943286
+ }
+ },
+ "pro": {
+ "gvr_v2": {
+ "points": 9,
+ "average_percent": 39.00405637364329,
+ "peak_percent": 68.39789370398464
+ },
+ "temporal_tiered": {
+ "points": 9,
+ "average_percent": 25.39201270408693,
+ "peak_percent": 58.92620571524122
+ },
+ "radix_cuda": {
+ "points": 9,
+ "average_percent": 7.832445928048165,
+ "peak_percent": 16.38668927402192
+ }
+ },
+ "v32": {
+ "gvr_v2": {
+ "points": 7,
+ "average_percent": 41.47747500688756,
+ "peak_percent": 66.53750601976955
+ },
+ "temporal_tiered": {
+ "points": 7,
+ "average_percent": 27.779057298223755,
+ "peak_percent": 51.13105668384031
+ },
+ "radix_cuda": {
+ "points": 7,
+ "average_percent": 8.28762774707847,
+ "peak_percent": 17.75730356940984
+ }
+ }
+ }
+ },
+ "evolution_vs_radix": {
+ "temporal_tiered": {
+ "geomean": 3.4718579420414657,
+ "wins_percent": 98.38908270059513
+ },
+ "gvr_v2": {
+ "geomean": 5.051863720455904,
+ "wins_percent": 100
+ }
+ }
+}
diff --git a/docs/source/blogs/media/gvr_v2/temporal_comparison.csv.gz b/docs/source/blogs/media/gvr_v2/temporal_comparison.csv.gz
new file mode 100644
index 000000000000..1f415da49af9
Binary files /dev/null and b/docs/source/blogs/media/gvr_v2/temporal_comparison.csv.gz differ
diff --git a/docs/source/blogs/media/gvr_v2/temporal_overlap.svg b/docs/source/blogs/media/gvr_v2/temporal_overlap.svg
new file mode 100644
index 000000000000..ce81feeafe3e
--- /dev/null
+++ b/docs/source/blogs/media/gvr_v2/temporal_overlap.svg
@@ -0,0 +1,5321 @@
+
diff --git a/docs/source/blogs/media/gvr_v2/v32_timings.csv.gz b/docs/source/blogs/media/gvr_v2/v32_timings.csv.gz
new file mode 100644
index 000000000000..b5ef92ff0946
Binary files /dev/null and b/docs/source/blogs/media/gvr_v2/v32_timings.csv.gz differ
diff --git a/docs/source/blogs/tech_blog/blog29_GVR_V2_Self_Sampling_Exact_TopK_for_Sparse_Attention.md b/docs/source/blogs/tech_blog/blog29_GVR_V2_Self_Sampling_Exact_TopK_for_Sparse_Attention.md
new file mode 100644
index 000000000000..ebb422c4ea35
--- /dev/null
+++ b/docs/source/blogs/tech_blog/blog29_GVR_V2_Self_Sampling_Exact_TopK_for_Sparse_Attention.md
@@ -0,0 +1,533 @@
+
+
+# GVR V2: Self-Sampling and Multi-Thresholding for Faster Exact Top-K
+
+*A Unified Selection Core for Prefill and Decode in TensorRT LLM*
+
+By NVIDIA TensorRT LLM Team
+
+Selecting 1,024 INT32 indices from 131,072 FP32 scores writes just **4 KiB of output**, yet one complete read of the scores moves **512 KiB**. Every additional full-row pass pays that input cost again. For a sparse-attention indexer, finding the Top-K boundary can therefore cost far more than emitting the winners.
+
+GVR V2 makes each full-row pass more useful without depending on the previous decode step to predict the current one. **Self-sampling estimates where the current row's Top-K boundary lies; multi-thresholding derives many exact population counts from one classification pass.** Together, they concentrate exact refinement on the small group of scores still competing for the final slots. Removing the Top-K prior also lets prefill and decode share a streaming selection core, with phase differences handled by row adapters.
+
+On B200, this design delivers **5.05× geometric-mean speedup over TensorRT LLM radix CUDA**, with **1.46× over GVR V1**. Both comparisons use the same 9,746 workloads spanning DeepSeek-V3.2, DeepSeek-V4 Flash, and DeepSeek-V4 Pro indexers.
+
+
+
+*Figure 1. Kernel time relative to GVR V2, geometrically averaged over the same workloads within each model; shorter is faster. GVR V1 uses a temporal hint. All three implementations cover the full 9,746-case grid.*
+
+**The operator contract.** Given FP32 indexer scores and valid-row metadata, Top-K returns unordered INT32 positions for sparse attention's KV selection. With finite scores and at least $K$ entries, it selects an exact value multiset through $K$ distinct indices; ties can choose different positions. [Enablement](#enable-gvr-v2) lists hardware, shape, and configuration requirements.
+
+[The original GVR blog](https://nvidia.github.io/TensorRT-LLM/blogs/tech_blog/blog21_Temporal_Correlation_Meets_Sparse_Attention.html) described a temporal shortcut: the previous decode step's selected indices predict the next step's winners. Experience with V1 exposed two limits: hint quality varies sharply, and maintaining the hint couples selection to the serving framework. V2 makes the current row the source of the guess, targeting **a stronger performance floor and better average latency**, while enabling **one selection core for prefill and decode**. The **Guess–Verify–Refine** exactness contract remains.
+
+**Table of Contents**
+
+- **[Motivation and Design Foundations](#motivation-and-design-foundations)**
+ - [From GVR V1 to V2: Why Move Beyond Temporal Hints?](#from-gvr-v1-to-v2-why-move-beyond-temporal-hints)
+ - [From Floyd–Rivest SELECT to GPU Top-K](#from-floydrivest-select-to-gpu-top-k)
+- **[Self-Sampling and Multi-Thresholding](#self-sampling-and-multi-thresholding)**
+ - [Why Threshold Quality Matters: Passes and Candidate Work](#why-threshold-quality-matters-passes-and-candidate-work)
+ - [Self-Sampling: Calibrate the Search to This Row](#self-sampling-calibrate-the-search-to-this-row)
+ - [Multi-Thresholding: Make Each Full-Row Pass Count](#multi-thresholding-make-each-full-row-pass-count)
+ - [Mapping Selection to Blackwell](#mapping-selection-to-blackwell)
+- **[Performance and Roofline Analysis](#performance-and-roofline-analysis)**
+ - [Performance Against GVR V1 and Radix CUDA](#performance-against-gvr-v1-and-radix-cuda)
+ - [The Roofline Model: Fewer Passes, More Useful Work](#the-roofline-model-fewer-passes-more-useful-work)
+- **[TensorRT LLM Integration and Takeaways](#tensorrt-llm-integration-and-takeaways)**
+ - [Decode and Prefill in TensorRT LLM](#decode-and-prefill-in-tensorrt-llm)
+ - [Conclusion](#conclusion)
+ - [Further Reading](#further-reading)
+
+## Motivation and Design Foundations
+
+### From GVR V1 to V2: Why Move Beyond Temporal Hints?
+
+GVR V1 gathers current scores at the previous step's Top-K indices to predict admission thresholds. **The GVR V1 baseline already uses multi-thresholding.** Its [streaming implementation](https://github.com/NVIDIA/TensorRT-LLM/pull/16877) uses sampled ladder counts to choose a pivot and a rescue rung, then verifies both exactly in a fused count/collect pass. This is the temporal-hint implementation compared with V2 below.
+
+Its admission objective is expressed through the monotone count function
+
+$$
+C(T)=\sum_{i=0}^{N-1}\mathbf{1}[x_i\ge T].
+$$
+
+An admission threshold should leave enough survivors to contain Top-K, but few enough to fit candidate capacity. A high-quality temporal hint can make admission very cheap. The difficulty is making that benefit reliable across inference workloads, even with several thresholds checked together.
+
+#### A Biased Sample with Variable Value
+
+Temporal hints are a **biased sample** of current scores at positions predicted from the previous step's winners. The hit rate is the fraction of the current Top-K covered by the mapped temporal hint. High, stable overlap makes that bias useful. Figure 2 shows why it is an unreliable assumption across layers and decode steps.
+
+
+
+*Figure 2. Temporal overlap on SWE-bench-64K workloads. Blue marks current selections matched after applying the temporal-hint index mapping; orange marks selections not predicted by that hint. V3.2 shifts prior indices by +1, while V4 Pro keeps the same compressed-bin indices. Upper panels show position crops; lower curves measure full-domain overlap across layers and steps, with means in parentheses. Even a high-mean layer can suffer an abrupt collapse.*
+
+The V3.2 +1 shift is a temporal prediction rule. Its overlap measures how well the shifted positions predict the current Top-K, rather than retention at identical token indices.
+
+Near-64K measurements also expose dependence on the input and layer:
+
+
+
+*Distribution of per-layer mean hit rates, with layer IDs matched between inputs within each model. Each layer is averaged over its decode steps and then weighted equally; P10–P90 and min–max describe those layer means, not individual transitions. V3.2 uses +1 alignment.*
+
+V4 Pro's mean changes from **71.5% to 57.9%** between these inputs, and its random-token P10–P90 spans **33.6–80.4%**. V3.2 has similar overall means across inputs, yet its weakest SWE-bench layer averages only **5.7%**. **An average overlap cannot serve as a dependable per-row performance assumption.** The table exposes variation across inputs and layers; Figure 2 adds the abrupt changes within a layer over time.
+
+A row's true hit rate is known only after the current selection is established. Verification can expose a poor threshold, but the hint gather and initial work have already been paid for. Conservative admission, repeated counts, capacity checks, and exact recovery keep weak hints safe; their overhead and extra reads reduce average speedup and make latency less predictable.
+
+V2 calibrates from the **current row**, removing dependence on temporal overlap and the read through old indices. It combines this calibration with histogram-based multi-threshold verification and crossing-bin refinement. **The defining change from GVR V1 is the source of the guess and the removal of temporal state; multi-thresholding is part of the design continuity.** These choices target a stronger practical performance floor and better average latency, without a fixed worst-case latency guarantee.
+
+#### A Hint That Crosses Framework Boundaries
+
+V1's prior also has a lifecycle outside the kernel. GVR V1 has no prefill engine: TensorRT LLM uses radix for prefill and can seed the decode prior from each request's last prefill selection. That phase-dependent history complicates a common selection architecture. V2's current-row calibration removes the Top-K prior dependency, allowing phase differences to stay in dispatch and row-interface adapters. The [integration section](#decode-and-prefill-in-tensorrt-llm) traces the consequences for CUDA Graph preparation and disaggregated serving.
+
+
+
+| Algorithm question | GVR V1 (temporal hint) | Streaming GVR V2 |
+| :---: | :---: | :---: |
+| Where does the guess come from? | Current scores gathered through previous-step indices | Packed sample windows spread across the current row |
+| What makes the guess useful? | High, stable overlap with previous winners | Coverage of the current row's score distribution |
+| What guides admission? | A hint-derived pivot and rescue rung | Sample-derived primary threshold, lower safety floor, and upper anchor |
+| What does verification learn? | Exact counts at multiple admission thresholds | Exact bin populations and counts at many boundaries |
+| Where does exact refinement start? | The admitted candidate set, with path-specific local refinement | The crossing bin containing rank $K$ |
+| What state crosses decode steps? | Per-layer prior indices | No Top-K prior in the self-sampling path |
+| How do prefill and decode relate? | Radix prefill; its last selection can seed temporal decode | Shared streaming selection with phase-specific row adapters |
+| How does a bad guess affect the result? | More admission/refinement work or recovery; membership remains exact | Lower admission or exact recovery; membership remains exact |
+
+
+
+
+
+*Figure 3. Multi-thresholding is shared by GVR V1 and V2. The flows emphasize their calibration and refinement choices on streaming paths; the bars compare complete GVR V1 and V2 implementations over radix CUDA. V2 removes the temporal-overlap dependency and prior-state lifecycle.*
+
+GVR V1 achieves **3.47×** speedup over radix CUDA; V2 reaches **5.05×**. V2 is **1.46× faster than GVR V1**. These gains compare complete implementations, including their calibration, verification, refinement, and execution paths.
+
+The fifth percentile remains above parity, while the minimum exposes workloads where V1 retains an advantage:
+
+
+
+P5 is the fifth percentile across workload-level speedups, each computed as V1 time divided by V2 time. The **0.689× minimum means V2 takes about 45% longer than V1** in the worst measured case. The 99.57% win rate supports broad improvement, while this minimum makes the remaining regressions explicit.
+
+**V1's calibration with temporal hints can still produce a better admission threshold.** An informative prior, validated against current-row samples, can let V1 admit a tighter candidate set than self-sampling. Diagnostics of the pronounced large-batch regressions show V1 accepting enough candidates on its first pass, while V2 initially admits fewer than $K$ scores and rescans at a lower threshold. This identifies extra admission work; differences in dispatch and refinement also contribute to total kernel cost.
+
+These local wins fit the motivation for moving beyond temporal hints. Their usefulness varies across inputs, layers, and decode steps, and their true overlap is unavailable before selection. V2 removes that unstable dependency to improve robustness and average latency with a shared current-row selection core. Self-sampling can still misestimate a tail, so its admission margin and recovery cost remain optimization targets. **A stronger practical performance floor is a design objective, not a guarantee that V2 beats V1 on every input.** This workload distribution also does not establish runtime P95/P99 latency or a fixed worst-case bound.
+
+### From Floyd–Rivest SELECT to GPU Top-K
+
+GVR V2's self-sampling was inspired by Floyd and Rivest's 1975 paper, [*Expected Time Bounds for Selection*](https://people.csail.mit.edu/rivest/pubs/FR75a.pdf). Its theoretical SELECT algorithm draws a random sample, chooses two sample order statistics to bracket the desired rank, partitions the full input, and continues exactly in the partition containing that rank. If the bracket misses, selection continues in the appropriate outer partition. Sampling reduces expected work without making the answer approximate.
+
+The classical expected comparison bound for ascending rank $i$ among $n$ elements is
+
+$$
+n+\min(i,n-i)+o(n).
+$$
+
+This belongs to a comparison model with random-sampling assumptions; the original treatment assumes distinct keys. [Kiwiel's later analysis](https://arxiv.org/abs/cs/0312055) establishes rigorous bounds for SELECT variants, including repeated keys. These results motivate **using sample ranks to narrow an exact selection problem**.
+
+GVR V2 carries that principle into a different cost model:
+
+
+
+| Design choice | Floyd–Rivest theoretical SELECT | GVR V2 streaming |
+| :---: | :---: | :---: |
+| Primary objective | Expected element comparisons | Kernel latency: input passes, memory traffic, and parallel work |
+| Calibration | Random sample and sample order statistics | Regularly spaced, packed sample windows and histogram quantiles |
+| Remaining selection | Exact partitioning and recursive selection | Exact bin counts, crossing-bin refinement, and recovery |
+| Result | An element at the requested rank | An exact set of $K$ indices, without requiring sorted output |
+
+
+
+On a GPU, doing more work on chip can be worthwhile if it avoids another full-row read. V2 couples sampling to candidate capacity, vectorized loads, and multi-threshold counts. **The inherited idea is sample-guided exact selection; the optimization target is GPU execution cost.** Its deterministic sampling policy does not inherit SELECT's randomized comparison bound. Exactness follows from full-row accounting and exact refinement or recovery.
+
+## Self-Sampling and Multi-Thresholding
+
+For nontrivial selection, the sampled streaming paths follow three steps. Register-resident families bypass sparse sampling and use a different initial bracket, while retaining exact classification and refinement; the dispatcher chooses among these execution families for each workload.
+
+1. **Guess:** sample packed windows of the current row to place admission and verification bins near its tail.
+2. **Verify:** account for every valid score and establish a complete admitted population containing at least K candidates. Too few survivors require wider admission; incomplete staging requires exact recovery.
+3. **Refine:** emit winners above the crossing bin and select the remaining slots exactly from that bin.
+
+
+
+*Figure 4. The sample histogram estimates a bracket; the verification histogram counts the complete admitted population. The lower panel shows how verification controls refinement and recovery. Bin heights and the 980/73/44 example are illustrative; exact recovery depends on the kernel family.*
+
+For example, with $K=1024$, suppose 980 scores lie above the crossing bin and 73 lie inside it. Emit the 980 directly and select the best **44 of those 73**. A useful sample keeps this boundary problem small; complete counts and exact refinement make the answer correct.
+
+### Why Threshold Quality Matters: Passes and Candidate Work
+
+A useful threshold balances **full-row passes and candidate work**. A loose threshold may save a scan yet admit so many candidates that processing them consumes the saving. Another pass is worthwhile only when the work it removes exceeds its cost.
+
+For a finite-score row, let $\tau$ be the exact K-th-largest score and $q\le\tau$ an admission threshold. The candidate population decomposes as
+
+$$
+C_p=C(q)=K+E+D(q,\tau),\qquad
+\frac{C_p}{K}=1+\frac{E+D(q,\tau)}{K}.
+$$
+
+Here $E$ counts excess entries tied at the boundary, and $D$ counts scores in the shell $q\le x_i\lt\tau$. Dense scores near the boundary can turn a small threshold error into large **candidate amplification**. Temporal overlap alone therefore cannot predict refinement work.
+
+
+
+*Figure 5.* A: exact tail counts identify thresholds satisfying $K\le C(q)\le B_r$, where $B_r$ is candidate capacity. The labeled segments split the admitted population into required winners, excess boundary ties, and the boundary shell. B: candidate amplification arises from the latter two components; their proportions are schematic. Exact handling of ties preserves the Top-K value multiset.
+
+Self-sampling aims to place admission near the current tail; multi-thresholding separates certain winners, the crossing bin, and lower bins. V2 balances full-row reads against candidate handling over $C(q)$ entries and exact refinement over $m$ crossing-bin candidates.
+
+### Self-Sampling: Calibrate the Search to This Row
+
+V2 calibrates inside the selection kernel, using the current row's layout to generate sample addresses. It needs neither previous winners nor a separate sampling kernel. The sample estimates a useful starting region; full-row verification determines whether that region contains enough candidates.
+
+#### Load Short Windows, Spread Them Across the Row
+
+A **sample window** is a contiguous group of scores. A CUDA thread block processes many such windows. The streaming families use these layouts:
+
+
+
+| Family | Scores per window | Loads and logical ownership |
+| :---: | :---: | :---: |
+| `main` | 8 FP32 scores = 32 bytes | One work item loads two adjacent `float4` vectors |
+| `clus` | 16 FP32 scores = 64 bytes | Two work items each load two vectors, covering the lower and upper halves |
+
+
+
+Regular spacing spreads the windows across the valid row. The sampling budget sets their count and spacing, while bounds and alignment handling keep loads valid for each phase. A thread block processes more windows in iterations when the sample exceeds its parallel capacity.
+
+**Why eight?** Eight FP32 scores fill a 32-byte memory sector when the window is sector-aligned. A scattered scalar sample can use only four bytes from each fetched sector. Two adjacent, 16-byte-aligned `float4` loads consume the whole window, giving more sample values per touched sector. This is local packing within each window; widely spaced windows still produce strided accesses across the warp. The [CUDA memory-access model](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#coalesced-access-to-global-memory) explains the sector granularity. Actual sector traffic also depends on the row's alignment and cache state.
+
+**Why sixteen in `clus`?** It groups two neighboring 32-byte pieces into one sampled location, divided between two work items. Each retains the same eight-score payload as `main`, avoiding a sixteen-score live payload in one worker. The engineering tradeoff is locality versus coverage: at a fixed sample budget, sixteen-score windows visit half as many locations as eight-score windows. Both aligned layouts can fully use their sectors; the clustered grouping favors wider local coverage at each visited location without doubling the per-worker sample payload. These are implementation tradeoffs, not a statistical guarantee or a universal optimum; a 64-byte window comprises four vector loads.
+
+
+
+*Figure 6. Each outlined memory box holds four FP32 scores. The 8/16-score grouping describes data layout, while CTA threads execute the work items. Register values feed a shared sample histogram; its rank crossings produce calibration anchors. The lower floor is computed where enabled.*
+
+#### Reduce, Histogram, and Extract the Anchors On Chip
+
+Each sampling worker keeps its first two vectors in registers. Local minima and maxima reduce within each warp; warp results are published through shared memory and combined after a CTA barrier. These extrema define **256 sample bins**. Workers then increment shared-memory counters for their sample values. Additional windows beyond the initially retained pair are reloaded for histogramming, limiting live register state.
+
+After another barrier, warp 0 scans the histogram using vector counter loads and warp shuffles. One scan finds all required rank crossings and clears the counters for the next phase, avoiding a sample sort or a separate search for every anchor.
+
+Let $S$ be the actual sample size and $A\ge K$ the desired full-row candidate population. The launch policy starts with a sample budget proportional to $N/A$, clamps its target between 256 scores and half the row, and rounds to legal windows. It computes target ranks from the resulting $S$, rather than the unrounded budget:
+
+$$
+r_A\approx\frac{AS}{N},\qquad r_K\approx\frac{KS}{N},\qquad r_{2A}\approx\frac{2AS}{N}.
+$$
+
+The intuition is to keep enough observations in the relevant tail: as its target fraction $A/N$ shrinks, increasing $S$ maintains roughly $AS/N$ tail observations, before clamping and window rounding. Descending histogram crossings supply:
+
+- **Primary threshold $T$:** the lower edge of the sample bin at rank $r_A$, aiming for roughly $A$ full-row survivors.
+- **Upper anchor $T_K$:** the bin edge near rank $r_K$, used to choose the verification histogram's upper bound with headroom for sampling error. Values beyond that bound remain candidates.
+- **Lower floor $T_{\mathrm{floor}}$:** the bin edge near rank $r_{2A}$, where enabled, admitting a larger population after an aggressive estimate.
+
+The sample histogram estimates these anchors without assuming a known score distribution. Regularly spaced windows can still miss an unusual tail or correlate with structured scores. Degenerate samples use conservative recovery paths; exact full-row accounting remains the authority.
+
+#### Overlap Calibration with the Upcoming Row Read
+
+In the variable-length `main` route, warp 0 prepares sampling geometry while other warps can prefetch their upcoming row slice. The schedule limits overlapping register lifetimes. Clustered streaming repeats the same sample in each CTA so that all ranks derive a common bracket before merging verification histograms. These choices trade a little repeated calibration work for simpler communication. Exact address formulas, prefetch placement, and anchor constants are retained in the [implementation companion](../media/gvr_v2/README.md#gpu-sampling-implementation).
+
+### Multi-Thresholding: Make Each Full-Row Pass Count
+
+A scalar verification pass answers one question: how many scores exceed $T$? GVR V1 already amortizes row reads across several admission thresholds. V2's streaming path uses a verification histogram to obtain a dense family of counts from the same classification work and locate the boundary for exact refinement.
+
+Conceptually, divide the bracket $[T,H]$ into $M$ ordered bins, with boundaries $t_0,\ldots,t_M$, and let $h_j$ be the exact population of bin $j$. A descending cumulative scan yields
+
+$$
+C(t_j)=\sum_{\ell=j}^{M-1}h_\ell,\qquad 0\le j\lt M.
+$$
+
+The streaming verification histogram has 256 bins. Each survivor is assigned to a bin once; a small on-chip cumulative scan then exposes counts at all its boundaries. **One classification contributes to an entire family of threshold counts.** The expensive score reads are shared, and the remaining scan touches only the bin counters. Verification can locate where the population crosses $K$ without issuing another full-row query for each trial threshold.
+
+Values above $H$ saturate into the top bin and remain candidates. Full-row accounting establishes whether enough survivors exist. The implementation may build the histogram during streaming, merge shard histograms through distributed shared memory, or reconstruct it from a complete bounded staging slab. Those execution choices preserve the same logical result: exact populations must cover the admitted set before its crossing is trusted.
+
+#### Turn a Threshold Search into a Boundary Problem
+
+Scanning bins from high scores to low identifies the crossing bin $j^{\ast}$ with
+
+$$
+a=\sum_{\ell\gt j^{\ast}}h_\ell\lt K,\qquad a+h_{j^{\ast}}\ge K.
+$$
+
+Every score in a higher bin is a certain winner. Every score in a lower bin is unnecessary. Only $K-a$ winners must be chosen from the crossing bin's $m=h_{j^{\ast}}$ candidates:
+
+$$
+\mathrm{TopK}(x)=\lbrace \text{all positions in higher bins}\rbrace
+\quad\cup\quad \mathrm{Top}_{K-a}(\text{crossing bin}).
+$$
+
+If the entire crossing bin is needed, it can be emitted directly. Small crossings use direct ranking; larger ones use exact order-preserving FP32 key refinement. The same decomposition underlies the 980-plus-44 example in Figure 4.
+
+The histogram discretizes the search region, not the selected scores. Exact comparisons within the crossing bin resolve its coarse boundaries, including ties. Thus fewer passes do not require approximate Top-K membership.
+
+#### How Verification Preserves Exactness
+
+**The sample focuses the bins; exact bin counts make the sample safe.** A useful bracket keeps the crossing small, while multi-threshold verification avoids a separate full-row query for every trial boundary. Figure 4 distinguishes the resulting paths: refine a complete admitted set, lower admission when too few scores survive, or recover exactly when staging or bracket checks fail.
+
+Two mechanisms have different jobs. The **admission ladder**—the primary threshold, lower floor, and conservative sentinel—widens the candidate region. The **verification bin boundaries** locate rank $K$ within that region. Non-split streaming can rescan at a lower threshold; split-row streaming can stage down to the lower floor within its scan. Overflow requires complete-set or whole-row exact recovery.
+
+Guess quality controls work, while complete counts and exact refinement control membership. There is no fixed one-pass guarantee. For example, [PR #18625](https://github.com/NVIDIA/TensorRT-LLM/pull/18625) repairs infinite-width brackets in register families with exact whole-row key selection, preserving `+inf` winners.
+
+The output contract is an exact selected **value multiset** with valid unique indices. Tied indices may differ from `torch.topk`; output order is unrestricted. Short rows return valid local indices followed by `-1` padding. NaN ordering remains implementation-specific.
+
+### Mapping Selection to Blackwell
+
+A single scheduling policy cannot serve both one short row and thousands of long rows efficiently. GVR V2 uses four kernel families, implemented in CuTe DSL:
+
+
+
+| Family | Where the scores or candidates live | Why it helps |
+| :---: | :---: | :---: |
+| `reg` | A row resides in one thread block's registers | Avoids repeated global loads when the row fits |
+| `reg_clus` | Register slices across cooperating blocks | Exposes more parallelism for medium rows at small batch sizes |
+| `clus` | Streaming shards; histograms and candidates shared within a hardware cluster | Merges through distributed shared memory |
+| `main` | Streaming scan with bounded candidate staging | Covers the remaining shapes, including long rows and large batches |
+
+
+
+A thread block is also called a cooperative thread array, or CTA. Blackwell thread-block clusters let cooperating CTAs exchange data through distributed shared memory. This reduces the need to materialize intermediate results in global memory for eligible shapes.
+
+Register families bypass sparse sampling but retain exact histogram crossing and refinement. In the normal case, the first $K$ current-row values establish the initial bracket; near the short-row regime, a whole-row bracket is used. Full-row classification and exact crossing refinement still determine the output. The `main` family can assign multiple CTAs to a row when a small batch would otherwise leave much of the GPU idle.
+
+This explains the two sources of performance improvement: a better starting threshold reduces selection work, and a suitable kernel family reduces the cost of executing that work. Neither eliminates the obligation to examine all valid scores.
+
+[PR #19076](https://github.com/NVIDIA/TensorRT-LLM/pull/19076) extends this execution policy through host dispatch while keeping the device kernels unchanged. It adds register plans for roughly 4K–8K-score rows, sizes register waves using the device's SM count, and includes targeted B300 routing. It also sets the direct-ranking gate to **96 candidates across register plans**, limiting quadratic work for difficult crossings. The streaming `main` and `clus` families retain their separate **288-candidate** gate. Dispatch therefore balances refinement cost with each family's execution strategy.
+
+## Performance and Roofline Analysis
+
+### Performance Against GVR V1 and Radix CUDA
+
+#### Benchmark Setup
+
+The benchmarks use FP32 indexer scores from the three models below on NVIDIA B200. Batch size $B$ counts score rows: each case repeats one captured row across 1 to 1,024 batch rows to measure kernel scaling. It does not represent heterogeneous serving concurrency. GVR V2 uses the merged [PR #19076](https://github.com/NVIDIA/TensorRT-LLM/pull/19076) implementation. Results report cold-L2 GPU kernel time, excluding compilation, input preparation, and Python overhead. Speedups are geometric means over matched workloads from separate benchmark runs. The grid uses single-token decode and case-matched launch envelopes; ragged batches, MTP, and prefill require separate evaluation. Historical serving results appear in the [integration section](#serving-gains-from-the-shared-engine).
+
+
+
+**$N$ is the indexer row width, not the original prompt length.** A roughly 512K-token V4 context yields a roughly 128K-wide indexer row because of 4× compression.
+
+#### Overall Results
+
+
+
+Both comparisons use the same 9,746 cases. The minimum column retains individual regressions, including the GVR V1 cases where V2 is slower; the [V1-to-V2 discussion](#from-gvr-v1-to-v2-why-move-beyond-temporal-hints) explains how informative temporal hints can retain an admission advantage. Figure 1 shows the model-level comparison over this same workload grid.
+
+#### The Gains Extend Beyond an Average
+
+The gains over TensorRT LLM radix CUDA vary with the model and shape. Figure 7 shows V2's speedup across every captured row length and all 11 batch sizes, with one panel per model.
+
+
+
+*Figure 7. TensorRT LLM radix CUDA time divided by GVR V2 time, geometrically averaged across layers at each shape. The panels share a 1–21× color scale; 1.0 means equal performance. Row lengths are rounded in the axis labels, and cell labels round to one decimal place.*
+
+The strongest shape-average gains occur around 4K scores at $B=1024$: **19.85× for V4 Flash and 20.18× for V4 Pro** at $N=4{,}099$, and **10.96× for V3.2** at $N=4{,}111$. The Pro result is the largest shape-average speedup in the grid.
+
+The advantage extends across all 275 plotted shapes. The smallest shape-average speedups are **1.55× for Flash, 1.52× for Pro, and 3.50× for V3.2**. These averages summarize layers at each shape; the overall table retains the lower minimum across individual workload cases.
+
+GVR V1 is the closer comparison. Figure 8 uses the same grid and aggregation, with a separate **1–3× color scale** to resolve the smaller differences.
+
+
+
+*Figure 8. GVR V1 time divided by GVR V2 time, geometrically averaged across layers at each shape. All three panels share a 1–3× color scale, with parity at 1.0× and cell labels rounded to two decimal places. An orange corner marks at least one constituent layer with speedup below 1×, even when the cell average is above parity.*
+
+The strongest shape-average gains are **2.87× for Flash, 2.83× for Pro, and 1.86× for V3.2**. Flash and Pro peak at $B=1024$ on short rows; V3.2 peaks near 128K scores at $B=8$.
+
+All 275 shape averages exceed parity, but layer averaging can hide local regressions. For Pro at $N=131{,}075$ and $B=512$, the shape average is **1.05×** even though one layer reaches **0.689×**. The orange markers preserve this distinction; the earlier [V1-to-V2 analysis](#from-gvr-v1-to-v2-why-move-beyond-temporal-hints) explains why informative temporal hints can still give V1 a local admission advantage.
+
+#### What Explains the Differences
+
+**TensorRT LLM radix CUDA.** The baseline uses the production dispatcher, including short-row insertion and long-row split-work paths. V2's **5.05×** advantage is consistent with reducing full-row selection passes and matching execution to the workload.
+
+**GVR V1 (temporal hint).** V1 already combines multiple admission thresholds through pivot/rescue verification. V2 replaces the temporal prior with current-row calibration and couples exact bin counts to crossing-bin refinement. Its **1.46×** gain compares complete implementations, including their execution policies; it does not isolate the contribution of self-sampling alone.
+
+#### Latency Across Row Length and Batch Size
+
+
+
+*Figure 9. Mean cold kernel time across all captured layers: 21 for Flash, 30 for Pro, and 61 for V3.2. Each row is a model; the columns contrast batch sizes 1 and 1,024. The solid line shows GVR V2; dashed lines show baselines measured in separate runs. Both axes are logarithmic; 1K means 1,024.*
+
+At $B=1$, keeping a short row in registers and exposing parallelism within a longer row matter more than saturating HBM. At $B=1024$, streaming throughput becomes more visible. The different shapes of these curves are why a single average cannot identify every useful operating region.
+
+### The Roofline Model: Fewer Passes, More Useful Work
+
+The bar chart shows how much time GVR V2 saves. The roofline compares that time with an ideal one-read, index-write traffic model. It connects the algorithm's goal—fewer full-row passes—to a bandwidth reference.
+
+#### Locate Top-K on the Hardware Roof
+
+For FP32 input and INT32 index-only output, define the ideal work and minimum traffic as
+
+$$
+W=BN,\qquad Q_{\min}=4B(N+K)\ \text{bytes},\qquad
+I=\frac{W}{Q_{\min}}=\frac{N}{4(N+K)}.
+$$
+
+Each work unit is an **abstract comparison per input score**, shared across implementations; it is neither an instruction count nor SELECT's comparison bound. For nontrivial selection, $1\le K\lt N$, intensity lies in **$0.125\lt I\lt 0.25$ compare/byte**.
+
+The theoretical and calibrated B200 roofs, in Tcompare/s, are
+
+$$
+P_{\mathrm{theory}}(I)=\min(37.225,8I),\qquad
+P_{\mathrm{calibrated}}(I)=\min(37.047,6.912I).
+$$
+
+The calibrated limits are **6.912 TB/s** sustained read bandwidth and **37.047 Tcompare/s** semantic comparison throughput. Their **5.36 compare/byte** intersection exceeds the maximum ideal Top-K intensity by over 21×, placing the workload band on Figure 10A's bandwidth slope.
+
+
+
+*Figure 10.* A: theoretical and calibrated roofs. B: Pareto curves at $B=1024$, plotting useful throughput $P=BN/t$ against ideal intensity $I=N/[4(N+K)]$. Green highlights V2; the dotted line is the calibrated bandwidth roof. All kernels share $Q_{\min}$; extra work remains in measured time. This measures useful work relative to ideal traffic, not actual DRAM utilization.
+
+#### Compare Pareto Curves and Reachable Rates
+
+Each operator's **Pareto curve** in Figure 10B is its measured intensity–throughput trace across row lengths at $B=1024$. At fixed $N$ and $K$, every implementation has the same horizontal position; a faster kernel moves **upward**, toward the calibrated roof. Figure 9 retains the contrasting single-row view, and Figures 7–8 cover all 11 batch sizes.
+
+The **reachable rate** is useful throughput divided by the calibrated roof, expressed as a percentage. Throughout the plotted bandwidth branch, its underlying ratio simplifies to
+
+$$
+\rho=\frac{P(I)}{P_{\mathrm{calibrated}}(I)}
+=\frac{Q_{\min}}{\mathrm{BW}\,t}.
+$$
+
+It measures efficiency relative to the ideal traffic bound. The table compares **average / peak reachable rate** along each Pareto curve: the average weights the plotted intensity points equally, and the peak is their maximum. Both use the same layer-averaged timings as Figure 10B.
+
+
+
+*Each cell shows average / peak.*
+
+GVR V2 leads both measures on all three models: its average reachable rate is **39.0–41.6%**, with peaks of **66.5–77.8%**. GVR V1 is the stronger baseline by both measures, averaging **25.4–27.8%**, with peaks of **51.1–63.3%**. On V3.2, it reaches **27.8% / 51.1%**, compared with V2's **41.5% / 66.5%**. Reporting both measures captures the best operating point and the performance sustained across the curve.
+
+#### Interpret the Remaining Gap
+
+The full ideal bound is $T_{\mathrm{roof}}=\max(Q_{\min}/\mathrm{BW},W/R_{\mathrm{compare}})$. Within the Top-K band, the bandwidth term sets this bound. A kernel's actual bottleneck can also involve histogram updates, synchronization, register pressure, candidate staging, or exact refinement. Attributing the measured gap to these mechanisms requires counters or ablations beyond the latency comparisons. Short-row identity/padding paths can skip score reads and lie outside this model.
+
+The pass/candidate tradeoff in Figure 5 explains two sources of this gap: another full-row scan adds traffic, while a loose admission threshold adds candidate work even when the scan count stays fixed. V2 uses exact bin populations to restrict the remaining selection to the crossing bin. These costs increase measured time; under the shared minimum-traffic model, they move useful throughput downward at the same intensity. The read-dominated roof is optimistic, especially when $K/N$ is large; the plotted throughput describes useful selection work rather than measured DRAM traffic.
+
+## TensorRT LLM Integration and Takeaways
+
+### Decode and Prefill in TensorRT LLM
+
+The integration follows three boundaries: remove the Top-K prior lifecycle, adapt each phase's row metadata, and prepare launch specializations before graph capture.
+
+#### Remove the Top-K Prior Lifecycle
+
+A temporal prior carries state from an earlier selection. Its buffers, initialization, request alignment, and write-back must stay valid through CUDA Graph warmup and replay. Disaggregated prefill/decode also needs a policy for providing that prior at the phase handoff.
+
+With `use_self_sampling_topk: true`, the [`TopK` wrapper](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/modules/top_k.py#L350-L392) selects V2's self-sampling decode [`run_varlen`](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode_self_sampling_host.py#L1435-L1445) entry point. Its public signature takes score, sequence-length, and output-index tensors plus row/launch metadata; it has **no `pre_idx` parameter**. The wrapper's [`needs_gvr_prior`](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/modules/top_k.py#L69-L75) property is false for this path.
+
+The separate V1 temporal dispatcher, [`tiered_topk`](https://github.com/NVIDIA/TensorRT-LLM/blob/be1b9885e8df9bf070e8cb68459e24a7119afaa9/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/top_k/gvr_topk_decode_dispatch.py#L334-L357), still requires valid `pre_idx` for its supported heuristic routes. Setting `use_self_sampling_topk: false` selects that temporal engine; callers must retain its prior input.
+
+V2 derives its bracket from the current scores in both phases, removing the previous-step Top-K buffer, prefill-to-decode prior seeding, and prior write-back from the self-sampling path. Request metadata and launch preparation remain, but they no longer maintain a selection history. [PR #18446](https://github.com/NVIDIA/TensorRT-LLM/pull/18446) confines prior ownership to the temporal engine.
+
+#### One Selection Core, Two Row Interfaces
+
+TensorRT LLM separates phase-specific row metadata from shared selection logic. The `TopK` module applies the same self-sampling configuration to supported decode and prefill paths, then passes each phase's valid score interval to the selected engine.
+
+
+
+*Figure 11. Current-row calibration lets both phases enter the same selection core without a temporal-prior lifecycle. The decode arrow shows its streaming route; register and cluster routes remain available under the same output contract. Prefill specializes the streaming implementation at compile time.*
+
+The adapters preserve each phase's indexing semantics. Decode derives valid prefixes from device KV lengths, multi-token prediction offsets, and compression. Prefill receives `[start, end)` in compressed columns and returns indices relative to `start`. Both write INT32 indices into caller-owned output, with identity indices and `-1` padding for short rows.
+
+Prefill's compile-time mode changes addressing, masking, and index origin in the same **`GvrMainKernel` streaming implementation** used by decode. Self-sampling, multi-threshold counting, collection, and exact refinement stay shared. One block per prefill row and specialized decode scheduling let the common algorithm serve different parallelism requirements. [PR #18702](https://github.com/NVIDIA/TensorRT-LLM/pull/18702) implements this reuse.
+
+This separation keeps phase-specific decisions in the wrapper, dispatch, and compile-time row adapters. Improvements to threshold calibration, counting, and exact recovery can serve both phases through the shared streaming body, reducing duplicated selection logic and the number of state transitions the framework must maintain.
+
+#### Stable Launches, Dynamic Row Lengths
+
+Host routing chooses a stable launch envelope; device metadata supplies each row's actual length. The physical memory bound remains separate from the bound used to select an execution plan. Before capture, warmup must prepare exact-row decode launchers and a bounded set of prefill tiers and width buckets, using the same envelope-dependent plan rule as execution and readiness checks.
+
+Prefill and decode have distinct specialization keys; decode launcher and warmup caches also distinguish SM count and architecture. This lets dispatch tune execution, including sampled prefill plans for qualifying small envelopes, while preserving the graph's prepared launch contract. These rules are implemented through [PR #18683](https://github.com/NVIDIA/TensorRT-LLM/pull/18683), [PR #18702](https://github.com/NVIDIA/TensorRT-LLM/pull/18702), and [PR #19076](https://github.com/NVIDIA/TensorRT-LLM/pull/19076).
+
+When scores fail the `TopK` module's dtype, stride, or alignment gate, it selects native insertion/radix. Prefill also selects radix for an all-short tile or a missing capture-time specialization. Decode requires its exact launcher key to be warmed before capture; a missing key raises an error. Invalid low-level tensor or storage contracts likewise raise, so fallback is a dispatch policy, not a catch-all for engine errors.
+
+#### Serving Gains from the Shared Engine
+
+In [PR #18702's B200 profiles](https://github.com/NVIDIA/TensorRT-LLM/pull/18702), the aggregate Top-K kernel time within measured Flash/Pro prefill windows improves by **1.84–2.61×** over radix CUDA. Adding V2 prefill to a deployment already using V2 decode improves serving throughput by **2.9–4.5%** on the tested long-input, batched workloads. These are historical integration results and measure the incremental benefit of the prefill change.
+
+For decode, [PR #18410](https://github.com/NVIDIA/TensorRT-LLM/pull/18410) reports **6–19% lower time per output token** on 8×B200 with TP8/EP8 and batch/concurrency 1. The serving benefit depends on Top-K's share of total execution time; the kernel comparisons above do not measure competing serving stacks.
+
+#### Enable GVR V2
+
+Use a TensorRT LLM revision containing [PR #19076](https://github.com/NVIDIA/TensorRT-LLM/pull/19076). The production V2 dispatch requires CUTLASS DSL, datacenter Blackwell SM100/SM103, $K\in\lbrace 512,1024,2048\rbrace$, and indexer compression ratios 1 or 4. Save the following as `gvr_v2.yaml`:
+
+```yaml
+sparse_attention_config:
+ algorithm: deepseek_v4
+ enable_heuristic_topk: true
+ use_self_sampling_topk: true
+```
+
+Use `algorithm: dsa` for DeepSeek-V3.2. The checkpoint supplies the model's Top-K width. With a prepared benchmark dataset in `dataset.jsonl`, run:
+
+```bash
+trtllm-bench --model deepseek-ai/DeepSeek-V4-Flash throughput \
+ --dataset dataset.jsonl \
+ --config gvr_v2.yaml \
+ --tp 8 --ep 8
+```
+
+`enable_heuristic_topk` defaults to `false`. Once it is enabled, `use_self_sampling_topk` defaults to `true`; the second field is explicit here for clarity. Supported prefill layers follow the same V2 selection. Setting `use_self_sampling_topk: false` selects the temporal-hint decode dispatcher, which chooses among temporal kernel routes; prefill remains radix.
+
+The fast path expects FP32 scores with unit inner stride, a row stride divisible by four floats, and a 16-byte-aligned base. A single-row decode input also needs its physical width divisible by four; its valid prefix may be shorter. The retired `TRTLLM_GVR_SELF_SAMPLING` environment variable is no longer the enablement mechanism, and `use_cute_dsl_topk` is not required to select V2.
+
+### Conclusion
+
+GVR V2 follows from a practical limit of temporal prediction: a biased hint can be excellent when overlap is high, yet expensive to rely on when quality fluctuates. **Self-sampling calibrates from the current row; multi-threshold counts locate the crossing; exact refinement resolves the remaining membership.** Together, they target both difficult-input performance and average latency while preserving exactness.
+
+The performance map shows the gains across shapes, and the Pareto curves relate them to the bandwidth roof. Removing the temporal prior also removes its framework lifecycle: a shared streaming implementation serves prefill and decode through phase-specific row interfaces. The result is one algorithmic core whose calibration depends on the input it is selecting now.
+
+### Further Reading
+
+[Benchmark methodology and figure reproduction](../media/gvr_v2/README.md) are available separately.
+
+Selection background:
+
+- [Floyd and Rivest, *Expected Time Bounds for Selection* (1975)](https://people.csail.mit.edu/rivest/pubs/FR75a.pdf), *Communications of the ACM* 18(3), 165–172: the sampling-based selection idea behind V2's calibration.
+- [Kiwiel, *Randomized Selection with Quintary Partitions*](https://arxiv.org/abs/cs/0312055): rigorous analysis of SELECT variants and repeated-key handling.
+
+Implementation milestones:
+
+- [PR #17821: original self-sampling decode integration](https://github.com/NVIDIA/TensorRT-LLM/pull/17821).
+- [PR #18410: current-row brackets and hint-free production API](https://github.com/NVIDIA/TensorRT-LLM/pull/18410).
+- [PR #18625: exact handling of positive infinity in register families](https://github.com/NVIDIA/TensorRT-LLM/pull/18625).
+- [PR #18646: device-side prior seeding for the temporal path](https://github.com/NVIDIA/TensorRT-LLM/pull/18646).
+- [PR #18683: physical envelopes and exact-row warmup](https://github.com/NVIDIA/TensorRT-LLM/pull/18683).
+- [PR #18446: configuration-based dispatch and prior-state removal for V2](https://github.com/NVIDIA/TensorRT-LLM/pull/18446).
+- [PR #18702: self-sampling prefill](https://github.com/NVIDIA/TensorRT-LLM/pull/18702).
+- [PR #19076: register-plan tuning, a unified crossing-bin gate, sampled prefill plans, and SM-aware B200/B300 dispatch](https://github.com/NVIDIA/TensorRT-LLM/pull/19076).
+
+For the surrounding model pipeline, see [Sparse Attention in TensorRT LLM](https://nvidia.github.io/TensorRT-LLM/blogs/tech_blog/blog17_Sparse_Attention_in_TensorRT-LLM.html) and [DeepSeek-V4 on NVIDIA Blackwell](https://nvidia.github.io/TensorRT-LLM/blogs/tech_blog/blog26_DeepSeek_V4_on_NVIDIA_Blackwell_Model_Specific_and_Agentic_Workload_Optimizations_in_TensorRT-LLM.html).
diff --git a/docs/source/conf.py b/docs/source/conf.py
index b60f66926686..9ac1821c2ca1 100644
--- a/docs/source/conf.py
+++ b/docs/source/conf.py
@@ -152,11 +152,11 @@
# -- Options for HTML output -------------------------------------------------
# https://www.sphinx-doc.org/en/master/usage/configuration.html#options-for-html-output
+# JSON files are downloadable data assets, not documentation source files.
source_suffix = {
'.rst': 'restructuredtext',
'.txt': 'markdown',
'.md': 'markdown',
- '.json': 'json',
}
html_theme = 'nvidia_sphinx_theme'