Skip to content

[DCP][Kernel][Perf] Fuse the empty-shard LSE mask into the A2A pack kernel - #54889

Merged
MatthewBonanni merged 5 commits into
vllm-project:mainfrom
rbrugaro-amd:rbrugaro/kk3-dcp-opt
Sep 10, 2026
Merged

MatthewBonanni merged 5 commits into
vllm-project:mainfrom
rbrugaro-amd:rbrugaro/kk3-dcp-opt

Conversation

@rbrugaro-amd

@rbrugaro-amd rbrugaro-amd commented Sep 2, 2026

Copy link
Copy Markdown
Contributor

[DCP][Kernel] Fuse the empty-shard LSE mask into the A2A pack kernel

Purpose

Under DCP, every MLA layer calls mask_dcp_empty_shards_ (vllm/v1/attention/ops/dcp.py)
immediately before packing the partial attention output and LSE for the all-to-all
combine. It sets the LSE of rows whose request holds no local KV shard — and of
CUDA-graph padding rows — to -inf, so those rows carry no weight in the reduction.

Written in eager PyTorch, that one line expands to eight kernels:

row_indices      = torch.arange(lse.shape[0], ...)              # arange
sequence_indices = torch.searchsorted(query_start_loc[1:], ...) # searchsorted
                     .clamp_max(seq_lens.shape[0] - 1)          # clamp
empty_rows = (row_indices >= query_start_loc[-1])               # compare
             | (seq_lens[sequence_indices] == 0)                # index, compare, or
lse.masked_fill_(empty_rows[:, None], float("-inf"))            # masked_fill

Six of the eight touch only ~100 elements. All eight are captured inside the full
CUDA graph, so each one pays the ~4.3 µs kernel-dispatch floor on every single
replay — roughly 35 µs per MLA layer, entirely independent of context length or
batch size.

The mask is a pure function of seq_lens and query_start_loc, both of which are
identical for every MLA layer within a step, and the pack kernel already loads the
very LSE value that needs masking
. So the mask can be folded into
_dcp_a2a_pack_send_kernel as a Triton-side search over query_start_loc plus a
tl.where, deleting all eight kernels for no additional memory traffic.

This is a parity fix, not a new optimization

vLLM's direct symmetric-memory DCP path already does exactly this inside its
kernel

csrc/libtorch_stable/attention/dcp_utils/dcp_direct_a2a_lse_reduce.cu:

__device__ int64_t find_sequence(const int32_t* query_start_loc, ...)  // binary search
...
peer_lse[...] = empty_kv ? -CUDART_INF_F : to_float(partial_lse[...]);

This PR brings the Triton fallback path to the same behaviour. The direct path is
gated on current_platform.is_cuda() (cp_common.py) and its .cu files are only
built under VLLM_GPU_LANG STREQUAL "CUDA", so the Triton path is what every ROCm
deployment runs, plus any CUDA deployment without symmetric memory. CUDA installs
that do take the direct path are unaffected — they never reach this code.

Design notes

  • CUDA-graph safe. The eight kernels are graph-captured (verified by correlating
    kernel IDs to hipGraphLaunch), so a data-dependent "skip when nothing is empty"
    branch is not a legal alternative — it would be resolved once at capture time and
    baked in. The mask has to be made free, not conditional.
  • num_seqs is a runtime argument, not a constexpr. Only the tl.arange bound
    is constexpr, so the kernel specializes on at most log2(max_num_seqs) variants.
    Measured across every CUDA-graph capture size 2–104: 7 compiled variants, versus
    103 if num_seqs were a constexpr.
  • Semantics are preserved exactly, including the row_indices >= query_start_loc[-1]
    padding term. That term is arguably redundant (padding slots are already zero-filled
    in the DCP local seq_lens buffer), but it is retained so the fused form is
    bit-identical rather than merely equivalent-in-practice.
  • mask_dcp_empty_shards_ is retained for the AG+RS combine path
    (_cp_lse_common), which is not exercised by --dcp-comm-backend a2a. Fusing that
    path is left for a follow-up.

Test Plan

tests/v1/attention/test_dcp_a2a_pack_mask.py compares the fused path against the
eager mask_dcp_empty_shards_ it replaces. The comparison is on bit patterns
(torch.equal over int16 views), not values: for a 2-element LSE pack the high
half of -inf bitcasts to a NaN payload, which never compares equal to itself.

Both sides of every comparison run the same pack kernel: the reference is the new
kernel with HAS_MASK=False fed by the old eager mask_dcp_empty_shards_, so the
only thing that differs is where the mask came from and everything else cancels.

Three tests, 353 cases:

