fix(vllm): carry patch making ep_gather BLOCK_D divide hidden_size - #6498
Open
ezhong0211 wants to merge 2 commits into
Open
fix(vllm): carry patch making ep_gather BLOCK_D divide hidden_size#6498ezhong0211 wants to merge 2 commits into
ezhong0211 wants to merge 2 commits into
Conversation
ezhong0211
force-pushed
the
kimi-k3-vllm-ep-gather-patch
branch
from
August 6, 2026 19:56
e867af4 to
f21486f
Compare
vLLM cannot serve Kimi-K3 without this. In ep_gather, BLOCK_D is set to min(hidden_size, 1024), which is not necessarily a divisor of hidden_size, and the next line asserts that it is. K3's routed experts are 3584 wide and 3584 % 1024 == 512, so the assert trips during memory profiling and the server never reaches ready. It is pure shape arithmetic, so no parallelism, quantization or backend flag works around it. The fix halves BLOCK_D until it divides hidden_size. Halving is the right shape here rather than a gcd or an explicit table: BLOCK_D feeds tl.arange(0, BLOCK_D) in the kernel below, which requires a power of two, and halving preserves that while terminating at worst at 1, which divides everything. For 3584 it lands on 512 in one step. The grid already computes cdiv(hidden_size, BLOCK_D), so a smaller block just means more blocks along the hidden dimension -- the kernel itself is unchanged. Blast radius is nil for anything serving today. Every hidden size currently in use (1024/2048/4096/5120/7168) is a multiple of 1024, so BLOCK_D stays at 1024 and the loop never executes. Only non-multiples move. Note the loop also does not fire for hidden_size < 1024 that is not a power of two (e.g. 896), where BLOCK_D is already non-power-of-two -- a pre-existing upstream quirk this change neither introduces nor fixes. This is the first real user of the scripts/docker/vllm/amzn2023/patches/ extension point, which until now held only a .gitkeep. Verified with git apply --check against both refs the vllm configs pin -- d223c900 (ec2, sagemaker) and 3f5bd482 (hyperpod) -- and the patched file parses. The assert is confirmed still present at d223c900, so upstream has not fixed this. Carried rather than upstream-only because K3 is unservable until it lands upstream. It should be deleted once the upstream fix is in a pinned ref; tracking issue to be linked here. Worth flagging separately, not fixed in this commit: the Dockerfile's patch loop ends in `; true`, so a patch that fails to apply is swallowed and the build goes green with an unpatched image -- exactly the failure this patch exists to prevent. Verified by hand here, but that loop should fail hard. CI signal is limited to the patch applying and the image building. K3 itself cannot be tested: ~1.51 TiB of GPU memory against 640 GB on the largest runner. Note also that vLLM's K3 failure is recorded as inferred, not observed -- vllm:server-cuda-v2 was never pulled on the B300 -- so the assert is read from the source, not from a captured traceback.
ezhong0211
force-pushed
the
kimi-k3-vllm-ep-gather-patch
branch
from
August 6, 2026 20:09
f21486f to
89d98b6
Compare
The patch loop ended in `; true`, so `git apply` failing produced a green build with an unpatched image -- discoverable only by inspecting the loaded source. That is the worst possible failure mode for a carried patch, and VLLM_REF has already moved once mid-work, which is exactly what makes a patch go stale. Skip non-existent paths explicitly so the no-patches case still passes, and let a real apply failure stop the build. Also drops a dead .gitkeep test that could never match the *.patch glob, and adds --verbose so the log shows which hunks applied. Add the missing .gitkeep to prebuilt_wheels/ and sccache-cache/. Both are COPY sources and neither existed in git, since git cannot track an empty directory, so those COPYs fail on a clean clone -- they only worked here because earlier local runs had created the directories. The comment claiming the directory is empty by default is now actually true.
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.
Purpose
vLLM cannot serve Kimi-K3 without this.
In
ep_gather,BLOCK_Dis set tomin(hidden_size, 1024), which need not dividehidden_size, and the next line asserts that it does. K3's routed experts are 3584 wideand
3584 % 1024 == 512, so the assert trips during memory profiling and the server neverreaches ready. It is pure shape arithmetic, so no parallelism, quantization or backend flag
works around it.
The fix halves
BLOCK_Duntil it divideshidden_size. Halving is the right shape hererather than a gcd or a lookup table, because
BLOCK_Dfeedstl.arange(0, BLOCK_D)in thekernel below, which requires a power of two. Halving preserves that and terminates at worst
at 1, which divides everything. For 3584 it lands on 512 in one step. The grid already
computes
cdiv(hidden_size, BLOCK_D), so a smaller block just means more blocks along thehidden dimension and the kernel itself is unchanged.
Blast radius is nil for anything serving today. Every hidden size currently in use
(1024, 2048, 4096, 5120, 7168) is a multiple of 1024, so
BLOCK_Dstays at 1024 and theloop never executes. Only non-multiples move. Note the loop also does not fire for
hidden_size < 1024that is not a power of two, for example 896, whereBLOCK_Dis alreadynon-power-of-two. That is a pre-existing upstream quirk this change neither introduces nor
fixes.
This is the first real user of the
scripts/docker/vllm/amzn2023/patches/extension point,which until now held only a
.gitkeep.Carried rather than upstream-only because K3 is unservable until a fix lands upstream. It
should be deleted once the upstream fix is in a pinned ref.
Two things for the reviewer:
trigger. I will add the link before merge.
; true, so apatch that fails to apply is swallowed and the build goes green with an unpatched image.
That is exactly the failure mode this patch exists to prevent. Verified by hand here, but
that loop should fail hard.
Test Plan
Kimi-K3 itself cannot be tested in CI. It needs roughly 1.51 TiB of GPU memory against
640 GB on the largest runner. So verification is that the patch applies to every pinned ref
and produces valid code.
git apply --checkagainst everyvllm_refthe configs pin, so no variant silentlyskips the patch. Note the ec2 and sagemaker ref moved from
7b3d595etod223c900in vllm: bump AL2023 images to upstream commit d223c90 #6467 while this work was in progress, so this must be re-checked at merge time.
needed and has not been fixed under us.
ast.parsethe patched file.serving changes behavior.
Deliberately not attempted: serving K3 on this image. That requires B300, which is not in
the CI fleet.
Test Result
Local checks pass:
git apply --checkpasses atd223c900(ec2, sagemaker) and3f5bd482(hyperpod).assert hidden_size % BLOCK_D == 0is confirmed still present atd223c900, soupstream has not fixed this.
ast.parseon the patcheddeep_gemm_utils.pysucceeds.BLOCK_D = 1024with theloop never executing. 3584 moves to 512 in one step.
CI result pending; will update once the PR is open.
One caveat stated plainly: vLLM's K3 failure is recorded as inferred, not observed. The
vllm:server-cuda-v2image was never pulled on the B300 during verification, so the assertis read from the source rather than from a captured traceback. The code path is
unambiguous, but no one has watched this specific server die.
Toggle if you are merging into main Branch
PR Checklist
pre-commit run --all-fileslocally before creating this PR. (Read DEVELOPMENT.md for details).