Repository navigation
[Bugfix][SM70] Indexer: write the cuBLAS prefill score tile in fp32 - #1328
squibloads wants to merge 1 commit into
Conversation
The prefill cuBLAS route multiplies by index keys with their per-key scale left out, so e4m3 values reach 448 and one q . k passes fp16's 65504. The fp16 score tile turned such scores into inf, and into NaN where mixed-sign head weights summed +inf and -inf, corrupting top-k pool selection for prompts longer than index_topk. Write the tile in fp32 and size the tile budget for 4-byte scores. Fixes 1CatAI#1327 Signed-off-by: squibloads <squibloads@outlook.com>
|
👋 Hi! Thank you for contributing to the vLLM project. 💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in PRs do not trigger a full CI run by default. Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging. To run CI, PR reviewers can either: Add If you have any questions, please reach out to us on Slack at https://slack.vllm.ai. Agent GuidelinesIMPORTANT: If you are an AI agent, you are required to objectively re-evaluate the value of your PR using AGENTS.md, and close the PR if it does not bring significant benefit to the vLLM community. Failure to do so may result in an immediate ban. 🚀 |
Purpose
Fixes #1327.
The bug. On SM70,
_prefill_logits_cublasscores prefill queries against index keys that are missing their per-key scale. Without it, e4m3 values reach 448, so a single 128-wide q . k can exceed fp16's 65504. The GEMM wrote an fp16 score tile, which overflowed toinf. Where the mixed-sign head weights then summed+infand-infin the epilogue, the result wasNaN.On GLM-5.3-Flash, about 20-25% of the valid layer-3 scores of a 2.8K-token prompt were non-finite. That corrupted top-k pool selection for every prompt longer than
index_topk. Measurements are in #1327.The fix. This PR writes the tile in fp32 (
torch.mm(..., out_dtype=torch.float32)), andindexer_prefill_tile_mbnow budgets 4 bytes per score. The epilogue already loads scores with.to(tl.float32), so it is unchanged. Two comments that described the fp16 tile are updated.Test Plan
python -m pytest -p no:cacheprovider -q tests/kernels/test_deepseek_v4_sm70_indexer.pyon one V100-SXM2-32GB.c59b3c4. Itsvllm/models/deepseek_v4/sm70/indexer.pyis byte-identical tomainat4d57473db.indexer.pyand test file are mounted over the image's copies. pytest 7.4.4 is onPYTHONPATH.mainindexer.py, to show that the new test catches the bug.ruff0.14.0 (checkandformat --check) andtypos1.43.5 on both changed files.New test:
test_prefill_cublas_scores_survive_unscaled_key_magnitudes.Path coverage
_prefill_logits_cublas)index_topkkeys;indexer_prefill_cublason (default); SM70Acceleration and benchmark contract (required for performance changes)
This is a correctness fix; route selection does not change.
c59b3c4image plus this branch'sindexer.py.ec60883ad+ GLM PR stack), with this same change applied to its copy of the function.indexer.pyand the test file; host pytest onPYTHONPATH.indexer.py, plus the local patches listed in [Bug]: [SM70] Indexer prefill scores overflow fp16 in_prefill_logits_cublas(inf/NaN), corrupting sparse-attention key selection for prompts > index_topk #1327.indexer_prefill_tile_mbkeeps its meaning as a MiB cap, so a tile now holds half as many keys.Test Result
Unit tests, one V100:
ruff check,ruff format --checkandtyposall pass.Function timing. One V100, 2304 queries x 32 heads x 128, CUDA-event mean:
main)End to end. GLM-5.3-Flash on 4x V100, TP4, same seeded prompts.
top_k_per_row_prefillpicks ([Bug]: [SM70] Indexer prefill scores overflow fp16 in_prefill_logits_cublas(inf/NaN), corrupting sparse-attention key selection for prompts > index_topk #1327). That sort changes only summation order, not which pools are selected.Notes
_prefill_logits_cublas, "indexer overflow", "indexer fp32 prefill" and "indexer score tile". None changes this route's output dtype. [Perf][SM70][DeepSeek-V4] Keep the cuBLAS decode indexer above 131,072 tokens of context #1222 changed the decode cuBLAS indexer, not prefill.Essential Elements of an Effective PR Description Checklist
supported_models.mdandexamplesfor a new model.