test cases what it covers
test_fused_mask_matches_eager 336 2 dtypes (bf16, fp16) × 4 world-size/head configs (8×2, 8×16, 4×4, 2×1) × 3 query lengths (1, 2, 3 — the latter two are the MTP multi-token-verify shape) × 2 padding-row counts × 7 shard layouts (no empty, interior, all, leading, trailing, the conc-52 decode batch, single request)
test_fused_mask_matches_eager_ragged (new test) 16 Non-uniform query_start_loc. All the cases above use a constant stride, so the in-kernel boundary search is never asked to handle unequal spans — which is exactly what MTP produces under partial acceptance. 8 seeds × 2 dtypes with random per-request query lengths in [1, 4] over 17 requests, a random mix of empty/non-empty shards, and 3 trailing padding rows. Every seed exercises all four query lengths and masks 50–65% of rows, so the comparison is not vacuous
test_mask_disabled_is_unmasked 1 seq_lens=None, query_start_loc=None leaves the LSE untouched
docker run --rm --entrypoint bash --device /dev/kfd --device /dev/dri \
  --group-add video --ipc=host -v $PWD:/work \
  vllm/vllm-openai-rocm:nightly-7c5dc571cbd1064ecc8a9b1045637ff647aa22cb -c '
  SP=$(python3 -c "import vllm,os;print(os.path.dirname(vllm.__file__))")
  cp /work/vllm/v1/attention/ops/dcp.py $SP/v1/attention/ops/dcp.py
  cd /work && python3 -m pytest -q tests/v1/attention/test_dcp_a2a_pack_mask.py'
353 passed in 8.40s

