Repository navigation
[ROCm][KV Offload] SimpleCPUOffload: batch copies on SDMA with hipMemcpyFlagExtPreferCE (HIP 7.16+) - #65559
Draft
benenzhu wants to merge 1 commit into
Draft
[ROCm][KV Offload] SimpleCPUOffload: batch copies on SDMA with hipMemcpyFlagExtPreferCE (HIP 7.16+)#65559benenzhu wants to merge 1 commit into
benenzhu wants to merge 1 commit into
Conversation
…cpyFlagExtPreferCE (HIP 7.16+) On HIP 7.16, hipMemcpyBatchAsync runs host <-> device batches as a blit kernel on the CUs, which slows the decode running next to the copies. Set hipMemcpyFlagExtPreferCE in the batch attributes on HIP >= 7.16 so the runtime hands the batch to the SDMA engines instead. Older runtimes keep flags 0. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Taoyu Zhu <zhutaoyu97@gmail.com>
ivanium
added a commit
to GirasoleY/vllm
that referenced
this pull request
Nov 5, 2026
…aries Allocate only the resolved CPU cache size and inherit the worker's NUMA memory policy. Keep pin/unpin paired per range and reuse HiSparse block-safe registration boundaries for chunked CPU offload pinning. Remove the independent HIP copy-engine optimization, tracked in vllm-project#65559. Validation: pin/unpin regression tests passed (2 cases); original registration leak reproduced; Ruff and diff checks passed. Registration-boundary test collection was blocked by missing FlashAttention CUDA extensions. GPU and end-to-end tests were not run. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
5 of 6 tasks
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Overview
vLLM's ROCm image now ships the HIP 7.16 runtime (
docker/Dockerfile.rocm_base,ROCM_RUNTIME_COMMIT25e14349f2). On it,hipMemcpyBatchAsyncrunsSimpleCPUOffloadConnector's host <-> device batches as a blit kernel on the CUs (__amd_rocclr_copyBufferBatch), so every load or store slows the decode that runs next to it. WithhipMemcpyFlagExtPreferCE(0x100) in the batch attributes, the runtime hands the whole batch to the SDMA engines instead and no kernel runs.build_paramsnow sets that flag on HIP >= 7.16 (_copy_flags_for_hip_version, next to the existingnumAttrsgate). Older runtimes keep flags 0, so nothing changes for them. CUDA is untouched.Claims
Validation
MI355X, one GPU, image
vllm/vllm-openai-rocm:nightly-40941f2c4ba96fba0a7e282a5355e89cd74d9ac5(main of 2026-11-03, HIP 7.16.26361). Before = the image's own vLLM (flags 0x0); after = this PR'scuda_mem_ops.pycopied over it (flags 0x100); same job, same GPU. The copies go throughbuild_params+copy_blocksexactly as the connector calls them, on one TP4 rank's KV layout of a 60-layer model: 131 per-layer tensors, 16 / 32 KiB pages (3.09 MiB per block), a registered host buffer, shuffled block ids, a low-priority stream. Decode is a CUDA graph of 32 bf16 GEMMs (M x 6144 x 4096 and back, 1.5 GiB of distinct weights, HBM-bound) replayed while N-block transfers loop; slowdown = per-replay time vs the same graph alone.Without the flag the slowdown is bimodal: some replays are clean and the rest (10-78% of them, depending on the row) run 1.2x to 5.6x slower, so some medians look fine while the means and tails do not. With the flag no replay is more than 1.2x slower.
pytest tests/v1/simple_kv_offload/test_hip_mem_ops.py: the newtest_copy_flags_for_hip_versioncovers the gate (unknown, ROCm 7.2, HIP 7.15, 7.16, 8.0); like the rest of that file it needs no HIP runtime.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.
AI assistance (Claude Code) was used to prepare this change.
🤖 Generated with Claude Code