Skip to content

[SM100] Paged MQA logits: optional exact coarse score histogram for LiteTopK decode - #94

Open
Heisenberg-Yin wants to merge 1 commit into
sgl-project:devfrom
Heisenberg-Yin:litetopk-decode
Open

Heisenberg-Yin wants to merge 1 commit into
sgl-project:devfrom
Heisenberg-Yin:litetopk-decode

Conversation

@Heisenberg-Yin

@Heisenberg-Yin Heisenberg-Yin commented Sep 29, 2026 •

Copy link
Copy Markdown

@DarkSharpness could you take another look? This revision is rebuilt on dev 7c4a95f. After deepseek-ai#462 / deepseek-ai#471, SM100 paged MQA logits are BF16-only, so the previous design (FP32 histogram, token Q tile, a separate BF16 producer) has nothing left to attach to. The histogram is now an optional output of dev's own kernel. It is the producer side of LiteTopK decode (sgl-project/sglang#40764).

Motivation

In DSA and DeepSeek-V4.1 decode, the top-k after the paged MQA logits starts with a pass that only builds a coarse histogram of each row. The logits kernel already holds every score in registers, so it can count that histogram while it writes the logits, and the SGLang selector skips the pass.

What this adds

  • API: fp8_fp4_paged_mqa_logits(..., indices=None, histogram=None). histogram is an optional int32 [rows, 1024] tensor (contiguous, 8-byte aligned).

    • The kernel adds to it: for every stored logit (column < context_len, not NaN) it increments one bin. Callers zero the tensor first, for example with a memset node in the same CUDA graph.
    • Bins run in descending order: FP16-RN exponent + 4 mantissa bits for |x| < 16, unit bins up to 223, a saturating bin at each end, and -0 counted as +0.
    • Logits are bit-identical with and without the histogram.
  • Coverage: SM100, MXFP8 and MXFP4, 32 or 64 heads, pages of 32 / 64 / 128, any number of rows per request. SM90 and other head counts are rejected with a clear error.

  • Without a histogram nothing changes. The generated source and SASS of all 49 compared kernels equal dev's. The only difference is that the parameter block grows by 8 bytes.

  • epilogue/coarse_histogram.cuh (new):

    • The bin encoding.
    • RowWindowHistogram: 16-bit counters in shared memory for a window of two Q blocks. They are flushed with packed 64-bit global adds when a task leaves the window, after 170 splits, and at exit.
    • Scores are binned two per register on a warp-uniform path (all fine, all unit, or mixed). Anything else takes an exact scalar path.
  • Kernel, scheduler and host:

    • sm100_mqa_logits.cuh gains if constexpr hooks in the shared core.
    • The paged scheduler caps a task at 170 splits, only when the histogram is on, so the 16-bit counters cannot wrap.
    • MXFP4 at head dim 128 keeps 7 KV stages instead of 8 to make room for the window. MXFP8 is unchanged.
  • Schema:

    fp8_fp4_paged_mqa_logits(Tensor q, Tensor? q_sf, Tensor fused_kv_cache, Tensor weights, Tensor context_lens,
        Tensor block_table, Tensor schedule_meta, int max_context_len, Tensor? indices, Tensor(a!)? histogram) -> Tensor
    

One commit: 12 files, +617 / -34. The previous revision had 19 files, +2150 / -44.