Existing DCP coverage is unchanged — tests/distributed/test_dcp_a2a.py gives the
identical result with and without this patch (the 7 failures are pre-existing on the
unpatched base and are dtype assertions in the test's own reference):

unpatched: 7 failed, 24 passed
patched:   7 failed, 24 passed

Measured effect

Kimi-K3, MI355X ×8, TP8 / DCP8, --dcp-comm-backend a2a, ROCM_AITER_MLA, fp8 KV
cache, FULL_AND_PIECEWISE CUDA graphs, 8k/1k ISL/OSL at concurrency 52, pure decode.
A and B are consecutive runs on the same node with identical configuration, differing
only in this patch. Values are averaged over the per-rank traces (A: 5 ranks, B: 6)
and normalised per MLA layer.

before:
image
after:
image

The eight kernels between kn_mla_reduce_v1_ps and _dcp_a2a_pack_send_kernel
disappear. Identifying them needed no name matching: selecting every kernel whose
invocations-per-MLA-layer drop by ~1.0 returns exactly eight distinct kernels,
matching the eight aten ops in the source function.

−35.00 µs/MLA layer removed, against +1.23 µs in the pack kernel _dcp_a2a_pack_send_kernel that absorbs
the work (6.81 → 8.04 µs) — a net −33.77 µs per MLA layer, or 0.81 ms/step
across 24 MLA layers, against a 35.7 ms decode step (2.3%). Because the saving is
dispatch-floor time it is roughly constant per step, so it is a larger fraction at
shorter contexts and smaller batches.

Nothing else moves. Per-invocation durations for unrelated kernels, A → B:

kernel before after
aiter MLA decode (mla_a8w8_qh16_qseqlen1_gqaratio16_lse_ps) 38.51 µs 38.56 µs
KDA decode fusion 19.60 µs 19.62 µs
DCP a2a (mscclKernel_Sum) 27.95 µs 27.59 µs
MLA split-K reduce (kn_mla_reduce_v1_ps) 7.72 µs 7.74 µs
DCP unpack/combine 4.78 µs 4.58 µs

End-to-end

Serving A/B at 128k/1k, TP8/DCP8, a2a, fp8 KV, decode attention on ROCM_AITER_MLA,
baseline arm run immediately before the fused arm in each pair. Metric is median_itl_ms
— the per-decode-step figure this change moves; throughput and TTFT are excluded because
at 128k both are dominated by prefill, which is untouched.

concurrency KV offload ITL off ITL on gain
52 LMCache 44.52 ms 43.60 ms 2.08%
20 LMCache 31.54 ms 30.96 ms 1.85%
2 none 20.07 ms 19.63 ms 2.21%

All six runs completed with full token counts and no worker deaths, and the C=52 baseline
reproduces an independent earlier measurement of the same configuration (44.70 ms).

Accuracy

GSM8K, 5-shot, same configuration, with the patch applied:

Filter exact_match Stderr
flexible-extract 0.9704 ± 0.0047
strict-match 0.9704 ± 0.0047

(Full 1319-sample test split. No speculative decoding in this configuration.)

Under DCP, every MLA layer calls mask_dcp_empty_shards_ just before packing
the partial attention output and LSE for the all-to-all combine. It sets the
LSE of rows whose request holds no local KV shard -- and of CUDA-graph padding
rows -- to -inf so they carry no weight in the reduction.

Written in eager PyTorch it expands to eight kernels (arange, searchsorted,
clamp, two compares, an index, an or, and the masked_fill), six of which touch
only ~100 elements. All eight are captured inside the full CUDA graph, so each
pays the ~4.3us dispatch floor on every replay: ~35us per MLA layer.

The mask is a pure function of seq_lens and query_start_loc, both identical
across every MLA layer in a step, and the pack kernel already loads the very
LSE value that needs masking. Fold the mask in as a Triton-side search over
query_start_loc plus a tl.where, deleting all eight kernels for no extra
memory traffic.

This brings the Triton path to parity with the direct symmetric-memory path,
which already does exactly this inside its kernel -- see find_sequence() and
`empty_kv ? -CUDART_INF_F : ...` in
csrc/libtorch_stable/attention/dcp_utils/dcp_direct_a2a_lse_reduce.cu.

num_seqs is a runtime argument rather than a constexpr so the kernel
specializes only on the arange bound: 7 variants across batch sizes 2..104
instead of one per distinct batch size.

Measured on Kimi-K3, MI355X, TP8/DCP8, --dcp-comm-backend a2a, fp8 KV,
8k/1k conc-52 decode, full CUDA graphs: the eight kernels drop from ~1.0 to
0.0 invocations per MLA layer, removing 35.0us/MLA layer against a 1.2us
increase in the pack kernel -- a net 33.8us per MLA layer, 0.81ms per step
over 24 MLA layers (2.3% of a 35.7ms step).

mask_dcp_empty_shards_ is retained for the AG+RS combine path, which is not
exercised by --dcp-comm-backend a2a.

Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>
The existing cases all build query_start_loc with a uniform stride, so the
gap between consecutive entries is constant and the in-kernel boundary
search is never asked to handle unequal spans. MTP with partial acceptance
produces exactly that: some requests contribute one row and others several,
in the same batch.

Add eight seeded cases per dtype with random per-request query lengths in
[1, 4] over 17 requests, a random mix of empty and non-empty shards, and
three trailing padding rows. Each seed exercises all four query lengths and
masks 50-65% of rows, so the comparison is not vacuous.

Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>

@claude claude Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

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

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

@rbrugaro-amd

Copy link
Copy Markdown
Contributor Author

@dllehr-amd @AndreasKaratzas can you please review and add the ready/ci-run plz?

@dllehr-amd dllehr-amd 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.

Looks okay to me. I'll wait for additional guidance from Matt or Lukas

@dllehr-amd dllehr-amd added the ready ONLY add when PR is ready to merge/full CI is needed label Sep 9, 2026
@github-actions

github-actions Bot commented Sep 9, 2026

Copy link
Copy Markdown

@rbrugaro-amd, CI is now available for this PR.

  • /ci run starts upstream CI; /amd-ci run starts AMD CI only.
  • /ci retry retries failed jobs in the CI build for the current PR head. If the current head has no CI build, it starts a new CI build for the current head containing only jobs that failed in the latest earlier CI build for this PR.
  • /amd-ci retry retries failed jobs in AMD CI for the current PR head. Use /amd-ci run when the current head has no AMD CI build.
  • /ci cancel cancels scheduled or running CI builds for this PR branch; /amd-ci cancel does the same for AMD CI only.

@AndreasKaratzas AndreasKaratzas left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

not sure of the numerics here, mayube i m wrong, just if you could please provide some info for that. also for the decorator on top of test file.

)

pytestmark = pytest.mark.skipif(
not torch.cuda.is_available(), reason="needs a GPU for the Triton pack kernel"

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Wait does this evaluate to true or false for us? I thought we should be using cuda_alike

@rbrugaro-amd rbrugaro-amd Sep 10, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

It does evaluate correctly to True, but it is the wrong predicate to express the intent, and I have switched it.

torch.cuda.is_available() returns True on ROCm, because the PyTorch ROCm build maps HIP onto the torch.cuda API. So the file was running on our hardware rather than being skipped. Measured in the ROCm image:

torch.cuda.is_available()        = True
torch.version.hip                = 7.2.53211
current_platform.is_cuda()       = False
current_platform.is_cuda_alike() = True

The trap is the neighbouring predicate: current_platform.is_cuda() is False on ROCm and would have silently skipped the whole file.

Switched to [current_platform.is_cuda_alike()] which is what the rest of the tree uses for "CUDA or ROCm"

No behavioural change; 353 tests still pass on MI355X.

num_heads = world_size * h_per_rank

out = torch.randn(num_rows, num_heads, head_dim, device=device, dtype=dtype)
lse = torch.randn(num_rows, num_heads, device=device, dtype=torch.float32)

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

we do we need 32 bit numerics here?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

The LSE is fp32 in production — the test is mirroring the real dtype, not choosing one. All of the following predate this PR:

  • _dcp_a2a_lse_pack_dim exists solely to pack an fp32 LSE into a narrower output dtype, and raises "Cannot pack fp32 LSE into output dtype {...}"
  • out_lse is allocated torch.float32
  • the direct symmetric-memory path allocates received_lse as torch.float32

It has to be fp32 for range — LSE is a log of summed exponentials, and this mask writes -inf into it. A 16-bit LSE would lose that.


# Compare bit patterns, not values: for a 2-element LSE pack the high half
# of -inf bitcasts to a NaN payload, which never compares equal to itself.
assert torch.equal(actual.view(torch.int16), expected.view(torch.int16)), (

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

so is this correct? (i.e. using 32 bit numerics in the beginning and then for the comparison only converting them to 16 bit ones)

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

I think my comment was misleading, and I have rewritten it in db07f2ab.

Two different tensors are in play. The LSE is fp32 and stays fp32. The send buffer is in the output dtype (16-bit), and _dcp_a2a_lse_pack_dim returns 2 for any 16-bit output, so the fp32 LSE is bit-split across two of its slots. .view(torch.int16) bitcasts the send buffer — reinterpreting 16-bit storage as 16-bit integers. Nothing is narrowed.

Those slots are not meaningful floats at all; they are an opaque container for half of an fp32. A masked row stores -inf = 0xFF800000, and the high half 0xFF80 reads differently depending on the format:

bf16 (1|8|7):   1 | 11111111 | 0000000      exp all-ones, mantissa 0     -> -inf
fp16 (1|5|10):  1 | 11111    | 1110000000   exp all-ones, mantissa 896   -> NaN

bf16 is the top 16 bits of fp32 (same 8-bit exponent), so the truncated half of -inf is still -inf. fp16 has only a 5-bit exponent, so three of fp32's exponent bits spill into fp16's mantissa — and an all-ones exponent with a non-zero mantissa is NaN, which never compares equal to itself.

Bit equality is also strictly stronger than allclose, which is what a bit-exactness claim should be asserting.

rbrugaro-amd and others added 3 commits September 10, 2026 05:45
torch.cuda.is_available() is True on ROCm -- the PyTorch ROCm build maps HIP
onto the torch.cuda API -- so the guard was behaving correctly. But it is the
wrong predicate to express the intent, and the neighbouring one is a trap:
current_platform.is_cuda() is False on ROCm and would silently skip the whole
file. Use current_platform.is_cuda_alike(), which reads as intended and covers
both backends.

No behavioural change; 353 tests still pass on MI355X.

Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>
…k test

The comment claimed the high half of -inf 'bitcasts to a NaN payload' for any
2-element LSE pack. That holds for fp16 but not bf16, and it invited the
reading that the fp32 LSE is being narrowed to 16 bits for the comparison.

Neither is what happens. The LSE stays fp32; it is the send buffer -- already
in the 16-bit output dtype -- that is bitcast, because its LSE slots hold
halves of an fp32 rather than meaningful floats. bf16 shares fp32's 8-bit
exponent, so the truncated half is still -inf; fp16's 5-bit exponent pushes
three of those bits into its mantissa, and an all-ones exponent with a
non-zero mantissa is NaN, which never compares equal to itself.

Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>
@rbrugaro-amd rbrugaro-amd changed the title [DCP][Kernel][Perf]Fuse the empty-shard LSE mask into the A2A pack kernel [DCP][Kernel][Perf] Fuse the empty-shard LSE mask into the A2A pack kernel Sep 10, 2026
@AndreasKaratzas

Copy link
Copy Markdown
Member

/ci run

@github-actions

Copy link
Copy Markdown

✅ Triggered Buildkite CI #88094 for commit 1de164b1cb09.

@AndreasKaratzas

Copy link
Copy Markdown
Member

@dllehr-amd my comments have been addressed

@MatthewBonanni MatthewBonanni left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

LGTM, thanks!

@MatthewBonanni
MatthewBonanni merged commit 6ee5bb0 into vllm-project:main Sep 10, 2026
132 of 133 checks passed
ItsRoy69 pushed a commit to ItsRoy69/vllm that referenced this pull request Sep 15, 2026
…ernel (vllm-project#54889)

Signed-off-by: Rita Brugarolas Brufau <rita.brugarolasbrufau@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ready ONLY add when PR is ready to merge/full CI is needed

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants