[DSv4 Perf] Optimize workspace reuse for eager break - #49236
Conversation
Signed-off-by: yewentao256 <zhyanwentao@126.com>
There was a problem hiding this comment.
Claude Code Review
This repository is configured for manual code reviews. Comment @claude review for a one-time review, or @claude review always to subscribe this PR to a review on every future push.
Tip: disable this comment in your organization's Code Review settings.
| if output_buffers is None: | ||
| index_weights_out = torch.empty_like(index_weights, dtype=torch.float32) | ||
| else: | ||
| assert use_fp4 |
There was a problem hiding this comment.
nit: Here we have a "output_buffers implies use_fp4", which could make for a "bumpy" user interface for this function - the function would have a nicer interface without this restriction. Either document this in the docstring or extend support
There was a problem hiding this comment.
Fixed, we extend the support
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Signed-off-by: yewentao256 <zhyanwentao@126.com>
yewentao256
left a comment
There was a problem hiding this comment.
In previous commit of this PR, the test passes https://buildkite.com/vllm/ci/builds/80503#019fa437-a168-432a-b60e-05bc623b2090
So the failed unit test is not related
Two upstream commits drive every one of the 20 conflict hunks: 38a466e [DSV4] Sequence Parallelism (vllm-project#46789) df71917 [DSv4 Perf] Optimize workspace reuse for eager break (vllm-project#49236) Sequence parallelism (dspark.py, model.py) ------------------------------------------ _use_sequence_parallel requires expert-parallel AND (deep_gemm_mega_moe or DP>1), so it is inert on the SM12x config (Marlin MXFP4 MoE, DP=1) even at 4-node TP=4 + EP. It is still wired correctly, because this branch serves multi-node users whose configs can reach it, and a half-applied SP is worse than none. nvidia/dspark.py resolves to ours byte-for-byte. PR #27 replaced the two-kernel mhc_post + hc_head sequence with the fused mhc_post_hc_head_tilelang, which never materializes the [T, hc_mult, H] intermediate -- and that intermediate is exactly what upstream's sp_all_gather needs between the two kernels. Upstream's DeepseekV4DecoderLayer/ModuleList shape is also incompatible with our DSparkLayer, and sp_shard on dim 0 would split DSpark blocks across ranks, breaking the batch_size = num_input_rows // block_size geometry. nvidia/model.py combines both sides, plus three SP correctness edits the merge did not produce on its own: the auto-merged sp_all_gather of hidden_states sits ahead of tail consumers that still take a sharded residual, so the deferred mhc_post path is now disabled under SP and materialization is forced before the gather. Upstream's `if layer is not None:` is not adopted; it is an unbound local when a PP rank owns no layers. Workspace reuse (csrc, attention.py, cache_utils.py) ---------------------------------------------------- Upstream split the fused qnorm-rope-kv-insert op into an allocating `..._insert` and a caller-buffered `..._insert_out` with an explicit q_head_padded argument. We adopt that shape wholesale and retire our local 53b6d11 variant: it solved the same problem, upstream's split is a strict superset, and the merged tree was already half-way there (ops.impl registered both, attention.py auto-merged to call `_out`). Keeping ours would mean reverting an upstream change on a file upstream actively develops -- the same reason we retired vllm-project#48304/vllm-project#48911/vllm-project#48959. ops.h had auto-merged into a broken hybrid, carrying our void/q_out declaration of the base name alongside upstream's `_out`; taking upstream's .cu without fixing it is an ODR mismatch at link time. torch_bindings.cpp's base schema had likewise auto-merged to our `Tensor! q_out` text. attention.py keeps PR #27's value: q is written in place whenever padded_heads == n_local_heads (true on SM12x TP=2), falling back to upstream's eager-scratch pool and then to our class scratch. All three satisfy the op's q_out.size(1) == q_head_padded contract. Side effect: amd/dspark.py:243 already called the 9-argument allocating form, which was broken against our out-form schema. Adopting upstream's shape fixes that pre-existing ROCm break. cache_utils.py takes upstream's `output_buffers` tuple and keeps our four dtype/device asserts on top of its shape-only pair. The d64074e block-table gather bound is untouched. Also updated because the op split is a cross-file contract: - fix_functionalization.py registered only the old name, so our hot path (now `_out`) would have silently lost its copy elision. Both forms are registered with the correct mutated-args maps. - flashinfer_sm120_decode.py and the two kernel-test helpers move to the new signatures.
…mprovement. (vllm-project#49236) Signed-off-by: yewentao256 <zhyanwentao@126.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
…mprovement. (vllm-project#49236) Signed-off-by: yewentao256 <zhyanwentao@126.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> Signed-off-by: root <root@smci355-ccs-aus-m02-09.cs-aus.dcgpu>
Purpose
Part of #45861
Optimize workspace reuse for eager break
Test
vllm serve deepseek-ai/DeepSeek-V4-Flash --tensor-parallel-size 4 --enable-expert-parallel --attention-backend FLASHMLA_SPARSE_DSV4 --attention-config '{"use_fp4_indexer_cache":true}' --kv-cache-dtype fp8 --tokenizer-mode deepseek_v4 --all2all-backend allgather_reducescatter --port 8003Acc
Perf