Changes since the previous revision (b8648e3 on dev b296ce3)

  • Rebuilt on dev 7c4a95f. Removed:
    • the FP32 histogram and the token Q tile (fp8_paged_mqa_logits and FP32 logits no longer exist);
    • the separate BF16 producer and tokens_per_request (dev's kernel is that kernel now);
    • the 16384-row assert.
  • Review comments:
    • The flush is now a single flush in coarse_histogram.cuh. It uses plain comparisons and adds bin pairs through pack_bin_pair.
    • There is now one histogram type, RowWindowHistogram, plus the disabled default NoHistogram.
  • Not included: the previous BF16 producer's 6-token Q tile. It will be a separate PR; its cost is shown under "Producer cost" below.

Validation (B200, CUDA 13.0, torch 2.14.1)

  • test_paged_mqa_logits_histogram (new, also in run_tests.sh's release sanitizer list): 128 / 128 cases, edge lengths 0-131,071 with 1-6 rows per request (PDL off and on), forced SM counts 1 / 7 / 100, and CUDA graph replay.
  • test_attention.py, release profile: prefill 4 / 240, paged 4 / 900, histogram 4 / 128, sparse 8 / 161.
  • compute-sanitizer memcheck and synccheck on the histogram test: 0 errors.
  • SASS against dev:
    • 49 / 49 kernels without a histogram are identical (38 paged, 8 contiguous, 3 metadata);
    • all 30 histogram variants use 128 registers with no spills.
  • Bin math, checked exhaustively: for all 65,536 BF16 values and all 2^32 packed pairs, the bins equal the FP32 definition.
  • [Perf] [NVIDIA] attention: Add opt-in LiteTopK decode top-k for DSA and DeepSeek-V4.1 on SM100 sglang#40764: its kernel test passes 23 / 23 with this build.

Producer cost (real data)

Method: B200, cold (1 GiB flush, then a single-call CUDA graph replay). The same package runs with and without the histogram, interleaved. The A/A spread is 0.98-1.01.

GLM-5.2, index layers 0 and 38: MXFP8 with 32 heads, converted the way sgl-project/sglang#44350 does. 383 cells: B 1-128 × L 8K-1M × next_n 1-4.

  • Geomean +2.28 %; median extra +0.42 µs. The maximum is +9.84 % (B 4, 8K), and no cell is above +10 %.
  • The cost falls with batch size: +4.0-4.3 % at B 1, +0.5-0.6 % at B 128. At ~1M it is at most +0.6 %.

DeepSeek-V4.1, layer 20: MXFP4 with 32 heads. 142 cells, averaged over two runs.

tokens per request 1 2 3 4 5 6
geomean +1.25 % +4.18 % +5.46 % +5.38 % +13.85 % +14.54 %
median extra (µs) +0.14 +0.64 +0.98 +0.69 +1.87 +2.66
  • 5-6 tokens per request: the cost is highest here, up to +21 % at B 64, ~1M (that cell was power-capped).
  • Against the previous revision's BF16 producer, both with the histogram:
    • at 1-4 tokens per request the two are on par (1.002-1.008×);
    • at 5 / 6 tokens, this revision is 1.135× / 1.114× slower, because it runs dev's 4-row Q tile where the old producer ran its 6-token tile;
    • the follow-up 6-token tile PR targets this;
    • the logits of the two producers are bit-identical in all 108 compared cells.

@Heisenberg-Yin

Heisenberg-Yin commented Sep 29, 2026 •

Copy link
Copy Markdown
Author

@DarkSharpness could you take a look? This is the producer side for LiteTopK decode. The paged MQA logits kernel emits the coarse histogram that the SGLang top-k v2 streaming and cluster paths currently build in their first pass over the logits.

@Heisenberg-Yin Heisenberg-Yin changed the title [SM100] Paged MQA logits: fused coarse score histogram for LiteTopK decode + per-next_n Q tile [SM100] Paged MQA logits: fused coarse score histogram for LiteTopK decode Sep 29, 2026
@Heisenberg-Yin
Heisenberg-Yin marked this pull request as ready for review September 29, 2026 08:33
Heisenberg-Yin added a commit to Heisenberg-Yin/sglang that referenced this pull request Sep 30, 2026
- select.cuh + coarse_bins.cuh: exact FP32 top-2048 selector driven by
  DeepGEMM's coarse score histogram, compiled on first use through
  load_jit. Covers decode (next_n = 1) and speculative verify
  (next_n <= 4). A stale histogram, or a crossing bin larger than the
  candidate capacity, falls back to an exact whole-row radix select
  instead of raising a device error. Ties go to the lower physical slot.
- fused.py: FusedDecodePlan(batch, next_n, device, candidate_capacity=8192)
  calls deep_gemm.fp8_fp4_paged_mqa_logits(histogram=...) directly.
- Drop the CuTe DSL selector, its AOT build and manifest, and the vendored
  DeepGEMM patch and private build: the histogram argument now comes from
  sgl-project/DeepGEMM#94.
- smoke_fused.py: decode and verify, eager and CUDA graph replay, plus a
  stale histogram and a candidate overflow.
- Add kernels/experimental/.clang-format, identical to kernels/jit's.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@Heisenberg-Yin Heisenberg-Yin changed the title [SM100] Paged MQA logits: fused coarse score histogram for LiteTopK decode [SM100] Paged MQA logits: exact coarse score histograms for LiteTopK decode (FP32 + BF16 DeepSeek-V4.1 producers) Oct 9, 2026
Heisenberg-Yin added a commit to Heisenberg-Yin/sglang that referenced this pull request Oct 9, 2026
One exact, histogram-driven decode top-k selector for the DSA indexer
(FP32 scores, top-2048, 64-slot pages) and the DeepSeek-V4.1 ratio-1/2
index layers (BF16 scores, top-512, 128-slot pages). DeepGEMM's paged MQA
logits kernel counts each row's scores into 1024 coarse bins while it
writes them; the selector derives the crossing bin from the histogram,
selects everything above it and ranks the bin exactly, falling back to an
exact whole-row select when the histogram disagrees with the scores or the
crossing bin overflows the candidate buffer. Rows of up to 10240 (FP32) /
16384 (BF16) scores are selected by one CTA from registers. Scratch returns
to rest after every call, so CUDA graph replays need no reset.

- kernels/jit/csrc/litetopk_decode/select.cuh: the selector, templated over
  the two shipped configurations.
- kernels/ops/attention/litetopk_decode.py: configs, storage / plan, and
  runtime detection of the device and the DeepGEMM producer API.
- srt/layers/attention/litetopk_decode.py: per-process selector buffers,
  plans per row count, CUDA-graph-safe growth.
- DSA: Indexer._get_topk_paged passes the plan's histogram to
  fp8_paged_mqa_logits and replaces the fused paged top-k transform for
  decode and verify (<= 4 tokens per request); when LiteTopK serves verify,
  verify tokens get their causal lengths in the producer schedule.
- DeepSeek-V4.1: FullTopKIndexer.topk_decode and
  SparseTableBackend.publish_decode use the BF16 producer and selector; the
  producer schedule is rebuilt once per decode / verify forward; candidate
  block maxima accept BF16 logits.

Off by default (SGLANG_OPT_LITETOPK_DECODE); unsupported devices, DeepGEMM
builds, modes and shapes keep the default top-k. The histogram producers
come from sgl-project/DeepGEMM#94.

Also covers DeepSeek-V4.1, replacing sgl-project#42351.
Heisenberg-Yin added a commit to Heisenberg-Yin/sglang that referenced this pull request Oct 9, 2026
One exact, histogram-driven decode top-k selector for the DSA indexer
(FP32 scores, top-2048, 64-slot pages) and the DeepSeek-V4.1 ratio-1/2
index layers (BF16 scores, top-512, 128-slot pages). DeepGEMM's paged MQA
logits kernel counts each row's scores into 1024 coarse bins while it
writes them; the selector derives the crossing bin from the histogram,
selects everything above it and ranks the bin exactly, falling back to an
exact whole-row select when the histogram disagrees with the scores or the
crossing bin overflows the candidate buffer. Rows of up to 10240 (FP32) /
16384 (BF16) scores are selected by one CTA from registers. Scratch returns
to rest after every call, so CUDA graph replays need no reset.

- kernels/jit/csrc/litetopk_decode/select.cuh: the selector, templated over
  the two shipped configurations.
- kernels/ops/attention/litetopk_decode.py: configs, storage / plan, and
  runtime detection of the device and the DeepGEMM producer API.
- srt/layers/attention/litetopk_decode.py: per-process selector buffers,
  plans per row count, CUDA-graph-safe growth.
- DSA: Indexer._get_topk_paged passes the plan's histogram to
  fp8_paged_mqa_logits and replaces the fused paged top-k transform for
  decode and verify (<= 4 tokens per request); when LiteTopK serves verify,
  verify tokens get their causal lengths in the producer schedule.
- DeepSeek-V4.1: FullTopKIndexer.topk_decode and
  SparseTableBackend.publish_decode use the BF16 producer and selector; the
  producer schedule is rebuilt once per decode / verify forward; candidate
  block maxima accept BF16 logits.

Off by default (SGLANG_OPT_LITETOPK_DECODE); unsupported devices, DeepGEMM
builds, modes and shapes keep the default top-k. The histogram producers
come from sgl-project/DeepGEMM#94.

Also covers DeepSeek-V4.1, replacing sgl-project#42351.
Heisenberg-Yin added a commit to Heisenberg-Yin/sglang that referenced this pull request Oct 16, 2026
…nd DeepSeek-V4.1 on SM100

One exact, histogram-driven decode top-k selector for the DSA indexer
(FP32 scores, top-2048, 64-slot pages) and the DeepSeek-V4.1 ratio-1/2
index layers (BF16 scores, top-512, 128-slot pages). DeepGEMM's paged MQA
logits kernel counts each row's scores into 1024 coarse bins while it
writes them; the selector derives the crossing bin from the histogram,
selects everything above it and ranks the bin exactly, falling back to an
exact whole-row select when the histogram disagrees with the scores or the
crossing bin overflows the candidate buffer. Rows of up to 10240 (FP32) /
16384 (BF16) scores are selected by one CTA from registers. Scratch returns
to rest after every call, so CUDA graph replays need no reset. Grid sizes
follow the GPU's SM count, which the storage passes to every call; the row
thresholds between launch shapes were measured on B200 (148 SMs).

- kernels/jit/csrc/litetopk_decode/select.cuh: the selector, templated over
  the two shipped configurations.
- kernels/ops/attention/litetopk_decode.py: configs, storage / plan, and
  runtime detection of the device and of every keyword the serving calls
  pass to the DeepGEMM producer API (declared Python parameters).
- srt/layers/attention/litetopk_decode.py: per-process selector buffers,
  plans per row count, CUDA-graph-safe growth.
- DSA: Indexer._get_topk_paged passes the plan's histogram to
  fp8_paged_mqa_logits and replaces the sgl-kernel fused paged top-k
  transform for decode and verify (<= 4 tokens per request); when LiteTopK
  serves verify, verify tokens get their causal lengths in the producer
  schedule. The torch paged MQA logits fallback and the flashinfer / torch
  top-k backends keep the default top-k.
- DeepSeek-V4.1: FullTopKIndexer.topk_decode and
  SparseTableBackend.publish_decode use the BF16 producer and selector in
  place of the default top-k (topk_transform_paged_v2 on FP32 logits, or
  DeepSelect on the BF16 logits of SGLANG_OPT_DSV41_BF16_INDEX_LOGITS); the
  producer schedule is rebuilt once per decode / verify forward.

Off by default (SGLANG_OPT_LITETOPK_DECODE); unsupported devices, DeepGEMM
builds, modes and shapes keep the default top-k. The histogram producers
come from sgl-project/DeepGEMM#94.

Also covers DeepSeek-V4.1, replacing sgl-project#42351.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
int* row_bins = global + static_cast<uint64_t>(row) * kBins + 2 * threadIdx.x;
#pragma unroll
for (uint32_t i = 0; i < kPairsPerThread; ++ i) {
if ((counts[i].x | counts[i].y) != 0) {

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

why not use normal count[i].x == 0 || count[i].y == 0 (or first pack into 1 unsigned long long)? this gives very poor readability.

// operations (8 * bin = 0x47f8 - (bits & 0xfff8) for positive, (bits & 0xfff8) - 0xa800 for negative scores); a split
// where some lane of the warp holds another value (NaN, +-Inf, 0, |x| < 2^-14 or >= 16) takes the exact scalar path.
template <uint32_t kMathThreads, uint32_t kBlockQ, uint32_t kWindowRows, bool kSwizzle = false>
struct CoarseHistogram {

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

it seems that there are 2 class with name CoarseHistogram? why?

…iteTopK decode

`fp8_fp4_paged_mqa_logits(..., indices=, histogram=)` can add an exact
1024-bin coarse histogram of every stored BF16 logit of each row to a
caller-zeroed int32 `[rows, 1024]` tensor, so a downstream top-k (LiteTopK
decode in SGLang) skips its own histogram pass over the logits. SM100,
MXFP8 and MXFP4, 32 or 64 heads; logits are bit-identical either way.

- epilogue/coarse_histogram.cuh: the bin encoding (FP16-RN exponent and 4
  mantissa bits below 16, unit bins to 223, saturating ends; BF16 form) and
  RowWindowHistogram: 16-bit counters for two Q blocks of rows in shared
  memory, published with packed 64-bit global adds when a task leaves the
  window, after 170 splits, and at exit. Per split the scores are binned two
  per register on a warp-uniform path: all fine (|x| < 16), all unit
  (|x| >= 16, e.g. GLM indexer layers), mixed, or else the exact scalar path.
  The packed paths are verified against the FP32 definition over all pairs
  of BF16 values.
- sm100_mqa_logits.cuh: `if constexpr` hooks in the shared core (initialize,
  prepare per task, observe per split, finish after the final barrier) and
  `kWithHistogram` + `int* histogram` on the paged entry. Without a
  histogram the generated source and SASS are unchanged; the parameter block
  grows by 8 bytes.
- Paged scheduler: optional per-task split cap, used only by the histogram.
- JIT host: the window follows the shared storage; MXFP4 with head dim 128
  keeps 7 KV stages (and 7-split chunks, so multi-row requests still reuse
  KV), MXFP8 is unchanged.
- API, TORCH_LIBRARY / pybind bindings, wrappers: `Tensor(a!)? histogram`,
  `histogram=None`; SM90 and other head counts are rejected.
- Tests: test_paged_mqa_logits_histogram (exact against a reference over the
  stored logits, accumulation, logits bit-identical; fine / unit / mixed /
  NaN / inf scores, empty and 1M rows, 1-6 rows per request, MXFP8 and MXFP4
  with 32 and 64 heads, forced SM counts, PDL on / off, CUDA graph replay),
  also in the release sanitizer run of run_tests.sh; README.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@Heisenberg-Yin Heisenberg-Yin changed the title [SM100] Paged MQA logits: exact coarse score histograms for LiteTopK decode (FP32 + BF16 DeepSeek-V4.1 producers) [SM100] Paged MQA logits: optional exact coarse score histogram for LiteTopK decode Oct 18, 2026
Heisenberg-Yin added a commit to Heisenberg-Yin/sglang that referenced this pull request Oct 18, 2026
…nd DeepSeek-V4.1 on SM100

One exact, histogram-driven decode top-k selector for the DSA indexer
(top-2048, 64-slot pages) and the DeepSeek-V4.1 ratio-1/2 index layers
(top-512, 128-slot pages), both reading the BF16 logits of DeepGEMM's
SM100 paged MQA logits kernel. The kernel counts each row's scores into
1024 coarse bins while it writes them (histogram=); the selector derives
the crossing bin from the histogram, selects everything above it and ranks
the bin exactly (ties to the lower physical slot), falling back to an exact
whole-row select when the histogram disagrees with the scores or the
crossing bin overflows the candidate buffer. Rows of up to 16384 scores
are selected by one CTA from registers. Scratch returns to rest after
every call, so CUDA graph replays need no reset. Grid sizes follow the
GPU's SM count, which the storage passes to every call; the row thresholds
between launch shapes were measured on B200 (148 SMs).

- kernels/jit/csrc/litetopk_decode/select.cuh: the selector, templated over
  the two shipped configurations (Bf16Top2048, Bf16Top512).
- kernels/ops/attention/litetopk_decode.py: configs, storage / plan, and
  runtime detection of the device and of DeepGEMM's BF16 paged MQA logits
  API with a histogram (declared Python parameters: indices and histogram,
  no logits_dtype).
- srt/layers/attention/litetopk_decode.py: per-process selector buffers,
  plans per row count, CUDA-graph-safe growth.
- DSA, on the MXFP8 (UE8M0) index-K cache: Indexer._get_topk_paged passes
  the plan's histogram through deep_gemm_mqa_logits.fp8_paged_mqa_logits,
  which then returns the BF16 logits without widening them, and replaces
  the sgl-kernel fused paged top-k transform for decode and verify; when
  LiteTopK serves verify, verify tokens get their causal lengths in the
  producer schedule. The torch paged MQA logits fallback and the
  flashinfer / torch top-k backends keep the default top-k.
- DeepSeek-V4.1: FullTopKIndexer.topk_decode and
  SparseTableBackend.publish_decode count the histogram with
  fp8_fp4_paged_mqa_logits and select in place of the default top-k
  (topk_transform_paged_v2 on the widened FP32 logits).
  Decode reuses the forward's DeepGEMM schedule; verify schedules the
  tokens of a request as one request (one KV pass), once per forward.

Off by default (SGLANG_OPT_LITETOPK_DECODE); unsupported devices, DeepGEMM
builds, modes and shapes keep the default top-k. Builds on sgl-project#44350 (the
fp8_fp4 MQA logits API); the histogram producer comes from
sgl-project/DeepGEMM#94.

Also covers DeepSeek-V4.1, replacing sgl-project#42351.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants