Repository navigation
[ROCm][Perf] Apply opt-in paged MXFP4 sparse indexer dense path for DeepSeek-V4 - #60834
Conversation
8e4f5df to
77f9dcf
Compare
| def _pooled_k_scratch(num_tokens: int, head_dim: int, device: str) -> torch.Tensor: | ||
| # One buffer for all indexer layers: they run one after another and each | ||
| # consumes its rows before the next layer writes. | ||
| return torch.empty(num_tokens, head_dim, dtype=torch.bfloat16, device=device) |
There was a problem hiding this comment.
Suggestion: Would it be better to do this with the workspace manager?
There was a problem hiding this comment.
Switched to the workspace manager. The indexer layers run one after another, so they all share one persistent buffer. forward claims it on the profiling run, before the workspace locks.
| @@ -0,0 +1,178 @@ | |||
| # SPDX-License-Identifier: Apache-2.0 | |||
There was a problem hiding this comment.
Reminder to self: I think we need to run :amd: (MI355 DPX) Attention Kernels Shard %N manually in CI to make it run on mi350 as it's optional
|
|
||
|
|
||
| @triton.jit | ||
| def _compress_pool_kernel( |
There was a problem hiding this comment.
Suggestion: can we check if we need to register this for warmup?
There was a problem hiding this comment.
With a fresh Triton cache, it was compiling during the CUDA graph warmup run, after JIT warmup. Registering it alone wasn't enough on ROCm: Triton's AMD backend gives pointers into storage of at most 2 GB a pointer_range=32 specialization, but TritonWarmupTensor always reported a small range, while the state cache views the KV pool, which is over 2 GB. So I added ptr_range_bytes to TritonWarmupTensor (default 0, only read by the ROCm backend) and warm up both pool sizes. With that, the variant used at runtime compiles during JIT warmup and nothing compiles after it.
There was a problem hiding this comment.
Good, I actually added the same fix in #60919 but I think we might merge this one first :)
…s kernel Signed-off-by: shen-shanshan <467638484@qq.com>
Signed-off-by: shen-shanshan <467638484@qq.com>
Signed-off-by: shen-shanshan <467638484@qq.com>
46e3e25 to
8f51c20
Compare
@simondanielsson Sorry for late reply, I'm busy with another work recently... Here is GPQA-Diamond A/B test results (also added into the PR body).
|
|
/ci run |
|
✅ Triggered Buildkite CI #96632 for commit |
…ooks vllm-project#60834 added ROCm MXFP4 indexing for DeepSeek V4 on aiter's Triton paged kernels, together with per-platform hooks for the indexer's cache, compressor, scorer and Q quantization. Move this PR's FlyDSL path onto those hooks, behind VLLM_ROCM_USE_AITER_FLYDSL_MXFP4_INDEXER=1, and return the shared indexer code to main: - DeepseekV4RocmFlyDslIndexer keeps the shared compressor, whose fused kernel pools, normalizes, ropes, quantizes and writes K straight into the FlyDSL page layout, and quantizes Q with bf16 weights. - RocmFlyDslIndexerMetadataBuilder plans the prefill windows and the scorers' CTA schedules once per step on top of the dense builder. - RocmFlyDslSparseAttnIndexer scores the paged cache in place. - Without the FlyDSL interfaces vLLM warns and keeps the Triton MXFP4 indexer instead of falling back to FP8. indexer.py, sparse_attn_indexer.py and attention.py are unchanged from main. Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com>
…eepSeek-V4 (vllm-project#60834) Signed-off-by: shen-shanshan <467638484@qq.com>
Slim this PR down to the K writer and drop the FlyDSL scorer, planner, ABI probe, env flag, docs and their tests. Merged vllm-project#60834 already scores the MXFP4 indexer cache with aiter's Triton kernel; this PR now only changes how that cache is written. On ROCm vllm-project#60834 writes indexer K in two launches: a Triton pooling kernel into a bf16 scratch buffer, then aiter's k_norm_rope_mxfp4_cache. Give the existing fused compress + RMSNorm + RoPE + MXFP4 kernel a ROCm branch that reproduces aiter's bf16 rounding points, ceil e8m0 scale, gfx950 FP4 packing and cache_format() page order, and let DeepseekV4RocmMxfp4Compressor launch it once. The written pages are byte-identical to the two-stage path, so the scorer and accuracy are unchanged; the pooling kernel and its scratch buffer are removed. Pass the layout as three integer compile-key fields so JIT warmup compiles exactly the launched variant, and warm it up with KV-pool-sized pointer ranges. Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com>
…ooks vllm-project#60834 added ROCm MXFP4 indexing for DeepSeek V4 on aiter's Triton paged kernels, together with per-platform hooks for the indexer's cache, compressor, scorer and Q quantization. Move this PR's FlyDSL path onto those hooks, behind VLLM_ROCM_USE_AITER_FLYDSL_MXFP4_INDEXER=1, and return the shared indexer code to main: - DeepseekV4RocmFlyDslIndexer keeps the shared compressor, whose fused kernel pools, normalizes, ropes, quantizes and writes K straight into the FlyDSL page layout, and quantizes Q with bf16 weights. - RocmFlyDslIndexerMetadataBuilder plans the prefill windows and the scorers' CTA schedules once per step on top of the dense builder. - RocmFlyDslSparseAttnIndexer scores the paged cache in place. - Without the FlyDSL interfaces vLLM warns and keeps the Triton MXFP4 indexer instead of falling back to FP8. indexer.py, sparse_attn_indexer.py and attention.py are unchanged from main. Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com>
Slim this PR down to the K writer and drop the FlyDSL scorer, planner, ABI probe, env flag, docs and their tests. Merged vllm-project#60834 already scores the MXFP4 indexer cache with aiter's Triton kernel; this PR now only changes how that cache is written. On ROCm vllm-project#60834 writes indexer K in two launches: a Triton pooling kernel into a bf16 scratch buffer, then aiter's k_norm_rope_mxfp4_cache. Give the existing fused compress + RMSNorm + RoPE + MXFP4 kernel a ROCm branch that reproduces aiter's bf16 rounding points, ceil e8m0 scale, gfx950 FP4 packing and cache_format() page order, and let DeepseekV4RocmMxfp4Compressor launch it once. The written pages are byte-identical to the two-stage path, so the scorer and accuracy are unchanged; the pooling kernel and its scratch buffer are removed. Pass the layout as three integer compile-key fields so JIT warmup compiles exactly the launched variant, and warm it up with KV-pool-sized pointer ranges. Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com>
Purpose
#58671 added a paged MXFP4 sparse indexer for DeepSeek-V4.1 on ROCm gfx950. The indexer K cache is stored as MXFP4 (68 B per token instead of 132 B for FP8) in the preshuffled page order that aiter's
pa_mqa_logits_mxfp4reads. That kernel scores the cache in place, so prefill no longer gathers K first. The PR has two layers:This PR brings the dense path to DeepSeek-V4 (e.g. DeepSeek-V4-Pro: 30 indexer layers at compress ratio 4, 64 heads × 128, top-k 1024). Today V4 on ROCm can only use the FP8 indexer cache, because:
dsa_indexer_uses_fp4()rejects every model exceptdeepseek_v41on ROCm;SparseAttnIndexer.forward_hipasserts that the cache is not FP4;cvt.rn.satfinite.e2m1x2.f32), so they cannot run on ROCm.The feature is opt-in, as on V4.1:
VLLM_ROCM_USE_AITER=1 vllm serve deepseek-ai/DeepSeek-V4-Pro -tp 8 \ --attention-config '{"indexer_kv_dtype": "mxfp4"}'The default path (FP8 indexer cache) is unchanged.
Changes
K write: ROCm pooling kernel + aiter's cache op
In V4 the indexer K is produced by the indexer's own
DeepseekCompressor. One fused Triton kernel does pool → RMSNorm → RoPE → MXFP4 → cache store, and nothing before the norm ever reaches memory. On ROCm we split it into two launches:_compress_pool_kernel, does the softmax-gated overlap pooling. It is copied from V4's fused kernel, up to the norm. It writes one bf16 row per boundary token into a scratch buffer that all indexer layers share ([max_num_batched_tokens, 128]).k_norm_rope_mxfp4_cache(..., shuffle=layout)then applies norm → RoPE → MXFP4 and stores into the preshuffled pages. This is the same op V4.1 uses, through the samerocm_mxfp4_indexer_k_storewrapper.This keeps the page layout and the quantization defined only in aiter, which is the principle #58671 follows (
rocm_paged_mxfp4_cache_layoutis the only place vLLM reads the layout).The cost is one extra small launch per indexer layer, plus a bf16 round trip of the pooled row before the norm. That rounding error (~2⁻⁹) is two orders of magnitude below the MXFP4 quantization step. A fully fused ROCm kernel can follow if profiles show the extra launch matters.
Q quant
The Q path uses aiter's
q_rope_mxfp4_quant, through the existingrocm_mxfp4_indexer_q_quant. It produces fp32 weights with the softmax and head scales folded in, as the aiter MQA-logits kernel expects.Scoring and metadata
Both are reused unchanged:
RocmSparseAttnIndexer, which callsrocm_mxfp4_sparse_attn_indexerwith no candidate blocks;RocmMxfp4IndexerMetadataBuilder.A new
DeepseekV4RocmMxfp4IndexerBackendsubclassesDeepseekV4IndexerBackend, so V4's 256-token blocks and ROCm adaptive-verification support carry over. At compress ratio 4 that gives 64-entry pages, which is aiter's preferred page size.Keeping CUDA untouched
Shared V4 code only gains extract-method / ClassVar hooks. The code inside them is moved verbatim, and the ROCm subclasses override the hooks.
Other changes
FusedKVCompressNormRopeInsertIndexerTritonKernel.get_warmup_keys()returns no keys on ROCm with an MXFP4 indexer cache. The warmup is a process-wide singleton that every non-CUDA compressor registers, including V4's head-512 attention compressors, so without this guard the PTX MXFP4 variant was compiled at startup and failed (error: unknown asm constraint 'f'). ROCm never runs that kernel: it writes the cache with aiter's op.dsa_indexer_uses_fp4(): the ROCm branch now acceptsdeepseek_v4as well asdeepseek_v41. The CDNA4 and aiter checks are unchanged, and the CUDA branch is untouched._aiter_indexer_cache_ops,rocm_mxfp4_indexer_k_storeandrocm_mxfp4_indexer_q_quantmove frommodels/deepseek_v41/amd/rocm.pytov1/attention/ops/rocm_paged_mxfp4_indexer.py, so V4 and V4.1 share them. They are still importable from the old module.Test Plan
Unit Test
pytest tests/kernels/attention/test_rocm_dsv4_mxfp4_indexer.py pytest tests/kernels/attention/test_rocm_paged_mxfp4_indexer.py \ tests/v1/attention/test_rocm_paged_mxfp4_indexer_plan.py pytest tests/models/test_deepseek_v4_rocm_compressor_gemm_fusion.py \ tests/models/test_deepseek_v4_rocm_wo_a.py \ tests/models/test_deepseek_v4_dspark_rocm.py \ tests/v1/attention/test_deepseek_v4_rocm_adaptive.py \ tests/kernels/test_compressor_kv_cache.py \ tests/kernels/test_deepseek_v4_cpu_kernels.pyBenchmark
Benchmark with SA InferenceX
8k1kandagenticworkload with server launch commands below.Fixed sequence length:
vllm serve /shared/models/deepseek-ai/DeepSeek-V4-Pro \ --port 8333 \ --tensor-parallel-size 8 \ --data-parallel-size 1 \ --max-model-len 9472 \ --max-num-seqs 512 \ --max-num-batched-tokens 16384 \ --async-scheduling \ --no-enable-prefix-caching \ --distributed-executor-backend mp \ --gpu-memory-utilization 0.8 \ --kv-cache-dtype fp8 \ --trust-remote-code \ --moe-backend aiter \ --tokenizer-mode deepseek_v4 \ --reasoning-parser deepseek_v4 \ --speculative-config '{"method": "mtp", "num_speculative_tokens": 2}' \ --compilation-config '{"mode":3,"cudagraph_mode":"FULL_AND_PIECEWISE"}' \ --load-format fastsafetensors \ --attention-config '{"indexer_kv_dtype": "mxfp4"}'Agentic replay:
vllm serve /shared/models/deepseek-ai/DeepSeek-V4-Pro \ --host 0.0.0.0 \ --port 8333 \ --trust-remote-code \ --async-scheduling \ --distributed-executor-backend mp \ --kv-cache-dtype fp8 \ --max-num-batched-tokens 8192 \ --tensor-parallel-size 1 \ --data-parallel-size 8 \ --prefill-schedule-interval 8 \ --long-prefill-token-threshold 16384 \ --gpu-memory-utilization 0.86 \ --moe-backend aiter \ --compilation-config '{"mode":3,"cudagraph_mode":"FULL_AND_PIECEWISE"}' \ --speculative-config '{"method": "mtp", "num_speculative_tokens": 3}' \ --tokenizer-mode deepseek_v4 \ --tool-call-parser deepseek_v4 \ --reasoning-parser deepseek_v4 \ --enable-auto-tool-choice \ --enable-prefix-caching \ --no-disable-hybrid-kv-cache-manager \ --max-num-seqs 64 \ --load-format fastsafetensors \ --attention-config '{"indexer_kv_dtype": "mxfp4"}'Accuracy Test
Accuracy test with
gsm8kdataset usinglm_eval:Test Result
Unit Test
Benchmark
Fixed sequence length:
In summary, this PR delivers a clear overall improvement over Main: throughput increases across all concurrency levels, Mean TTFT decreases substantially—especially at low and medium concurrency—and Mean TPOT remains slightly better across the board. The results suggest that the optimization is beneficial for both latency-sensitive workloads and overall serving throughput, with the most pronounced gains appearing at lower concurrency.
Agentic replay:
Accuracy Test
GSM8K:
The MXFP4 indexer in this PR shows very small accuracy degradation compared with the FP8 indexer in main. On GSM8K, the flexible-extract exact-match score decreases from 94.92% to 94.77% (-0.16%), while the strict-match score decreases from 95.00% to 94.84% (-0.17%). Overall, the accuracy impact is negligible (<0.2% relative), indicating that the MXFP4 indexer largely preserves the model's original accuracy.
GPQA Diamond (newly added):
Pull Request Checklist
I used vLLM's
/pr-checklistskill. (Mandatory for agents, optional for humans).AI assistance was used during the creation of this PR.
Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.
Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.
Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.
Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.