Skip to content

[Perf][SM70][DeepSeek-V4] Long-prefill indexer and sparse attention: vectorized epilogue, fused softmax, Turing scorer - #1280

Open
Peuqui wants to merge 20 commits into
1CatAI:mainfrom
Peuqui:pr-dsv4-prefill-stacked
Open

Peuqui wants to merge 20 commits into
1CatAI:mainfrom
Peuqui:pr-dsv4-prefill-stacked

Conversation

@Peuqui

@Peuqui Peuqui commented Oct 27, 2026 •

Copy link
Copy Markdown
Contributor

Update 2026-10-28: rebased onto main c59b3c4, which contains #1222 and #1227; their commits are no longer part of this branch, which now holds 20 commits (the 14 of #1279 and 6 new). The six are unchanged apart from the regenerated reference (git range-diff), with one addition: the fused-softmax commit now also updates the two DeepSeek-V4.1 callers of attend_gathered_keys in vllm/models/deepseek_v41/attention_ops.py (from #1277), which still passed the removed logits buffer and relied on the function to zero unused keys; the CI mypy run on the earlier push caught it. Against main's function on random data with NaN in the unused slots, 99.69-100 % of their outputs are bit-identical and the rest differ by one fp16 step, on a Tesla V100 and a Quadro RTX 8000, and the DeepSeek-V4.1 test files give the same results as on main. Test Plan files on the head before that addition (6dc35fd), Quadro RTX 8000 / Tesla V100, all passing: test_deepseek_v4_sm70_indexer.py 31 / 30 (1 skipped), test_deepseek_v4_sm70_routes.py 30 / 30, test_cudagraph_dispatch.py 51 / 51, test_dsv4_sm70_sparse_bmm.py 12 skipped / 12 passed, test_dsv4_sparse_prefill_bmm.py 17 / 17, test_deepseek_v4_sm70_qnorm_rope_kv_insert.py 4 / 4, test_compressor_kv_cache.py 32 / 32 (6 skipped each), test_sparse_indexer_decode_seq_lens.py 11 / 11. pre-commit run --from-ref origin/main --to-ref HEAD on this head: every hook passes except check-layering, which now reports vllm/config/speculative.py (model coupling 142 -> 144) next to vllm/model_executor/layers/fused_moe/runner/moe_runner.py (platform coupling 37 -> 39); main c59b3c4 reports the same two counts, so this PR adds none. pre-commit run mypy-3.10 --hook-stage manual: passed.

Purpose

Long DeepSeek-V4-Flash prefills on SM70/SM75 pipeline stages spend much of their non-MoE time in memory-bound glue around the lightning indexer and the sparse MLA attention. In a per-stage nsys profile of a 150k-token prefill (PP=5 over three Tesla V100 and two Quadro RTX 8000), the indexer's relu/head-sum epilogue took 14.0 s and 17.6 s of the two RTX stages, the pipeline's slowest, and six elementwise torch ops took 820 of 1434 us of each sparse attention pass on the RTX 8000.

Stacked PR. The branch sits on #1279 (TP prefill, itself on #1278 and #1216). #1222 (cuBLAS decode indexer above 131,072 tokens, same indexer file), on which it was built, is merged. The copy of the #1279 stack is identical to the open PR. Only the last six commits are new here, so please review those (git log 3f48315..32e4372, where 3f48315 is the head of #1279). The Turing indexer kernel uses the m16n8k8 fragment helpers of #1278 and shares the sm_75 source list in CMakeLists.txt and neighbouring lines of csrc/ops.h and csrc/torch_bindings.cpp with #1279.

The six commits make five changes (commit 2 has its test update in a commit of its own):

  1. Vectorized indexer prefill epilogue (77fd5d1). BLOCK_K 512 over 2 warps reads eight fp16 scores (16 bytes) per thread instead of two. On the [508 tokens x 64 heads, 1024 keys] score tile it reads at the cards' plain read rate of that tile. The head sum stays fp32; its reduction order changes in the last bits (up to 1.5e-5 on unit-scale scores), so the top-k can differ only at near ties. Prompts up to index_topk * 4 tokens select every key and are unaffected.
  2. Fused sparse MLA softmax (459edf6; test update 1ba5d92). The gather writes zeros for unused slots and never reads their indices. One Triton kernel reads the fp32 scores once and writes fp16 probabilities: scale, mask, sink column and the fp32 softmax with libdevice exp, the same expf torch's softmax uses. The fp32 logits and probability buffers go away (140 -> 110 MiB per 128-token pass at 640 keys). 99.2-99.8 % of the outputs are bit-identical, the rest differ by at most 3e-5 (reduction order); the error against float64 is unchanged.
  3. Fused lightning-indexer prefill scores on Turing tensor cores (37febaf). indexer_prefill_scores_sm75 keeps one token's queries for 32 heads as m16n8k8 A fragments in registers, stages e4m3 key tiles widened to fp16 in shared memory, rounds the scores to fp16 as the cuBLAS route stores them, and keeps relu, head weights and the head sum in registers; only the [tokens, keys] result reaches memory.
    • Taken when the indexer runs its relu route on an SM75 device (the capability of the query's own device) with 32 or 64 heads. Volta and other head counts keep the cuBLAS route plus the epilogue of commit 1, or the per-head Triton kernel.
    • kernel_config.sm70_sparse.indexer_prefill_cublas keeps its meaning (cuBLAS or the per-head Triton kernel) for the chunks that do not take the fused kernel; its docstring now says so.
    • The source builds for sm_75 only. Each prefill scoring route logs its choice once per device ("SM70 indexer prefill: ... on cuda:N").
  4. Indexer logits freed per prefill sub-chunk (8477763). The loop kept the previous sub-chunk's scores alive until the next scorer had allocated its own, so the peak under VLLM_SPARSE_INDEXER_MAX_LOGITS_MB was two sub-chunks while the profile run reserves one. SparseAttnIndexerKpool has the same loop and gets the same release; it was not measured.
  5. Compressed KV gather with 2048 two-warp workers per request instead of 128 (437ad09). Same kernel and arithmetic, bit-identical output. This Triton gather also serves the ROCm backend and NVIDIA GPUs without CuteDSL (dequantize_and_gather_k_cache); only the V100 and the RTX 8000 were measured.

Not a duplicate. I compared the changed files with those of all 82 open PRs (complete file lists through the API), searched the added and removed lines of every open PR for this PR's symbols (indexer_prefill_scores, _relu_weight_headsum, _EPILOGUE_BLOCK_K, sparse_prefill_bmm, dequantize_and_gather_k_cache_triton, NUM_WORKERS, VLLM_SPARSE_INDEXER_MAX_LOGITS_MB, del logits, indexer_prefill_cublas), ran a three-way merge with every PR that shares a file, and searched the open PRs for indexer, "sparse MLA", softmax, epilogue, DeepSeek-V4, DSv4, "lightning indexer" and gather. Findings:

AI assistance was used for this change. I reviewed every line and ran the tests and measurements below.

Test Plan

Native build of this branch in place (TORCH_CUDA_ARCH_LIST=7.0; indexer_prefill_scores_sm75.cu builds for sm_75 only, like the sm_75 sources of #1278 and #1279).

One pytest process per file on a Quadro RTX 8000 (SM75) and on a Tesla V100 (SM70):
python -m pytest tests/kernels/test_deepseek_v4_sm70_indexer.py
python -m pytest tests/models/test_deepseek_v4_sm70_routes.py
python -m pytest tests/kernels/attention/test_dsv4_sparse_prefill_bmm.py
python -m pytest tests/kernels/attention/test_dsv4_sm70_sparse_bmm.py
python -m pytest tests/kernels/test_deepseek_v4_sm70_qnorm_rope_kv_insert.py
python -m pytest tests/kernels/test_compressor_kv_cache.py
python -m pytest tests/v1/attention/test_sparse_indexer_decode_seq_lens.py
python -m pytest tests/v1/cudagraph/test_cudagraph_dispatch.py

pre-commit run --from-ref origin/main --to-ref HEAD; pre-commit run mypy-3.10 --hook-stage manual --from-ref origin/main --to-ref HEAD

Commit 3 in a model, A/B/A/B on an 8-layer DeepSeek-V4-Flash mini checkpoint (layers 0-7 of the checkpoint byte for byte, MTP off), PP=2 with a Tesla V100 as stage 0 and a Quadro RTX 8000 as stage 1, this branch in place with stock Triton, --max-num-batched-tokens 1024, --max-model-len 307200: one uncounted boot, then without and with the fused Turing scorer twice each. "Without" is this branch with the fused branch of sm70_indexer_prefill_logits removed, so the RTX stage takes the cuBLAS route plus the epilogue of commit 1. Per boot: prefill of 8k, 32k, 150k and 280k tokens, decode time per token, and the top-5 logprobs of 16 greedy tokens after a short and after a 4k-token prompt.

One full-model boot of this branch in place (stock Triton, no kernel_config Triton switch): DeepSeek-V4-Flash PP=5 as in the contract below, but without DSpark, since Model Runner V2 on main rejects speculative method 'dspark', and with CUDA graphs for one decode token. Per boot: prefill and decode client, eight quality questions (1500 tokens, read by hand), a long-context probe up to the entry's maximum length, and the route lines of the indexer prefill scorer in the log.

Path coverage

Path / provider Format Key conditions (shape, TP/EP, graph, capability) Expected selection / fallback Evidence (source / static / executed / measured) Unverified scope
SM70-family indexer prefill epilogue fp16 scores cuBLAS prefill scoring (Volta; Turing chunks outside commit 3) BLOCK_K 512, 2 warps executed: test_deepseek_v4_sm70_indexer.py on RTX 8000 and V100; measured: CUDA graphs, model A/B none
sparse MLA bmm attention, prefill and decode FP8 KV (DeepSeek-V4) SM70 family gather with zeroed slots + fused Triton softmax executed: test_dsv4_sparse_prefill_bmm.py, test_dsv4_sm70_sparse_bmm.py (decode reference Volta-only), test_deepseek_v4_sm70_routes.py; measured: CUDA graphs, model A/B none
indexer prefill scores on SM75 e4m3 keys, 32 or 64 heads relu route, SM75 (query's own device) indexer_prefill_scores_sm75; otherwise cuBLAS + epilogue or the per-head Triton kernel executed: test_deepseek_v4_sm70_indexer.py on RTX 8000 (fp64 reference with fp16 score rounding, 32/64 heads, partial tiles); measured: mini checkpoint A/B on the RTX stage (Test Result); executed: full-model boot of this branch, route lines on both RTX stages (Test Result) SM8x+ (never selected there); head counts other than 32 and 64 (not selected)
indexer logits per sub-chunk any VLLM_SPARSE_INDEXER_MAX_LOGITS_MB below one chunk's scores previous sub-chunk freed before the next is scored executed: mini checkpoint runs with 1, 2 and 5 sub-chunks per step (SparseAttnIndexer) SparseAttnIndexerKpool (same change, not run)
compressed KV gather FP8 KV (DeepSeek-V4) prefill; Triton gather (SM70 family, ROCm, NVIDIA without CuteDSL) 2048 workers, bit-identical executed: test_deepseek_v4_sm70_qnorm_rope_kv_insert.py, test_compressor_kv_cache.py on RTX 8000 and V100; measured: CUDA graphs, every grid compared bit for bit with the 128-worker launch ROCm and NVIDIA GPUs other than V100 and RTX 8000 (not run)

Acceleration and benchmark contract (required for performance changes)

  • Default enabled or opt-in: default enabled on the matching capability (per device); no new switch. The fused Turing scorer cannot be switched off separately.
  • Required CLI options and environment switches: none for the change. Measured entry: DeepSeek-V4-Flash (MXFP4 experts, block-FP8 linears), PP=5 over the five cards (CUDA_VISIBLE_DEVICES=0,1,4,3,2: RTX 8000 on the first and last stage), VLLM_PP_LAYER_PARTITION=10,8,8,8,9, DSpark with 5 speculative tokens, --dtype half, --kv-cache-dtype fp8, --max-num-seqs 1, --max-model-len 307200, VLLM_USE_V2_MODEL_RUNNER=1, NCCL_P2P_DISABLE=1. The two model A/B series below used --max-num-batched-tokens 512 and --num-gpu-blocks-override 3000; the full boot of this branch used --max-num-batched-tokens 1024 and --num-gpu-blocks-override 3500, without DSpark (Test Plan).
  • KV cache dtype used for the benchmark: fp8.
  • Wheel SHA or source commit:
    • Model A/B series: fork builds of github.com/Peuqui/1Cat-vLLM. Reference 4da0bcf (the fork's production at the time); changed 5669da1 (commit 1) and e1cb50f (commits 1-2). Both changed builds also carry 6c0ef8e and 0314585 (QPN8 decode split-K by occupancy, M <= 8). Net of the two, that rule only changes Turing grids whose split-16 launch spills fewer than half the SMs into a second wave (output widths of about 4,600 to 5,760); DeepSeek-V4-Flash's block-FP8 projections have widths 512, 1024, 4096, 8192 and 32768.
    • Full boot: this branch at 437ad09 in place. Mini checkpoint runs: an earlier, unpublished state of this branch that differs only in the scope of the route log lines and the wrapping of two strings.
    • Commits 3, 4 and 5 were not in either model A/B series.
  • PYTHONPATH, source overlays, or external native libraries used:
    • Model A/B series: the changed arm ran its fork source tree on PYTHONPATH with its own compile cache; both arms used the stock Triton 3.6.0 of the environment.
    • The logits measurements of commit 4 ran on the mini checkpoint with a fork tree, the patched Triton build of a later PR of this series and its tensor-core switch.
    • Mini checkpoint A/B and full boot of this branch: PYTHONPATH on this tree only, stock Triton 3.6.0, no kernel_config Triton switch.
  • User-entry route-hit, speed, and output-quality evidence: see Test Result.
  • Existing kernel/backend registry and capability-rejection reasons: selection stays in the DeepSeek-V4 SM70 indexer and sparse attention routes, keyed on the capability of the layer's own device. indexer_prefill_scores_sm75 rejects other devices, dtypes, strides and misaligned inputs with an error.
  • Per-engine configuration and deprecated-variable compatibility: no setting added or changed; the docstring of sm70_sparse.indexer_prefill_cublas now names the Turing chunks it does not decide.
  • New/changed environment metadata and generated reference check: no environment metadata. docs/design/architecture/runtime_reference.md is regenerated (one Sm70SparseConfig anchor moved by the docstring line) and checked by the check-architecture-reference hook.
  • Route snapshot command and intentional changed rows: none; the routes are chosen inside the indexer and sparse attention ops.
  • SM70 registered variables / unregistered reads / model-parameter restrictions, before -> after: no change.

Test Result

Model A/B/A/B on DeepSeek-V4-Flash PP=5 (Wheel SHA above): order one uncounted cold boot of the changed build, then reference a, changed a, reference b, changed b. Prefill time in s, lower is better (reference a / reference b -> changed a / changed b):

  • Epilogue alone (commit 1): 32.7k tokens 14.66 / 14.66 -> 14.16 / 14.15 (-3.4 %); 178k 112.79 / 112.76 -> 93.94 / 93.96 (-16.7 %); 287k 222.35 / 222.50 -> 173.07 / 173.14 (-22.2 %). Decode step unchanged (87-88 ms at 178k, 93-94 ms at 287k). Greedy output and the eight quality answers byte-identical in all five boots.
  • Epilogue and fused softmax (commits 1-2): 32.7k 14.67 / 14.66 -> 12.94 / 12.94 (-11.8 %); 178k 112.71 / 112.70 -> 82.20 / 82.14 (-27.1 %); 287k 222.26 / 222.19 -> 149.73 / 149.72 (-32.6 %). Decode step 1.9-3.4 ms shorter (178k 86.9 / 86.8 -> 84.9 / 84.5 ms, 287k 93.6 / 93.8 -> 90.4 / 90.8 ms). Texts byte-identical within each arm and different between the arms (reduction order); all read by hand and equivalent.

Commit 3, mini checkpoint A/B/A/B (Test Plan). Prefill time in s, lower is better (without a / without b -> with a / with b):

  • 8k tokens 1.116 / 1.127 -> 1.122 / 1.113 (-0.4 %); 32k 4.130 / 4.163 -> 4.152 / 4.246 (+1.3 %); 150k 22.906 / 22.914 -> 20.871 / 20.881 (-8.9 %); 280k 51.255 / 51.413 -> 43.537 / 43.516 (-15.2 %).
  • Decode time per token 16.86 / 17.01 -> 16.98 / 17.05 ms (58.8-59.3 -> 58.7-58.9 tok/s); the indexer decode path is unchanged.
  • Greedy tokens identical in all four counted boots for both prompts. The largest top-5 logprob gap against the first "without" boot is 0.062 for the "with" boots and 0.125 for the second "without" boot, so the fused scorer stays within the boot-to-boot spread of the cuBLAS route.
  • Peak memory.used of the RTX stage 30,025 MiB with the fused scorer and 30,165 MiB without; the V100 stage 30,374 MiB in both.
  • The mini checkpoint's text is not meaningful (8 of 43 layers); only equality and distances are evaluated.

Full boot of this branch (Test Plan), one boot, no A/B:

  • Route lines: "SM70 indexer prefill: fused SM75 scorer" on cuda:0 and cuda:4 (the two RTX 8000 stages), "SM70 indexer prefill: cuBLAS route" on cuda:1 to cuda:3 (the V100 stages).
  • Eight quality answers (1500-token limit), read by hand: all correct and all ended with stop (four-stroke engine, a Python function for the n-th prime, a speed calculation with the right result of 84 km/h, a question with a misspelled name declined instead of answered with invented facts, an ordering puzzle, a code bug, a summary, an English explanation).
  • Long-context probe, 400 decoded tokens per prompt; time to first token in s (lower is better), decode ms per step and tok/s (higher is better; one token per step without DSpark): 32,668 tokens 10.00 s, 114.7 ms, 8.7 tok/s; 178,480 tokens 60.68 s, 197.6 ms, 5.1 tok/s; 287,039 tokens 107.66 s, 260.0 ms, 3.8 tok/s; 302,833 tokens 115.14 s, 268.9 ms, 3.7 tok/s. All four answers are coherent summaries of their prompt.
  • These decode numbers are those of main plus this stack without speculative decoding; apart from commit 2 (model A/B above) this PR does not change the decode path.
  • KV cache 317,261 tokens. Peak memory.used per card (GPU 0 RTX 8000, 1 V100, 2 RTX 8000, 3 V100, 4 V100), MiB: boot and client 39,745 / 32,202 / 35,575 / 32,202 / 32,202; quality and long-context probe 39,231 / 32,284 / 35,635 / 32,284 / 32,284.

Kernel microbenchmarks, from the commit messages (CUDA graphs, lower is better):

  • Indexer epilogue, [508 tokens x 64 heads, 1024 keys] fp16 tile (66.6 MB): RTX 8000 632 -> 118.6 us (105 -> 562 GB/s), V100 362 -> 85.1 us (184 -> 782 GB/s).
  • Sparse attention pass, us per call: prefill T128 H64 W640 RTX 8000 1438.3 -> 755.7, V100 1178.7 -> 632.0; W2176 RTX 8000 4666.8 -> 2367.7, V100 3760.3 -> 1983.5; decode T6 H64 (128 + 512 keys) V100 146.5 -> 96.5, SWA only (128 keys) 55.1 -> 36.6.
  • Turing indexer scores, RTX 8000, 512 tokens x 64 heads, ms per call, cuBLAS + epilogue -> fused: 1024 keys 0.326 -> 0.151, 4096 keys 1.270 -> 0.586, 18750 keys 5.551 -> 2.305, 37500 keys 11.018 -> 4.518 (57-70 TFLOPS); max error against fp64 equal to the cuBLAS route in all shapes.
  • Compressed KV gather, us per call (rows 2048 / 18750 / 37500): RTX 8000 47.2 -> 7.2 / 532.3 -> 107.3 / 1061.9 -> 215.6; V100 45.8 -> 7.7 / 616.4 -> 87.4 / 1229.3 -> 167.7 (about 280 and 345 GB/s at 18750 rows).
  • Logits per sub-chunk (DeepSeek-V4 8-layer mini checkpoint, PP=2, 1024-token chunks, 280k-token prompt, V100 stage, torch allocator peak in MiB): budget 160 before the change 16759 (two sub-chunks alive at once, 160 + 113 MiB), after 16646; budget 512 (one sub-chunk of 273 MiB) 16759. Prefill time with budget 512 against 64, A/B/A/B, s: 8k 0.854 / 0.843 vs 0.845 / 0.837, 150k 15.593 / 15.593 vs 15.596 / 15.599, 280k 32.849 / 32.758 vs 32.888 / 32.871.

Unit tests on this branch at 437ad09, native build in place (_C carries 73 sm_70 and 4 sm_75 cubins), one process per file:

  • Quadro RTX 8000: test_deepseek_v4_sm70_indexer.py 31 passed; test_deepseek_v4_sm70_routes.py 30 passed; test_dsv4_sparse_prefill_bmm.py 17 passed; test_dsv4_sm70_sparse_bmm.py 12 skipped (its decode reference kernel is Volta-only); test_deepseek_v4_sm70_qnorm_rope_kv_insert.py 2 passed; test_compressor_kv_cache.py 32 passed, 6 skipped; test_sparse_indexer_decode_seq_lens.py 11 passed; test_cudagraph_dispatch.py 51 passed.
  • Tesla V100: test_deepseek_v4_sm70_indexer.py 30 passed, 1 skipped (the fused-route case needs SM75); test_deepseek_v4_sm70_routes.py 30 passed; test_dsv4_sparse_prefill_bmm.py 17 passed; test_dsv4_sm70_sparse_bmm.py 12 passed; test_deepseek_v4_sm70_qnorm_rope_kv_insert.py 2 passed; test_compressor_kv_cache.py 32 passed, 6 skipped; test_sparse_indexer_decode_seq_lens.py 11 passed; test_cudagraph_dispatch.py 51 passed.
  • The two model-head-count cases of test_deepseek_v4_sm70_indexer.py (32 and 64 heads) compare with an fp64 reference that rounds the per-head scores to fp16 as every prefill route does.

pre-commit run --from-ref origin/main --to-ref HEAD: every hook passes except check-layering, which reports vllm/model_executor/layers/fused_moe/runner/moe_runner.py (platform coupling 37 -> 39); this PR does not touch that file and main reports the same. pre-commit run mypy-3.10 --hook-stage manual: passed. The pre-commit CI job fails on main's own run at 899e9ed (run 40532180681) with mypy errors in vllm/v1/kv_cache_interface.py:156-158 and the same check-layering count; none of these files is touched here.

🤖 Generated with Claude Code

https://claude.ai/code/session_0172gtUiymppDzziDqKfy78Q

Peuqui added a commit to Peuqui/1Cat-vLLM that referenced this pull request Oct 27, 2026
- indexer_prefill_scores_sm75 also requires an even query token stride:
  it loads query pairs as 32-bit words.
- Each indexer prefill scoring route logs its choice once per process
  (fused SM75 scorer, cuBLAS route, Triton scorer with or without tensor
  cores). With the default scope "local" only the first local rank,
  that is pipeline stage 0, logged.
- SparseAttnIndexerKpool frees each prefill sub-chunk's logits before
  the next one is scored, like SparseAttnIndexer.
- The compressed KV gather comment gives the measured rates instead of
  "the memory rate"; the indexer_prefill_cublas docstring names the
  Turing chunks that the fused scorer takes before that choice.
- The workspace test says why it sets sparse_tensor_cores = False.

Tests: tests/kernels/test_deepseek_v4_sm70_indexer.py 31 passed, 8
skipped (tensor-core Triton cases need the sm7x Triton) and
tests/models/test_deepseek_v4_sm70_routes.py 30 passed on an RTX 8000.
The kernel source equals the one built and booted in the 1CatAI#1280 branch;
this tree's _C is not rebuilt.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0172gtUiymppDzziDqKfy78Q
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Peuqui and others added 15 commits October 28, 2026 06:24
….512

Turing ran block-FP8 QPN8 linears through the Volta m8n8k4 packed GEMM for
M <= 8 (half tensor rate on SM75) and above M=8 reconstructed the whole FP16
weight on every call before cuBLAS. A new m16n8k8 GEMM decodes and
block-scales the packed codes once per CTA while staging them in shared
memory (bit-identical to the dense reconstruction), reads both operands
through ldmatrix, rasterises row tiles in groups for L2 reuse and splits K
across CTAs with a fixed-order FP32 reduction. Tile and split come from a
wave model over the kernel instances' occupancy. No second weight layout.

Turing runs M=9..512 through it; M <= 8 keeps the packed GEMM and larger M
the dense prefill. Timed without Python dispatch (20 calls per CUDA graph,
replays interleaved, Quadro RTX 8000), the tensor-core GEMM took
1.01-1.28x the time of the packed GEMM at M=1..8 on 10 of 13 DeepSeek-V4
and Qwen3.8-27B projection shapes and 0.78-1.02x on the other three; at
M=16..64 it took 0.16-0.60x the time of the previous path. A DeepSeek-V4
A/B (PP5, DSpark K=5, decode M=6 on the RTX stages) with the GEMM also at
M <= 8 had 1 ms slower decode steps and 0.5 % faster prefill.

The FP8 decoder and lane map move into fp8_qpn8_codes_sm70.cuh, shared by
both kernels. The Turing source ships sm_75 code even in a 7.0-only build.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0172gtUiymppDzziDqKfy78Q
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
The scale of a K step was read from global memory in store_shared, after
the MMAs of the previous step, so its latency sat on the path to the next
__syncthreads. It is now loaded in load_global together with the codes and
carried in registers. Same arithmetic, bitwise equal outputs (840 tile and
split configurations on the 27B TP2 shapes); the BM128 variant also drops
its 32-byte stack frame. RTX 8000, CUDA graphs against the previous kernel:
M 128-512 with BM128 7-12 % faster; BM16, BM32 and BM64 within 2 %.
Prefetching two K steps was measured as well and rejected: more registers
(BM16 54 -> 64, BM32 72 -> 82), fewer resident blocks, slower for every
tile height.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0172gtUiymppDzziDqKfy78Q
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
A CTA of the M <= 8 QPN8 GEMM covers 32 output columns with split-K warps
along K, so it holds 32 * split_k threads. On Turing (1024 threads per SM)
a split-16 CTA fills half an SM, so N = 5120 (160 CTAs) spills 16 CTAs
past the 144 slots into a second, almost empty wave. Split-K 8 fits the
grid into one wave: o_proj/out_proj (K 3072) 0.78-0.90 and down_proj
(K 8704) 0.73-0.87 of the split-16 time (CUDA graphs, M = 1..8, RTX 8000).
The V100 holds 320 split-16 CTAs per wave and loses with split-K 8 on
spills (1.18-1.24), so Volta keeps its table.

The fitted split applies to the M <= 8 launch only, where it was measured:
the per-tensor path keeps the table split for its M16 and M32 bodies (the
native M=17..32 candidate takes split-K 12 or 16), and the block path
reaches the packed GEMM only up to M 8.

Split-K 32 for grids with no more CTAs than SMs was faster in isolation
(DeepSeek-V4 wq_a/wkv 0.79-0.90 of split 16) but changed no DeepSeek-V4
step time in an A/B/A/B (62/65/70 ms both, prefill 250k 222.25-222.28 vs
222.16-222.37 s) and lowered the sampled draft acceptance in all eight
comparisons (0.1-1.5 points), so it is not used. Qwen3.8-27B NVFP4 (FP8
attention projections, TP2 on two RTX 8000) gains 1.4-1.8 % decode tok/s
in an A/B/A/B, texts byte-identical in all boots. Output error against FP64
is unchanged (per-tensor 4.1-4.9e-4 at M 1-100).

Tests: tests/quantization/test_turing_qpn_kernel_selection.py 33 passed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0172gtUiymppDzziDqKfy78Q
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
The grouped QPN kernel launched one grid row per routing slot (tokens *
top_k). A batch never has more groups than experts, so in prefill most
rows were padding groups that only read goff and exit: DSv4 at 512 tokens
launched 3072 rows for at most 256 groups, Flash-Next at 748 tokens 7480
for 512. The grid now spans min(slots, experts) rows, still a static
bound for CUDA graphs; decode (slots < experts) launches as before.

Bit-identical output. CUDA-graph microbench of the kernel source with only
this line changed against the previous build (sm_70 SASS on both cards,
synthetic weights, uniform routing; us per call, lower is better):
- DSv4 T512 (E256 top-6): RTX 8000 w13 7070 -> 6506, w2 3546 -> 3250;
  V100 w13 5524 -> 5004, w2 3004 -> 2442
- Flash-Next PP T748 (E512 top-10): RTX w13 3800 -> 3297, w2 2682 -> 2002;
  V100 3127 -> 2567, 2557 -> 1634
- Flash-Next TP2 T748: RTX 1961 -> 1653, 2113 -> 1329; V100 1663 -> 1275,
  1977 -> 1052
- T6 decode: unchanged

Tests: tests/kernels/moe/test_skinny_sm70_moe.py 77 passed on V100 and RTX
8000 with the rebuilt _C; pre-commit passed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019x3uEMYxiUMDKrA6E8o8sP
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Turing executes Volta's mma.m8n8k4 at half its tensor rate. Expert
weights are now packed at load in the fragment order of the card's
MMA: m8n8k4 on SM70, m16n8k8 from SM75 on (a pure permutation, same
size, no second copy; each layer lives on one card). The new kernel
uses the weights as the 16-row A operand and the expert's rows as the
8-column B operand, reads each activation slice once per warp (no
shared-memory stage), shares each activation load across two
32-column tiles and each weight decode across four 8-row tiles in
prefill. The decoded weights equal the m8n8k4 path bit for bit; only
the tensor core's summation order differs.

Microbenchmark on a Quadro RTX 8000 (CUDA graphs, per call, same
checkpoint weights in both orders):
  DSv4 512 tokens  w13 6555 -> 4303 us, w2 3271 -> 2238 us
  GLM-5.3 512      w13 8566 -> 5316 us, w2 4205 -> 2760 us
  decode (4-6 tokens) at parity (-1.3 % .. +1.0 %)
The m8n8k4 order of the new source is bit-identical to the previous
kernel; the m16n8k8 order stays within 6.3e-4 of the output maximum
and matches the fp32 reference as closely.

The shared header mma_fragments_sm70.cuh holds the e2m1 decoder (derived
from TurboMind's cvt_f16x8_e2m1), the m16n8k8 MMA and ldmatrix, replacing
the copies in nvfp4_qpn4_sm70.cu, skinny_moe_qpn_sm70.cu and
fp8_qpn8_mma_sm75.cu; below SM75 the MMA and ldmatrix helpers trap. The
split-K choice per fragment order stays in grouped_splitk.

skinny_moe_qpn_sm70.cu ships sm_75 code next to its sm_70 code. The rest
of the SM70 family stays sm_70 and runs on Turing through binary
compatibility: a Qwen3.8-27B A/B/A/B of the whole family built for 7.0 and
7.5 against this split showed no gain from the extra images.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019x3uEMYxiUMDKrA6E8o8sP
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
The Turing sweep's structure carries over to the m8n8k4 kernel: a CTA
takes two 32-column tiles that share each staged activation fragment,
and both GEMMs prefer 8 split-K warps (16 or 10 only where 8 does not
divide K). Prefill keeps two 8-row blocks per weight load (four helped
DSv4 by 1 % and slowed GLM-5.3 by 6 %).

V100, CUDA graphs, us per call against the previous kernel:
  DSv4 512 tokens  w13 4979 -> 4374 (0.88), w2 2441 -> 2395 (0.98)
  GLM-5.3 512      w13 6878 -> 5954 (0.87), w2 3567 -> 3218 (0.90)
  decode (4-6 tok) 0.80 .. 0.97
With the same split-K the tile counts change no bit; the new split
regroups the reduction.

The two-accumulator variant (nacc 2) had no caller and leaves the
kernel and the op signature; the N % 64 check and the grid are shared
by both fragment orders.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019x3uEMYxiUMDKrA6E8o8sP
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Real routers skew: in a DeepSeek-V4 prefill chunk a few experts take
hundreds of rows while others take a handful. The large MXFP4 dispatch
test only drew uniform ids (about 256 rows per expert); it now also runs
a skewed draw that checks its own shape (one expert over four 32-row
passes, one under 16 rows), so the hand kernels' multi-pass groups and
single partial blocks are compared against the reference in one call.

Tests: tests/kernels/moe/test_skinny_sm70_moe.py -k large_mxfp4: 2
passed on V100 (m8n8k4), 4 passed on RTX 8000 (both orders).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
… without P2P

For two-rank tensor-parallel groups whose GPUs NCCL cannot connect peer to
peer (P2P disabled or unavailable), one kernel exchanges the FP16 payload
through a pinned host buffer both ranks map: each block writes its chunk,
fences, raises a flag, waits for the peer's flag and adds the peer's chunk.
One FP16 rounding and commutative, so the result has NCCL's bits. No CPU
work, capturable in CUDA graphs. It sits after the SM70 ring in the
all-reduce chain and takes FP16 payloads up to
kernel_config.host_reduce.max_bytes (256 KiB); larger ones stay on
NCCL. kernel_config.host_reduce.enabled turns it off.

The exchange kernel is adapted from kernels/skinny_ar.cu of
mzen17/v100-skinny-unify (MIT, notice kept in the source), turned into
handle-based channels so a process can own more than one two-rank group.

Co-authored-by: Mike Zeng <92705460+mzen17@users.noreply.github.com>
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_014ZLMmGB8mjc43VpwPxLQNn
(cherry picked from commit f53a668)
FA2's Turing build runs causal head-dim-256 prefills over an fp16
paged cache slowly: its split-KV kernel keeps Q in shared memory next
to K and V, a CTA fills the SM's 64 KB, and without cp.async every
tile's loads stall the math. A 2048-token chunk after 150k tokens of
context takes 165.8 ms there, about 23 TFLOPS.

The new kernel holds 128 packed query rows per CTA in registers
(eight warps of 16; the query heads of a KV head share each loaded
key), double-buffers K/V tiles of 32 keys in shared memory, and
fetches the next tile in two register-staged halves while the
current one computes. 32-bit cache offsets and a compile-time page
size keep the MMA loop free of spills; the O rescale is skipped while
every factor of the warp is 1. The source is built for sm_75 only.

Quadro RTX 8000, Qwen3.8-27B per-rank shapes (12 query heads, 2 KV
heads), causal chunk at the end of the context, ms per call:
  2048 tokens  32k ctx  35.3 -> 13.1  150k ctx 165.8 -> 61.5 (61 TFLOPS)
  512 tokens   150k ctx  52.5 -> 20.0
  96 tokens    150k ctx  24.6 -> 19.8
  64 tokens    150k ctx   5.3 vs 19.9  (FA2 splits the context there)
Errors against an fp32 reference equal FA2's on every case.

The FLASH_ATTN backend takes the kernel on SM75, by the capability of
the query's own device, once a request in the batch brings at least
attention_config.flash_attn_turing_prefill_min_query_len (65) query
tokens: causal FA2 attention with an fp16 query over a non-quantized
cache, without sliding window, ALiBi, sinks, softcap or batch
invariance. The launcher takes pages of 16, 32 or 64 tokens and forms
32-bit cache offsets, so the selection keeps other page sizes and
cache views of 2^31 elements or more on FA2.

Tests: tests/kernels/attention/test_paged_prefill_attention_sm75.py
(fp32 reference cases on an SM75 GPU; the selection keeps pages of
128 and 800 tokens and a cache beyond 32-bit offsets on FA2).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019x3uEMYxiUMDKrA6E8o8sP
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
MLA prefill backend selection asks
flash_v100_dense_prefill_lse_available() while the model is built,
before the engine binds its Flash-V100 runtime. The check went through
get_flash_dense_forward(), which binds the operation and raised
"Flash-V100 resources must be bound before execution": the first
GLM-5.3-Flash boot failed with it. The loader is now separate from the
bound accessor; availability only loads, execution still binds.

Tests: tests/v1/attention/test_flash_v100_dense_prefill_availability.py,
2 passed (availability without a bound runtime; execution still
requires it).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019x3uEMYxiUMDKrA6E8o8sP
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Without peer access, a TP2 all-reduce goes through host memory: on two
Quadro RTX 8000, NCCL's all-reduce kernels took 96 of 208 s of a
150k-token Qwen3.8-27B prefill that already ran the Turing prefill
attention. kernel_config.tp_prefill_overlap_min_tokens (0 = off) now
runs eager prefill batches of at least that many tokens, without draft
tokens or padding, as two interleaved halves, after the DBO
micro-batch mechanics: both halves run the model on one compute stream
in two threads, and each all-reduce (GroupCoordinator's
_all_reduce_out_place and all_reduce_sum2) runs on a comm stream while
the other half computes its layer. The field stays out of the compile
hash: the compiled graph is the same, only called twice.

A half is a token range of the batch
(prefill_overlap.token_range_batch): a request cut at the split ends
there in the first half and continues in the second, as a following
chunk would, so the second half's attention reads the first half's KV
writes of the same layer. Each half builds its attention metadata with
its own builder (builders 1 and 2; builder 0 keeps the whole batch's
metadata for the speculator). In mamba align mode both halves must use
the recurrent state slot of the chunk end: align_ubatch_split moves the
split so that the first half also ends inside that state block (801 +
799 tokens for a 1600-token chunk of 800-token blocks), and the second
half continues the state the first half wrote in place. The output
concatenation moved to ubatching (cat_ubatch_outputs) for both runners.

Needs TP without PP, DP or decode context parallelism. Errors at
startup reject LoRA, KV transfer, PLE offload, mamba cache mode "all",
encoder-decoder models, models with FusedMoE layers (their modular
kernel asserts outside DBO when prepare/finalize is not async), and
models with deepstack visual features (Qwen3-VL's deepstack_num_level,
Granite 4 Vision's _ds_layer_indices), which sit in a model buffer that
each half would read from token 0. Inside the halves the
SM70 fused all-reduce + RMSNorm ops raise, since they would run a
collective on the compute stream beside the other half's all-reduce on
the comm stream, and model inputs other than input_ids, inputs_embeds
and positions raise.

A failing micro-batch now sets a shared abort flag before it wakes the
next one, which raises MicroBatchAborted at its next yield instead of
waiting forever for the failed thread; PrefillOverlap re-raises the
original error. The existing DBO micro-batches use the same contexts.

Tests: tests/v1/worker/test_prefill_overlap.py, 14 passed on an RTX
8000 (token ranges, M-RoPE input split, align-mode split, rejection of
DCP, encoder-decoder, MoE and deepstack models, and on a GPU: two
halves with a slow stand-in all-reduce on the comm stream equal the
whole batch bit for bit, and a failure in either half stops both and
reports its error); tests/v1/worker/test_gpu_model_runner_v2.py, 5
passed (its two failing
test_sm70_v2_mtp_profile_reports_from_last_pp_stage cases fail the
same on main).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Without peer access, NCCL moves TP all-reduces through host memory with
SM copies (LL protocol), and so does the host reduce's kernel mode.
Their polling of host memory slows the kernels running beside them. In
an nsys profile of a 100k-token Qwen3.8-27B prefill with the TP prefill
overlap on two Quadro RTX 8000, the NVFP4 weight dequantization kernels
ran 31 and 7 % longer per call, the FP8 ones 18 to 37 %, and the
largest cuBLAS GEMM 22 % longer per tile row than in the baseline
profile without overlap, so overlapping gained little.

host_reduce.copy_engine_max_bytes (0 = off; otherwise above max_bytes)
adds a mode for payloads above max_bytes: the copy engines move this
rank's tensor into its host slot and the peer's back in up to 16
pieces, ordered by stream memory operations (cuStreamWriteValue32 once
a piece landed, cuStreamWaitValue32 before reading the peer's), then
one FP16 add gives the same bits as the kernel mode and NCCL's
two-rank sum. Sending runs on the channel's own copy stream, receiving
and the add on the caller's stream; slots alternate by call parity and
are reused only after the peer acknowledged reading them. The mode
pins four slots of copy_engine_max_bytes in shared host memory.

The call sequence lives on the host, so the mode cannot be captured:
while a stream is capturing, payloads above max_bytes are rejected
(copy_engine_mode_not_capturable) and the next all-reduce in the chain
serves them, as before. Decode payloads stay on the kernel mode.

RTX 8000 pair, 8,192,000 B per call, eager: 3.15 ms (2.60 GB/s)
against NCCL's 3.93 ms (2.08 GB/s); a GEMM loop beside it changed its
time by -2.8 % against +18.2 % beside NCCL, a memory-bound loop by
+1.8 % against +6.0 %. In the 100k-token profile with this mode the
prefill took 95.85 s instead of 118.22 s, the dequantization kernels
ran 4 to 11 % and the largest GEMM tile 0.3 % longer than without
overlap.

Tests: tests/distributed/test_host_reduce.py on the RTX pair with
NCCL_P2P_DISABLE=1, 3 passed (new: copy-engine mode bit-identical to
NCCL for 0.3 to 21 MB, back to back on a side stream beside GEMMs and
mixed with the kernel mode; a captured payload above max_bytes is
rejected without failing the capture; the config rejects a
copy_engine_max_bytes at or below max_bytes).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
One CTA of paged_prefill_attention_sm75 fills a Turing SM (64 KB of
shared memory), so its grid loses up to a whole wave to the last one: a
1600-token chunk of Qwen3.8-27B per TP rank is 150 CTAs on the RTX
8000's 72 SMs, three waves for 2.1 waves of work, and an 800-token half
of the TP prefill overlap is 76 CTAs, two waves for 1.06.

Each CTA's key range can now be split into whole-tile ranges; the
splits write unnormalized fp32 partial outputs with their row maxima and
sums, and merge_splits_kernel combines them by log-sum-exp. The launcher
picks the split count (up to 4) that finishes the CTAs in the fewest
full-CTA times on this device (SMs times the kernel's occupancy); the op
takes num_kv_splits, 0 for that choice, for tests and benchmarks.

RTX 8000, 12 query and 2 KV heads, head dim 256, CUDA graph, ms per call
(without split-KV / automatic): 1600 tokens after 30k 12.40 / 9.71,
after 100k 40.42 / 31.75, after 184k 75.59 / 58.47; 801 tokens after
100k 27.08 / 17.14, after 184k 49.65 / 32.07. One forced split keeps
the time without split-KV (40.41 at 100k); 1600 after 1600 is
1.28 / 1.08.

Tests: tests/kernels/attention/test_paged_prefill_attention_sm75.py,
21 passed on an RTX 8000 (the fp32 reference cases with 0, 1, 2 and 4
splits, including splits without visible keys, and the selection
cases).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
The prefill path of the Turing NVFP4 linear dequantizes each weight into
an fp16 buffer for cuBLAS. Its Triton kernel took 64 groups per program,
the fastest tile on Volta, which holds 128 values per thread on Turing,
past Triton's register budget there: about 200 GB/s on the RTX 8000.

Sweep at the Qwen3.8-27B TP2 shapes per rank (GROUPS_PER_BLOCK 4..64 x
1..8 warps, every configuration bit-identical): Turing runs fastest at
16 groups and 8 warps, ms per call today -> new: 17408 x 5120
1.090 -> 0.496 (460 GB/s), 5120 x 8704 0.580 -> 0.250, 5120 x 2560
0.166 -> 0.075, 8192 x 5120 0.525 -> 0.230. On the V100 64 groups stay
the fastest for the large shape (17408 x 5120 0.930 ms), so Volta keeps
them; the tile is chosen per device.

Tests: tests/kernels/quantization/test_nvfp4_qpn2_dequant.py on an RTX
8000 and a V100, 11 passed each.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
The relu -> weight -> head-sum epilogue of the cuBLAS prefill scoring ran
BLOCK_K 256 over 4 warps, two fp16 keys (4 bytes) per thread and load. On
the [508 tokens * 64 heads, 1024 keys] score tile of a 150k-token prefill
it read 66.6 MB in 632 us on RTX 8000 (105 GB/s) and 362 us on V100
(184 GB/s); BLOCK_K 512 over 2 warps (eight keys, 16-byte loads) takes
118.6 us (562 GB/s) and 85.1 us (782 GB/s), the cards' plain read rate of
that tile (565 / 780 GB/s). Measured in CUDA graphs, 20 launches each.

In a per-stage nsys profile of the DSv4 PP5 entry (150k-token prefill)
this epilogue took 14.0 s and 17.6 s of the two RTX stages, the
pipeline's slowest, and 7.7 s of each V100 stage.

The head sum is fp32 either way; the different thread layout changes its
reduction order in the last bits (up to 1.5e-5 on unit-scale scores), so the
selected top-k can differ only at near ties. Prompts up to index_topk * 4
tokens select every key and are unaffected. The fp32 decode scores already
load 8 bytes per thread; their launch is unchanged.

Tests: tests/kernels/test_deepseek_v4_sm70_indexer.py 26 passed on RTX
8000 and V100; pre-commit and mypy-3.10 passed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01YaG8VjXn5eN3Aq7eZrA3dp
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Peuqui and others added 5 commits October 28, 2026 08:20
…ather

The batched-matmul sparse MLA attention ran six elementwise torch ops
around its two BLAS calls: a broadcast masked_fill over the gathered
keys, then scale, copy into the sink-padded logits, mask, sink column,
softmax and the fp16 cast of the probabilities. On a 128-token pass with
64 heads and 640 keys they took 820 of 1434 us on RTX 8000 in the
per-stage nsys profile of the DSv4 PP5 entry (150k-token prefill); the
masked_fill alone, a non-vectorized broadcast kernel, took 395 us, twice
the gather.

Now the gather writes zeros for unused slots and never reads their indices
(a Triton row gather for prefill; the decode dequantizing gather takes the
same mask), and one Triton kernel reads the fp32 scores once and writes fp16
probabilities: scale, mask, sink column and the fp32 softmax with libdevice
exp, the same expf torch's softmax uses. The fp32 logits and probs buffers
go away (140 -> 110 MiB per 128-token pass at 640 keys).

DeepSeek-V4.1 (1CatAI#1277) calls attend_gathered_keys as well. Its two callers in
vllm/models/deepseek_v41/attention_ops.py now zero the unused slots
themselves, since their gathers leave them as read, and pass the
probabilities as [T, H, W] in the key dtype. Against main's function on
random data with NaN in the unused slots, 99.69-100 % of their outputs are
bit-identical and the rest differ by one fp16 step (at most 2.4e-4), on a
Tesla V100 and a Quadro RTX 8000.

Old against new in CUDA graphs (us per call, 10 calls per graph):
  prefill pass T128 H64 W640:  RTX 8000 1438.3 -> 755.7, V100 1178.7 -> 632.0
  prefill pass T128 H64 W2176: RTX 8000 4666.8 -> 2367.7, V100 3760.3 -> 1983.5
  decode T6 H64 C4 (128 + 512 keys): V100 146.5 -> 96.5
  decode T6 H64 SWA only (128 keys): V100 55.1 -> 36.6
Error against a float64 reference is unchanged (max and mean equal to three
digits); 99.2-99.8 % of the outputs are bit-identical, the rest differ by at
most 3e-5 (softmax reduction order).

Tests: tests/kernels/attention/test_dsv4_sparse_prefill_bmm.py and
test_dsv4_sm70_sparse_bmm.py 29 passed on V100; on RTX 8000 17 passed, 12
skipped (the decode reference kernel is Volta-only); pre-commit and
mypy-3.10 passed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01YaG8VjXn5eN3Aq7eZrA3dp
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
The previous commit fused the sink softmax: decode and prefill keep
FP32 scores and FP16 probabilities of the selection width, with no
extra sink column and no separate logits buffer. The two workspace
tests still expected four buffers with width + 1 and failed since then.
The decode test now takes the last three BMM arguments; the prefill
test expects the gathered keys, the FP32 scores and the FP16
probabilities.

Tests: tests/models/test_deepseek_v4_sm70_routes.py, 30 passed on an
RTX 8000 and a V100.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
…ores

The SM70-family relu route scores prefill chunks by stacking the heads into
M for one cuBLAS GEMM, writing the [tokens * heads, keys] fp16 scores (64 MiB
per 512-token chunk and 1024 keys) and reading them back in a second kernel
for relu, head weights, head sum and key scales. On the DSv4 RTX stages that
pair took 6.9 s of a 51 s 150k prefill (nsys, cuBLAS 4.07 s + head sum 2.80 s),
and the RTX stages bound the pipeline (V100 stages 47-48 s).

indexer_prefill_scores_sm75 does it in one pass: each warp holds one token's
queries for 32 heads as m16n8k8 A fragments in registers, a CTA stages e4m3
key tiles widened to fp16 in shared memory (double buffered), scores are
rounded to fp16 as the cuBLAS route stores them, and relu, head weights and
the head sum stay in registers with a shuffle reduce-scatter; only the
[tokens, keys] result reaches memory. 32 or 64 heads (GLM pads 16 to 32).

Microbench (RTX 8000, CUDA graphs, 512 tokens x 64 heads): 1024 keys
0.326 -> 0.151 ms, 4096 keys 1.270 -> 0.586 ms, 18750 keys 5.551 -> 2.305 ms,
37500 keys 11.018 -> 4.518 ms (57-70 TFLOPS); max error against fp64 equal
to the cuBLAS route in all shapes, 32-51 % of values bit-identical (head-sum
order).

The chunk takes it when the indexer runs the relu route on an SM75
device (the query's own) with 32 or 64 heads; other chunks keep the
cuBLAS or per-head Triton route, and indexer_prefill_cublas decides only
between those two. The source builds for sm_75 only, and the op checks
the device, dtypes, strides and alignment it relies on. Each route logs
its choice once per device.

Tests: tests/kernels/test_deepseek_v4_sm70_indexer.py, 31 passed on an
RTX 8000; 30 passed and 1 skipped (the fused-route case needs SM75) on a
V100. The model-head-count cases (32 and 64 heads, 37 queries, 1100
keys) compare with an fp64 reference that rounds the per-head scores to
fp16 as every prefill route does.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019x3uEMYxiUMDKrA6E8o8sP
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
The prefill loop rebinds `logits` per sub-chunk, so the previous
chunk's scores stay alive until the next scorer has allocated its own:
under VLLM_SPARSE_INDEXER_MAX_LOGITS_MB the peak was two sub-chunks,
while the profiling run reserves one. Freeing them after the top-k
bounds the peak by the budget. SparseAttnIndexerKpool runs the same
loop and gets the same release; it was not measured.

DeepSeek-V4 mini checkpoint (8 layers, PP2), 1024-token chunks, one
280k-token prompt, V100 stage, torch allocator in MiB:
  budget 512 (one sub-chunk of 273 MiB): allocated peak 16759
  budget 160 before this change:         allocated peak 16759 (160 + 113)
  budget 160 with this change:           allocated peak 16646
Prefill time, A/B/A/B, budget 512 vs 64 (s, less is better): 8k
0.854/0.843 vs 0.845/0.837, 150k 15.593/15.593 vs 15.596/15.599,
280k 32.849/32.758 vs 32.888/32.871.

Tests: tests/models/test_deepseek_v4_sm70_routes.py, 30 passed on an
RTX 8000 and a V100; the mini runs exercised the loop with 1, 2 and 5
sub-chunks per step.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
Every prefill chunk dequantizes its request's whole compressed cache
(18.75k rows on average over a 150k prefill at ratio 4, 37.5k at its
end) into the workspace the sparse attention gathers from. The kernel
gave each request 128 workers that walk their rows one by one, which
left most SMs idle. 2048 two-warp workers keep the kernel and its
arithmetic and only spread the rows; the output is bit-identical.

CUDA graphs, us per call (less is better), 128 x 4 warps -> 2048 x 2:
  rows          2048    18750    37500
  RTX 8000  47.2 -> 7.2   532.3 -> 107.3   1061.9 -> 215.6
  V100      45.8 -> 7.7   616.4 ->  87.4   1229.3 -> 167.7
The DSv4 150k-prefill profile spends 0.76 s per RTX stage and 0.82 s
per V100 stage in this kernel.

The Triton gather also serves the ROCm backend and NVIDIA GPUs without
CuteDSL (dequantize_and_gather_k_cache); only V100 and RTX 8000 were
measured.

Tests: tests/kernels/test_deepseek_v4_sm70_qnorm_rope_kv_insert.py and
tests/kernels/test_compressor_kv_cache.py on an RTX 8000 and a V100, 2
and 32 passed (6 skipped) each; the microbenchmark above compared every
grid bit for bit with the 128-worker launch.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0141pGKkAa2PjDMhrUwfX3ox
Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com>
@KewaiiGamer

Copy link
Copy Markdown
Screenshot_20261028_193106 insane-all-in-his-mind

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