Skip to content

[AMD] Enable GLM DSA prefill top-k to the v2 kernel - #37889

Open
EricKing626 wants to merge 16 commits into
sgl-project:mainfrom
EricKing626:amd/topk-v2-prefill
Open

EricKing626 wants to merge 16 commits into
sgl-project:mainfrom
EricKing626:amd/topk-v2-prefill

Conversation

@EricKing626

@EricKing626 EricKing626 commented Sep 4, 2026

Copy link
Copy Markdown
Contributor

Motivation

Extend the DSA top-k v2 kernel to packed score rows and route GLM-5.x prefill through it. The prefill top-k kernel drops ~73%; at ISL 70000 / OSL 300 that is +4.9% token throughput per GPU, −3.5% median TPOT and −4.8% median TTFT (geomean over concurrency 4-64); GSM8k 0.927.

image

PR #36684 and PR #36851 turned the v2 fused top-k on for GLM-5.x on ROCm, but only decode ever reached it. Prefill still ran the legacy paged transform, which gathers the wide page_size=1 table.

The blocker was addressing, not math. The v2 kernel assumed each row starts at column 0 of its own score row and owns one page-table row. DSA extend breaks both: all requests share one packed score buffer, so a row starts at some column in the middle, and every row of a request reads the same page-table row.

Two small lookup tables fix that -- one for where a row starts, one for which page-table row it uses. No copy of the score buffer, and no per-row page table, which at 327K context would be hundreds of MB per forward.

Modifications

topk_v2.cuh, dsv4/topk.py -- add two optional (rows,) int32 indirections to topk_transform_512_v2:

omitted (decode layout) supplied (packed extend)
score window of row i scores[i, :seq_lens[i]] scores[i, row_starts[i] : row_starts[i] + seq_lens[i]]
page-table row page_tables[i] page_tables[row_to_batch[i]]

Both default to null and give exactly the old addressing, so nothing changes for existing callers. Shifting the score pointer keeps the selected index row-local, which is what the transform already expects, so the output path is untouched. The page-table row count has to match the score row count only when row_to_batch is absent; with it, the caller owns that bound.

dsa/dsa_topk_backend.py -- route packed PAGED extend to _topk_transform_v2_paged, passing ks as row_starts and token_to_batch_idx as row_to_batch. GLM-5.x already qualifies as PAGED for EXTEND, and the plan needs no new work: it is already built once per forward over dsa_seqlens_expanded, whose row count is what v2 sees.

Two shapes still take the legacy path. The dispatch checks for them itself, so they fall back instead of tripping the helper's asserts:

  • Chunked-prefix-cache extend, where the plan covers the whole forward but each call sees only one chunk. Not to be confused with scheduler-level chunked prefill: under --chunked-prefill-size each chunk is its own forward with its own plan, and does take the v2 path.
  • Row stride not a multiple of 4, which the vectorized load needs. An extend row stride is the batch's total KV length, so it is aligned only by luck.

The decode condition is left byte-identical, because dsa_drop_wide_page_table drops the page_size=1 table on exactly that condition and the two must stay in sync.

dsa_backend.py -- comment only, recording that packed PAGED extend is now a plan consumer.

Accuracy Tests

GSM8k, 8x MI355X (gfx950), GLM-5.2-MXFP4, TP=4, non-MTP, with this PR:

$ python3 benchmark/gsm8k/bench_sglang.py --port 8000 --num-questions 1200 --parallel 1200
Accuracy: 0.927
Invalid: 0.000
Latency: 39.462 s
Output throughput: 3095.724 token/s

test/registered/kernels/ops/attention/test_topk_v2.py on gfx950: 250 passed, 3 warnings in 19.99s, including the new test_topk_v2_packed_rows cases that check the row_starts / row_to_batch layout against a row-local reference.

Speed Tests and Profiling

Baseline is #29xxx as merged (top-k v2 on, prefill still legacy); this PR is the prefill routing on top of it. Same build, same flags, non-MTP, 8x MI355X, TP4.

Kernel-level -- topk_main_kernel now replaces topk_transform_prefill_kernel on every prefill launch; the old kernel is gone from the traces entirely. ISL 70000 / OSL 300, per TP rank, matching launch grids:

conc grid.x baseline this PR Δ
4 16384 2710.2 µs/launch 721.9 µs/launch −73.4%
4 4615 (tail) 1390.3 µs/launch 359.1 µs/launch −74.2%
64 16384 2694.2 µs/launch 726.4 µs/launch −73.0%

Per rank at conc 4: 882.9 ms → 234.9 ms of top-k over the profiled window.

End-to-end -- ISL 70000 / OSL 300, --max-running-requests 8. Geomeans are over the per-config ratios.

conc Interactivity (tok/s/user) Token TPUT per GPU (tok/s) Δ TPUT Median TTFT (ms) Median TPOT (ms) Median ITL (ms)
4 29.34 → 30.35 4573.4 → 4761.6 +4.1% 3276 → 3128 34.08 → 32.95 10.63 → 10.52
8 14.13 → 15.08 4938.7 → 5139.9 +4.1% 5774 → 5646 70.77 → 66.33 12.28 → 12.24
16 12.44 → 13.08 5064.3 → 5296.6 +4.6% 27576 → 26563 80.38 → 76.43 12.21 → 11.99
32 12.94 → 13.06 5004.1 → 5301.0 +5.9% 79535 → 74147 77.25 → 76.57 12.20 → 12.30
64 12.72 → 13.00 4996.4 → 5297.1 +6.0% 181367 → 169622 78.60 → 76.91 12.11 → 12.01
geomean +3.7% +4.9% −4.8% −3.5% not claimed

Interactivity is 1000 / median TPOT; Token TPUT per GPU is (input + output) tok/s / 4 for one TP4 server. Concurrency above 8 exceeds the admission cap, so TTFT there is mostly queueing time.

Reproduce:

export SGLANG_ROCM_FUSED_DECODE_MLA=0
export ROCM_QUICK_REDUCE_QUANTIZATION=INT4   # FP4 model

python3 -m sglang.launch_server --model GLM-5.2-MXFP4 \
  --tp 4 --trust-remote-code --tool-call-parser glm47 --reasoning-parser glm45 \
  --mem-fraction-static 0.85 --kv-cache-dtype fp8_e4m3 \
  --disable-radix-cache --chunked-prefill-size 16384 \
  --dsa-prefill-backend triton --dsa-decode-backend triton \
  --enable-aiter-allreduce-fusion --tokenizer-worker-num 8 \
  --max-running-requests 8

python3 -m sglang.bench_serving --backend sglang --dataset-name random \
  --random-input-len 70000 --random-output-len 300 --random-range-ratio 0.8 \
  --max-concurrency $CONC --num-prompts $((5 * CONC))

Checklist

Review and Merge Process

  1. Ping Merge Oncalls to start the process. See the PR Merge Process.
  2. Get approvals from CODEOWNERS and other reviewers.
  3. Trigger CI tests with comments or contact authorized users to do so.
    • Common commands include /tag-and-rerun-ci, /tag-run-ci-label, /rerun-failed-ci
  4. After green CI and required approvals, ask Merge Oncalls or people with Write permission to merge the PR.

CI States

Latest PR Test (Base): 🚫 Run #34846434360
Latest PR Test (Extra): ❌ Run #34846433791
Latest PR Test (AMD ROCm 10): ❌ Run #34846434075

The v2 fused top-k + page-table transform assumes the decode layout: row i
selects over scores[i, :seq_lens[i]] and maps through page_tables[i]. DSA
extend cannot describe itself that way -- every request's scores are packed
into one row-major buffer, so a row's window starts at a batch-global column
offset, and a request contributes many rows that all share its single
page-table row.

Add two optional per-row indirections, row_starts (score column offset) and
row_to_batch (page-table row), so that layout needs neither a row-local copy
of the score buffer nor a per-row expansion of the page table -- at 327K
context the latter would be hundreds of MB per forward. Offsetting the score
pointer also keeps the selected index row-local, which is what the transform
already expects, so the emit path is untouched.

Both default to null and reproduce the previous addressing exactly, so this
commit is a capability addition with no behavior change. The page-table row
count is only required to match the score row count when the mapping is
absent; with it the caller owns that bound, which is the one invariant the
kernel cannot check.
GLM-5.2 runs prefill with a dsa_prefill_backend outside the flashmla_sparse
family (tilelang), so get_topk_transform_method already returns PAGED for
EXTEND rather than RAGGED. Despite that, prefill never reached the v2 fused
transform: the dispatch also required row_starts to be absent and the score
row count to equal the page-table row count, and extend violates both -- it
passes ks and has many rows per request. It therefore fell to the legacy
transform, which gathers the wide page_size=1 table.

Now that the kernel can address packed rows, dispatch extend to it as well,
feeding ks as row_starts and the metadata's token_to_batch_idx as
row_to_batch. The plan needs no new work: it is already built per forward
over dsa_seqlens_expanded, whose row count is exactly what v2 sees.

Two shapes deliberately stay on the legacy path. Chunked extend, because its
plan spans the whole forward while each call sees one chunk, and any extend
whose row stride is not a multiple of 4, because the kernel's vectorized load
requires that and an extend row stride is the batch's total KV length -- only
aligned by luck. Both are checked in the dispatch rather than left to the
helper's assertions, so they fall back instead of raising.

The decode condition is left byte-identical: dsa_drop_wide_page_table drops
the page_size=1 table for exactly that condition, and the two must stay in
sync or the legacy transform would read a table that no longer exists.
1am9trash and others added 4 commits September 4, 2026 17:01
A window start is a running KV length, so it lands on a 16-byte boundary only
by luck, and the previous commit offset the score pointer by it directly. The
vectorized load then faults: CI aborted on the ragged case [1, 13, 2], whose
starts are 0 / 1537 / 3086.

Do what topk_ragged_kernel already does for the same reason. Round the read
base down to the vector boundary, mask the <= 3 columns that pulls in ahead of
the window (they belong to the preceding request, so they are finite scores
that would otherwise win the selection), widen seq_len by the residue, and
subtract it back at the page lookup. The mask lands after the PDL wait, or the
indexer overwrites it.

The residue is a read-window artifact, so every decision stays on the row's
real length: the trivial path takes the un-rounded problem and reads no scores
at all, and the cluster routing threshold is compared against the same length
the plan used. The widened seq_len cannot leave the dispatch level's bound
because the window end is unchanged and the level comes from the score column
count.

The cluster path is excluded on the host whenever row_starts is set: there one
row is split across the blocks of a cluster, so the head mask would need a
cluster-wide barrier to be visible, whereas every other path has one block per
row and publishes it with the __syncthreads() that already opens forward().

Zero residue reproduces the previous addressing exactly, so nothing changes for
callers that do not pass row_starts.
The packed-row extension of the paged top-k v2 transform (`row_starts` /
`row_to_batch`, plus the 16-byte read-base alignment it needs) only exists
because ROCm has no other route: on CUDA, DSA extend reaches the fused
transform through `topk_transform_ragged`, while on ROCm `get_topk_transform_
method` returns PAGED for EXTEND and the ragged kernel is unreachable.

Compile the extension out on non-ROCm builds so the CUDA paths are unchanged:

- `TopKPagedParams`: `scores` goes back to `const float*`, and `row_starts` /
  `row_to_batch` (with `head_residue` / `mask_head`) exist only under
  `USE_ROCM`. `problem()` builds the original problem and re-points it only on
  ROCm when packed rows are supplied.
- `TopKProblem::index_shift` and its use in `transform_output` are `USE_ROCM`
  only; `transform_output` keeps its original form otherwise.
- `topk_main_kernel`: the residue / head-mask block is `USE_ROCM` only; the
  CUDA build keeps the plain `seq_len <= topk` trivial check and the cluster
  dispatch loses the `row_starts == nullptr` exclusion, which is now dead
  there.
- Host `transform_paged` rejects `row_starts` / `row_to_batch` with a
  `RuntimeCheck` on non-ROCm instead of wiring them up.

Python side, the `is_xpu()` branch of `topk_transform_paged_v2` is restored to
its original body; the packed-row precondition moves out into its own check
that asserts `is_hip()` only when packed rows are actually passed. The PAGED
extend dispatch in `dsa_topk_backend` is gated on `is_hip()`, and
`test_topk_v2_packed_rows` is skipped off ROCm.

No functional change on ROCm.
row_starts shifts the read window and records the correction in
index_shift, but index_shift is only applied by transform_output. In
INDICES mode (page_table absent) emit() writes the raw index, so the
selected positions come back off by up to kVecSize-1 with no diagnostic.

No caller needs that combination, so reject it on the host instead of
leaving it unguarded.
mask_head writes the pulled-in residue columns back into scores. If the
caller passes a view whose row stride is smaller than its row width, one
row's head mask lands in the previous row's tail and corrupts scores the
kernel has not read yet. Assert stride >= width on the packed path.
topk_transform_paged_v2 masks the head-alignment residue columns to -inf
inside the score tensor, but the docstring described the call as
read-only. Say so explicitly, along with the two constraints it implies:
scores cannot be reused afterwards, and its rows must not overlap.
@EricKing626 EricKing626 changed the title [AMD] Route GLM DSA prefill top-k to the v2 kernel [AMD] Enable GLM DSA prefill top-k to the v2 kernel Sep 9, 2026
@KingRei

KingRei commented Sep 12, 2026

Copy link
Copy Markdown
Contributor

/rerun-failed-ci

1 similar comment
@EricKing626

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

The packed-row path rounds each row's read base down to a 16-byte
boundary, so a selected position is up to kVecSize-1 too large. That
correction lived in its own TopKProblem field applied inside
transform_output, which forced a USE_ROCM branch into a code path CUDA
also compiles.

bias already exists for the same shape of correction: ragged mode adds
the row's offset into the flattened output there. The two never coexist
- paged callers have no output offset, and ragged never rounds down - so
bias = -residue carries the round-down instead. The correction now
happens once in emit(), which both output modes share, and
transform_output goes back to its single unbranched form.
Conflicts came from upstream's new DUAL_OUTPUT mode (raw_indices),
which lands on the same lines as the packed-row (row_starts) support.

- params: keep both head_residue/mask_head and get_raw_output_ptr
- topk_main_kernel: keep the ROCm residue prologue, and pass
  get_raw_output_ptr(blockIdx.x) to trivial_transform on both branches
- transform_paged: signature takes raw_indices, then row_starts and
  row_to_batch, so existing positional callers are unaffected
- reject row_starts + raw_indices (host-side, C++ and Python): the raw
  output is written straight from the index register and never goes
  through emit, so it would miss the residue correction that bias carries

@DarkSharpness DarkSharpness left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Hi. Could you please temporarily hold on until #38798 ? We have some critical update on v2 algorithm, which introduce many new mechanisms for that (like padding).

Also, some comments on this PR: please try to avoid adding too many comments in the code (especially for parts that remain unchanged). Please keep the comments brief within 1 ~ 2 lines

pre-commit reflows this call onto one line (114 cols, under the limit).
Per review: keep added comments brief and drop the ones attached to
code this PR does not change (the cluster-path note under #ifndef
USE_ROCM, and a stale duplicate above the page_table RuntimeCheck).
Restore the original bias comment; the field's behaviour is unchanged.
@EricKing626

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

1 similar comment
@EricKing626

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

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.

4 participants