Skip to content

[kernel] Add opt-in deterministic CuTe-DSL DSA top-k backend (--dsa-topk-backend cutedsl)#31115

Open
zkyue wants to merge 3 commits into
sgl-project:mainfrom
zkyue:deterministic-dsa-topk-cutedsl
Open

[kernel] Add opt-in deterministic CuTe-DSL DSA top-k backend (--dsa-topk-backend cutedsl)#31115
zkyue wants to merge 3 commits into
sgl-project:mainfrom
zkyue:deterministic-dsa-topk-cutedsl

Conversation

@zkyue

@zkyue zkyue commented Jul 14, 2026

Copy link
Copy Markdown

Motivation

The deterministic-inference effort (#15041) currently has a gap at DSA indexer top-k selection:

  • The default sgl-kernel fast_topk_v2 radix kernel is very fast, but by design it emits winners in shared-memory-atomic arrival order, resolves the k-th-boundary tie set by arrival order, and (as documented in Reduce topk kernel shared memory from 128KB to 32KB for better occupancy #17747) stages threshold-bin candidates in a fixed 4K-entry smem window where overflow candidates are gracefully skipped. These are reasonable throughput tradeoffs, but they mean the output — and under exact score ties even the selected token set — can differ run to run for identical input.
  • This matters in practice because DSA indexer logits (sum_h relu(q·k) * w) contain a large exact-0.0 tie plateau wherever every head's dot product is negative. On such realistic inputs we measured the default kernel returning a different top-k index set on every invocation (see A/B below). Downstream, a different gather set breaks bitwise reproducibility of sparse attention even when everything else is deterministic.
  • The only deterministic option today is the non-default flashinfer path (SGLANG_DSA_TOPK_FLASHINFER_DETERMINISTIC=1), which routes through the unfused masking wrapper.

This PR adds an opt-in deterministic top-k backend so a deterministic DSA stack does not need to leave the default kernel family: --dsa-topk-backend cutedsl.

Modifications

  • python/sglang/jit_kernel/cutedsl_topk.py (new): CuTe-DSL SIMT top-k selector, one CTA per row, 512 threads. Coarse 8-bit histogram pass → MSB-first radix refinement to the exact threshold key → smallest-index tie resolution at the k-th boundary via an index radix-select → canonical descending bitonic sort of unique (score, index) composites. Properties:
    • Bitwise deterministic: no atomic whose arrival order can affect any output bit (histogram atomics only add; compaction atomics only permute buffers that feed order-independent histograms or the canonical sort). Output order is canonical: score descending, index ascending on bitwise-equal scores.
    • Exact under any distribution: threshold-bin overflow falls back to full-row rescans (identical results, slower) instead of dropping candidates.
    • Contract-compatible with DSATopKBackend.topk_func: (score fp32 (B, L), lengths i32, topk, row_starts)(B, topk) i32 local indices, front-packed, −1 padded; exact -inf exclusion matching the torch backend.
    • topk ≤ 4096 (production DSA topk = 2048), any row width < 2^24, SM90+, CUDA-graph-safe after one eager warmup per config (replay verified bitwise-equal to eager).
    • Occupancy-aware survivor-stage sizing: small grids (B ≤ #SMs, i.e. decode) take a 20480-entry smem stage for free (1 CTA/SM either way; keeps 64k-wide rows off the fallback, 3.1× faster there); full grids keep the 4096-entry stage and 4 CTA/SM.
  • python/sglang/srt/layers/attention/dsa/dsa_topk_backend.py: DSATopKBackend.CUTEDSL enum value + topk_func dispatch.
  • python/sglang/srt/server_args.py: cutedsl added to --dsa-topk-backend choices. Like the existing torch backend, cutedsl currently requires SGLANG_DSA_FUSE_TOPK=false (it implements selection, not the fused page-table transforms); help text updated accordingly.
  • test/registered/jit/test_cutedsl_topk.py (new): contract tests (tie-robust exact-set parity vs a torch.topk oracle across randn / realistic-relu / quantized-tie / -inf inputs, with and without row_starts, edge lengths 0/topk/full-width), bitwise-determinism repeat test, CUDA-graph capture/replay test, empty-batch test.

Accuracy Tests

Tie-robust checker (per row: index range/uniqueness/count, selected-score multiset bitwise-equal to the torch.topk reference — a unique answer even under exact fp32 ties, unlike index comparison): all green for the new kernel across randn, realistic-relu plateau, quantized ties, all-equal rows (L=32k), narrow-range distributions (L=32k), sprinkled/all -inf rows, 128k-wide decode rows, topk ∈ {512, 1024, 2048, 4096}, with and without row_starts.

Determinism A/B on B200 (30 repetitions per input, identical input tensor, busy sibling stream to perturb scheduling; "sets" compares row-wise sorted index sets):

input backend distinct bitwise outputs distinct index sets
realistic relu logits (exact-0.0 plateau), 256×16k fast_topk_v2 30/30 30/30
cutedsl 1 1
quantized ties, 256×16k fast_topk_v2 30/30 30/30
cutedsl 1 1
all-equal rows, 64×32k fast_topk_v2 30/30 7/30
cutedsl 1 1
narrow-range distinct values, 64×32k fast_topk_v2 30/30 30/30 (+64/64 rows fail exact-set parity — #17747 overflow skip)
cutedsl 1 1
randn (no ties), 256×16k fast_topk_v2 30/30 (order only) 1
cutedsl 1 1

Speed Tests and Profiling

B200, topk=2048, paired both orders, µs/call. Caveat: the host was shared (74–89% background GPU load from other jobs) — values are min-of-rounds as the best uncontended estimate; ratios were stable across three runs. flashinfer-det is the existing deterministic path measured as sglang dispatches it (_topk_unfused masking + flashinfer.top_k(deterministic=True)).

shape (rows×width) cutedsl fast_topk_v2 flashinfer-det vs default vs flashinfer-det
prefill 4096×8k (causal ragged) 438 165 903 0.38× 2.06×
prefill 8192×16k (causal ragged) 1195 511 2952 0.43× 2.47×
prefill 4096×32k (causal ragged) 1257 458 2384 0.36× 1.90×
prefill uniform 8192×8k 910 367 1921 0.40× 2.11×
decode 64×64k 101 34 129 0.34× 1.28×
decode 128×128k 268 57 267 0.21× 1.00×
decode 32×8k 24 10 53 0.43× 2.20×
short rows 256×1k 20 2.3 51 0.11× 2.56×
verify-like 512×64k 486 101 565 0.21× 1.16×
realistic relu prefill 8192×16k 1926 700 4148 0.36× 2.15×
realistic relu decode 64×64k 89 32 104 0.36× 1.17×

Honest summary: the default kernel remains 2.3–4.8× faster on real-work shapes (more on short rows, where it writes identity indices without reading scores) — this PR does not propose changing the default. As the deterministic backend, cutedsl is faster than the current flashinfer deterministic path on 9/11 shapes and at parity on the rest, while additionally being exact under tie/overflow regimes.

Known follow-up headroom: a 1024-thread variant (largest remaining lever for the decode gap), fp16 coarse first pass, host-launch pre-binding.

Related

#15041 (deterministic inference), #17747 / #16858 (top-k smem/perf background), flashinfer#2215 (deterministic top-k precedent).

Checklist


CI States

Latest PR Test (Base): ❌ Run #29315598173
Latest PR Test (Extra): ❌ Run #29315597939

--dsa-topk-backend cutedsl: bitwise-deterministic exact top-k selection
(canonical order: score desc, index asc on bitwise-equal scores;
smallest-index tie subset at the k-th boundary; threshold-bin overflow
falls back to full-row rescans instead of dropping candidates).
Complements the deterministic-inference effort: the default radix kernel
emits arrival-order results and can return a different index set per run
under exact score ties, which realistic DSA indexer logits (relu zero
plateau) hit routinely.

Selection-only for now: like the torch backend it requires
SGLANG_DSA_FUSE_TOPK=false.
@gemini-code-assist

Copy link
Copy Markdown
Contributor

Warning

You have reached your daily quota limit. Please wait up to 24 hours and I will start processing your requests again!

…tests

Host wrapper: replace asserts with actionable errors and enforce the
supports() envelope (topk <= 4096, L < 2**24, B*L and B*topk < 2**31) so
out-of-envelope calls fail loudly instead of silently mis-tie-breaking
(L >= 2**24 wraps nidx24) or mis-addressing; move lengths/row_starts to
score.device (int32) and validate their size, matching the torch backend;
refuse first-touch allocation of the cached default-row_starts buffer
under CUDA graph capture and make its zero-fill stream-safe; serialize
cute.compile behind a lock, key the compile cache by device, and compile
under score.device; validate SGLANG_DSA_TOPK_CUTEDSL_SURVCAP; document
NaN ordering (bitwise total order, not torch.topk NaN-maximal) and the
trusted-metadata window clamps.

Tests: verify the advertised smallest-index tie subset at the k-th
boundary (previously only order, count, and score multiset were checked);
stable data seeds (zlib.crc32 instead of PYTHONHASHSEED-randomized
hash()); skip cleanly when nvidia-cutlass-dsl is unavailable; cover an
all--inf row with a nonzero window.
@zkyue

zkyue commented Jul 14, 2026

Copy link
Copy Markdown
Author

Pushed 48bce24 — pre-review hardening pass over the wrapper and tests:

  • cute_topk_func now enforces its documented envelope with actionable errors (CUDA fp32 input, topk <= 4096, L < 2^24, B*L/B*topk < 2^31) instead of asserts; out-of-envelope shapes previously failed late or, for L >= 2^24, could silently mis-tie-break at the k-th boundary.
  • lengths/row_starts are moved to score.device (int32) and size-validated, matching the torch backend.
  • CUDA-graph hardening: the cached default-row_starts buffer refuses first-touch allocation under graph capture (same eager-warmup contract as compilation) and its zero-fill is stream-safe; the compile cache is device-keyed and compilation is serialized and pinned to score.device.
  • SGLANG_DSA_TOPK_CUTEDSL_SURVCAP is validated; NaN ordering is now documented (deterministic bitwise total order — differs from torch.topk NaN-maximal; DSA indexer logits are NaN-free).
  • Tests: the advertised smallest-index tie subset at the k-th boundary is now asserted directly (previously only count/order/score-multiset); data seeds are stable across processes; clean skip when nvidia-cutlass-dsl is unavailable; added an all--inf row with a nonzero window.

All 12 registered tests pass on B200, and the determinism A/B harness stays bitwise-stable across runs on tie-heavy inputs.

@DarkSharpness

Copy link
Copy Markdown
Collaborator

Intersting observation. I never expect there will be that many 0 in the indexer output (this is a degeneration case since we cannot distuiguish 0 values).

@zkyue

zkyue commented Jul 14, 2026

Copy link
Copy Markdown
Author

Right — and the zeros are structural rather than a fixture artifact: the DSA lightning-indexer score is a sum of ReLU-ed per-head dot products, so every non-positive contribution lands on exactly 0.0. In realistic captures a sizable fraction of each row sits on that 0.0 plateau, and whenever a row's k-th boundary falls inside the plateau the tie set is huge — that is exactly where the default kernel's selection varies run-to-run (and where the #17747 overflow-drop bites hardest). Bitwise determinism therefore needs a total order among equal scores; this backend breaks ties by smallest index (a second radix pass over ~idx), which is what keeps replays byte-identical. Happy to add a short note about the tie/0.0-plateau semantics to the backend docs if that would help.

@DarkSharpness

DarkSharpness commented Jul 14, 2026

Copy link
Copy Markdown
Collaborator

@zkyue Could you please share some realistic data on this (e.g. score distribution for the indexer output in online serving)? I wonder whether our topk_v2 kernel (JIT kernel) is safe under this setting.

@zkyue

zkyue commented Jul 14, 2026

Copy link
Copy Markdown
Author

Working on exactly that — I will follow up here with per-row score-distribution stats (fraction of exact-0.0 entries, tie-set size at the k-th boundary) from a reproducible setup, plus a small harness so you can measure the same on any live deployment, and a focused check of how topk_v2 behaves as the boundary tie-set grows. Should have numbers up shortly.

@zkyue

zkyue commented Jul 14, 2026

Copy link
Copy Markdown
Author

Thanks for taking a look — here is the data we promised. Short answer to "is topk_v2 safe under this setting": in our tests it is memory-safe on arbitrarily large tie sets (no crashes, no illegal addresses, no -1 short-fills, no duplicate/out-of-range indices), but (a) the selected set changes run to run whenever the k-th boundary falls inside a tie plateau, (b) on ReLU-shaped rows we observed rare drops of strictly-positive candidates (1/256 rows in one config), and (c) the narrow-value-range regime of #17747 still drops candidates systematically. Details and reproduction below.

Provenance, first. We cannot share score captures from our internal serving deployments, so everything here is synthetic and reproducible: the "realistic relu logits" used in this PR's tests are relu(randn - 1.28) (generator is in test_cutedsl_topk.py), and the distribution study below sweeps the plateau fraction as a parameter instead of claiming one number. To get real "online serving" numbers — which depend on model and traffic — the harness at the bottom runs on scores captured from any live deployment with a 2-line hook; we'd genuinely be interested in what it reports on your traffic.

Why exact-0.0 plateaus are structural. In the DSA/lightning-indexer scoring as implemented in sglang (fp8_index in dsa/tilelang_kernel.py): score = k_s * sum_h w_h * relu(q_h . k) with nonnegative head weights w_h and an e8m0 (power-of-two) k_s, so score == 0.0 exactly iff all head dots are <= 0, and the scale multiply preserves the exact zero. The architecture guarantees the plateau exists; only its size is model/traffic dependent.

When does the plateau reach the top-k boundary? It is a threshold phenomenon. The k-th boundary sits inside the 0.0 plateau iff nnz(row) < k, i.e. frac_zero > 1 - k/L; otherwise the boundary tie set is ~1. For k=2048 that means: L=4096 needs >50% zeros, L=8192 >75%, L=32768 >93.75%, L=131072 >98.4%. So mid-length rows (a large share of typical online traffic) hit it with moderate sparsity, while long rows need extreme sparsity. Measured on the synthetic indexer model (64 heads, per-head logits equicorrelated Gaussians with correlation rho and mean mu; sweep in the harness):

L frac(score==0.0) boundary tie-set size (k=512 / 1024 / 2048) rows straddling
2048 0.61 1 / ~1247 / (L<=k) 64/64 at k=1024
2048 0.78 ~1599 / ~1599 / (L<=k) 64/64
2048 0.91 ~1859 / ~1859 / (L<=k) 64/64
8192 0.78 1 / 1 / ~6382 64/64 at k=2048
8192 0.91 1 / ~7416 / ~7416 64/64
32768 0.61–0.91 1 / 1 / 1 0
131072 0.61–0.91 1 / 1 / 1 0

Either the tie set at the boundary is 1, or it is thousands of tokens — nothing in between. When it is thousands, which zeros pad the tail of the top-k is underdetermined, and those tokens still participate in sparse attention (an indexer score of exactly 0.0 does not mean zero attention weight), so the model output can differ between identical runs.

topk_v2 behavior vs boundary tie-set size (sgl_kernel.fast_topk_v2 v0.4.4, B200, torch 2.11/CUDA 13.3; k=2048 — the only k the kernel supports; 25 replays per config with a concurrent noise stream; "wrong score-set" = selected-score multiset differs from torch.topk's, i.e. a candidate at/above the true boundary was dropped for a strictly-worse one):

input (fp32) B x L boundary tie set distinct outputs distinct index SETS -1 / dup / OOB wrong score-set rows
all exactly 0.0 64 x 32768 32768 25/25 21/25 0 0
0.0 plateau + 256 distinct positives 64 x 32768 32512 25/25 25/25 0 0
0.0 plateau + 1024 distinct positives 64 x 32768 31744 25/25 25/25 0 0
same, L=131072 (B=8) 8 x 131072 130048 25/25 25/25 0 0
tie at 1.0 straddling boundary, T=64…16384 64 x 32768 64…16384 25/25 25/25 0 0
relu(randn-1.28)*11.3 (nnz ~1640 < k) 256 x 16384 ~14740 @ 0.0 25/25 25/25 0 1/256
relu(randn-1.28)*11.3 (nnz ~3280 > k) 64 x 32768 1 25/25 1/25 0 0
1.0 + 1e-7*randn (narrow distinct, #17747) 64 x 32768 1 25/25 25/25 0 64/64
same, L=131072 8 x 131072 1 25/25 25/25 0 8/8
control: boundary among well-separated values 64 x 32768 1 25/25 1/25 0 0

Reading of this: on pure exact-tie plateaus the kernel is well-behaved in the correctness sense — every run returns a valid top-k (dropped candidates are ties, so the score multiset stays right) — and the failure is reproducibility only: a different token set nearly every run once the boundary is inside the plateau. The realistic ReLU shape (plateau plus a spray of small positives) additionally produced a rare row where a strictly-positive candidate was replaced by a 0.0 one, and the narrow-range regime from #17747 drops true candidates on every row. No illegal behavior anywhere.

The cutedsl backend in this PR does not change which regime traffic lands in — it only makes the tie-break defined (score desc, index asc: the smallest-index subset of the boundary tie set), so the selected set is a pure function of the input, run to run and across batch compositions.

Both scripts are below; the first one is the interesting one to run on real serving traffic — capture the score argument of DSATopKBackend.topk_func (e.g. torch.save(score.float().cpu(), f"/tmp/idx_scores_{time.time_ns()}.pt") right before the topk call) and feed it --input. If its boundary==0.0 rows / straddle counts are nonzero on your deployment, the run-to-run selection variance above applies there.

dsa_topk_score_stats.py — per-row tie/plateau stats for captured or synthetic indexer scores
# dsa_topk_score_stats.py — per-row tie/plateau statistics for DSA indexer scores.
#
# Answers: "how big is the tie set at the top-k boundary of the scores that
# feed fast_topk_v2 / the DSA topk backend?"
#
# Two modes:
#   1) Real capture:   python dsa_topk_score_stats.py --input scores.pt [--lengths lengths.pt]
#      scores.pt = torch.save'd fp32 tensor (B, L), the exact `score` argument of
#      DSATopKBackend.topk_func. Capture in a live deployment with a 2-line hook, e.g.
#      in dsa_topk_backend.py right before the topk call:
#          torch.save(score.float().cpu(), f"/tmp/idx_scores_{time.time_ns()}.pt")
#   2) Synthetic model of the indexer (no weights needed):
#      python dsa_topk_score_stats.py --synthetic
#      score[s] = ks[s] * sum_h w_h * relu(x[h,s]) with x[h,s] = per-head logit,
#      modeled as correlated Gaussians: x[h,s] = mu + sqrt(rho)*c[s] + sqrt(1-rho)*z[h,s].
#      rho = cross-head correlation of "is this key relevant", mu = mean logit.
#      Exact-0.0 plateau fraction = P(all H head-logits <= 0); rho/mu sweep spans
#      "no plateau" -> "row is mostly exact 0.0" so no single number is assumed.
import argparse

import torch


def row_stats(row: torch.Tensor, topks=(512, 1024, 2048)) -> dict:
    """row: 1-D fp32 scores (only the valid window). Returns tie stats per topk."""
    n = row.numel()
    out = {"len": n, "frac_zero": (row == 0.0).float().mean().item()}
    vals = torch.sort(row, descending=True).values
    for k in topks:
        if n <= k:
            out[k] = dict(boundary=None, tie=0, above=n, straddle=False, zero_bound=False)
            continue
        v = vals[k - 1].item()  # k-th largest = boundary score
        tie = (row == v).sum().item()  # tie set containing the boundary
        above = (row > v).sum().item()  # rows strictly above: always selected
        out[k] = dict(
            boundary=v,
            tie=tie,
            above=above,
            straddle=(tie > 1 and above + tie > k),  # selection within tie set is underdetermined
            zero_bound=(v == 0.0),
        )
    return out


def synthetic_rows(L, H=64, rho=0.9, mu=-1.0, n_rows=64, seed=0, device="cpu"):
    """Documented synthetic indexer: per-head logits are equicorrelated Gaussians."""
    g = torch.Generator(device=device).manual_seed(seed)
    c = torch.randn(n_rows, 1, L, device=device, generator=g)  # shared relevance factor
    z = torch.randn(n_rows, H, L, device=device, generator=g)  # per-head noise
    x = mu + (rho**0.5) * c + ((1 - rho) ** 0.5) * z
    w = torch.rand(n_rows, H, 1, device=device, generator=g) / H  # nonneg head weights
    score = (w * torch.relu(x)).sum(dim=1)  # (n_rows, L)
    ks = torch.exp2(torch.randint(-2, 3, (1, L), device=device, generator=g).float())
    return score * ks  # e8m0 per-key scale: preserves exact 0.0


def report(rows_iter, topks):
    import statistics as st

    agg = {k: {"tie": [], "straddle": 0, "zero_bound": 0, "n": 0} for k in topks}
    fz = []
    for row in rows_iter:
        s = row_stats(row.float(), topks)
        fz.append(s["frac_zero"])
        for k in topks:
            if s[k]["boundary"] is None:
                continue
            a = agg[k]
            a["tie"].append(s[k]["tie"])
            a["straddle"] += s[k]["straddle"]
            a["zero_bound"] += s[k]["zero_bound"]
            a["n"] += 1
    print(f"  frac(score==0.0): mean {st.mean(fz):.3f}  min {min(fz):.3f}  max {max(fz):.3f}")
    for k in topks:
        a = agg[k]
        if not a["n"]:
            print(f"  topk={k}: all rows shorter than k")
            continue
        t = sorted(a["tie"])
        print(
            f"  topk={k}: boundary-tie-set size med {t[len(t)//2]} p90 {t[int(len(t)*0.9)]} "
            f"max {t[-1]} | straddling rows {a['straddle']}/{a['n']} "
            f"| boundary==0.0 rows {a['zero_bound']}/{a['n']}"
        )


if __name__ == "__main__":
    p = argparse.ArgumentParser()
    p.add_argument("--input", help=".pt tensor (B,L) of captured indexer scores")
    p.add_argument("--lengths", help="optional .pt int tensor (B) of valid lengths")
    p.add_argument("--synthetic", action="store_true")
    p.add_argument("--topks", default="512,1024,2048")
    args = p.parse_args()
    topks = tuple(int(x) for x in args.topks.split(","))
    if args.input:
        sc = torch.load(args.input, map_location="cpu").float()
        ln = torch.load(args.lengths, map_location="cpu") if args.lengths else None
        rows = [sc[i, : (ln[i].item() if ln is not None else sc.shape[1])] for i in range(sc.shape[0])]
        print(f"captured scores: {sc.shape[0]} rows")
        report(rows, topks)
    elif args.synthetic:
        for L in (2048, 8192, 32768, 131072):
            for rho, mu in ((0.0, -1.0), (0.5, -1.0), (0.9, -1.0), (0.9, -2.0), (0.99, -1.0)):
                n_rows = max(8, min(64, 2**22 // L))
                print(f"L={L} H=64 rho={rho} mu={mu} ({n_rows} rows)")
                report(synthetic_rows(L, rho=rho, mu=mu, n_rows=n_rows), topks)
    else:
        p.error("need --input or --synthetic")
topk_v2_tie_safety.py — the replay/tie-sweep used for the table above
# topk_v2_tie_safety.py — sgl_kernel.fast_topk_v2 behavior vs boundary tie-set size.
#
# For each synthetic config we build B identical-distribution rows where the
# tie set containing the k-th boundary has a controlled size, replay the
# kernel N times (with a noisy sibling stream perturbing SM scheduling), and
# check per rep:
#   - bitwise output hash + per-batch sorted-index-set hash (selection stability)
#   - contract: -1 short-fill count (rows have >= k valid candidates, so any -1
#     is a dropped candidate), duplicate indices, out-of-range indices
#   - selected-score multiset vs torch.topk (tie-tolerant correctness):
#     mismatch = candidates at/above the true boundary were dropped and
#     replaced by strictly-worse ones (issue #17747 mode)
#
# Families:
#   plateau:  m distinct positive scores, rest exactly 0.0 (ReLU indexer shape).
#             If m < k the boundary sits in the zero plateau: tie = L - m.
#   tie@v:    T scores exactly v=1.0 straddling the boundary, distinct values
#             above and below (generic quantization-tie shape).
import argparse
import hashlib
import sys

import torch
from sgl_kernel import fast_topk_v2


def make_plateau(B, L, m, device, seed):
    g = torch.Generator(device=device).manual_seed(seed)
    sc = torch.zeros(B, L, device=device)
    if m:
        vals = 0.5 + torch.arange(m, device=device) * 1e-3  # distinct positives
        for b in range(B):
            pos = torch.randperm(L, device=device, generator=g)[:m]
            sc[b, pos] = vals[torch.randperm(m, device=device, generator=g)]
    return sc


def make_tie(B, L, k, T, device, seed):
    g = torch.Generator(device=device).manual_seed(seed)
    above = max(0, k - T // 2)  # boundary lands inside the tie set
    sc = torch.empty(B, L, device=device)
    for b in range(B):
        pos = torch.randperm(L, device=device, generator=g)
        sc[b, pos[:above]] = 2.0 + torch.arange(above, device=device) * 1e-4
        sc[b, pos[above : above + T]] = 1.0
        nb = L - above - T
        sc[b, pos[above + T :]] = 0.01 + torch.arange(nb, device=device) * 1e-7
    return sc


def make_relu_pr(B, L, device, seed):
    """The (public) PR fixture: relu of shifted gaussian, ~10% nonzero.
    At L=16384 nonzeros ~1638 < 2048 so the boundary sits in the 0.0 plateau."""
    g = torch.Generator(device=device).manual_seed(seed)
    return torch.relu(torch.randn(B, L, device=device, generator=g) - 1.28) * 11.3


def make_narrow(B, L, device, seed):
    """Distinct values in a narrow range (all within ~1e-6 of 1.0): the radix
    threshold-bin staging-cap overflow regime (issue #17747)."""
    g = torch.Generator(device=device).manual_seed(seed)
    return 1.0 + torch.randn(B, L, device=device, generator=g) * 1e-7


def h(t):
    return hashlib.sha1(t.cpu().numpy().tobytes()).hexdigest()[:12]


def check(sc, idx, k):
    """Returns (short_fill_rows, dup_rows, oob_rows, wrong_set_rows)."""
    B, L = sc.shape
    short = dup = oob = wrong = 0
    ref = torch.sort(sc, dim=-1, descending=True).values[:, :k]
    for b in range(B):
        row = idx[b]
        valid = row[row >= 0]
        if (row >= L).any() or (row < -1).any():
            oob += 1
            continue
        if valid.numel() < k:
            short += 1
        if valid.numel() != valid.unique().numel():
            dup += 1
        got = torch.sort(sc[b, valid.long()], descending=True).values
        if got.numel() != k or not torch.equal(got, ref[b]):
            wrong += 1
    return short, dup, oob, wrong


def run(name, sc, k, reps, noise_mat, results):
    B, L = sc.shape
    ln = torch.full((B,), L, dtype=torch.int32, device=sc.device)
    ns = torch.cuda.Stream()
    outs, sets = set(), set()
    worst = (0, 0, 0, 0)
    err = ""
    try:
        for r in range(reps):
            with torch.cuda.stream(ns):
                for _ in range(4):
                    noise_mat @ noise_mat
            idx = fast_topk_v2(sc, ln, k)
            torch.cuda.synchronize()
            outs.add(h(idx))
            sets.add(h(torch.sort(idx, dim=-1).values))
            worst = tuple(max(a, b) for a, b in zip(worst, check(sc, idx, k)))
    except RuntimeError as e:
        torch.cuda.synchronize()
        err = str(e).splitlines()[0][:60]
    results.append((name, B, L, k, len(outs), len(sets), *worst, err))
    print(results[-1], flush=True)


if __name__ == "__main__":
    p = argparse.ArgumentParser()
    p.add_argument("--gpu", type=int, default=0)
    p.add_argument("--reps", type=int, default=25)
    args = p.parse_args()
    torch.cuda.set_device(args.gpu)
    dev = "cuda"
    noise_mat = torch.randn(2048, 2048, device=dev)
    results = []
    print("name B L k distinct_out distinct_set short dup oob wrong_set err", flush=True)
    # fast_topk_v2 supports only topk=2048 (asserted in sgl_kernel wrapper;
    # kernel is specialized for the DSv3.2 production config)
    for B, L in ((64, 32768), (4, 32768), (8, 131072)):
        for k in (2048,):
            for m in (0, 256, 1024, 4095, 8192):
                tie = L - m if m < k else 1
                run(f"plateau_m{m}_tie{tie}", make_plateau(B, L, m, dev, 7), k, args.reps, noise_mat, results)
            for T in (64, 512, 4096, 16384):
                run(f"tie1.0_T{T}", make_tie(B, L, k, T, dev, 11), k, args.reps, noise_mat, results)
    # PR-public relu fixture: L=16384 -> ~1638 nonzeros < 2048, boundary in plateau;
    # L=32768 -> ~3277 nonzeros > 2048, boundary among distinct nonzeros (control)
    for B, L in ((64, 16384), (256, 16384), (64, 32768)):
        run(f"relu_pr_L{L}", make_relu_pr(B, L, dev, 13), 2048, args.reps, noise_mat, results)
    # narrow-range distinct values: #17747 staging-cap overflow regime
    for B, L in ((64, 32768), (8, 131072)):
        run(f"narrow_L{L}", make_narrow(B, L, dev, 17), 2048, args.reps, noise_mat, results)
    import csv

    with open("topk_v2_tie_safety.csv", "w") as f:
        w = csv.writer(f)
        w.writerow(["name", "B", "L", "k", "distinct_out", "distinct_set", "short", "dup", "oob", "wrong_set", "err"])
        w.writerows(results)
    print("done", len(results), "configs -> topk_v2_tie_safety.csv")

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants