Feature/Fix: support GLM for all engines mix and vllm moriio pd & adding all to test - #25
Merged
Merged
Conversation
dorado269
force-pushed
the
yihou.dev.vllm.glm
branch
from
July 23, 2026 08:35
4b214b5 to
77128d0
Compare
dorado269
marked this pull request as ready for review
July 23, 2026 11:09
dorado269
requested review from
JohnQinAMD,
jiejingzhangamd,
limou102 and
xiaobochen-amd
as code owners
July 23, 2026 11:09
Collaborator
|
we add DOC in release check, so please add sign off in git commit. |
patch_moriio_dsa_write.py, patch_moriio_hetero.py and patch_vllm_moriio_blocksize.py targeted the pre-v0.25.1 MoRIIO connector. v0.25.1 refactored that connector into moriio_layout.py, which natively handles all three concerns (full-layer wait_for_save, per-layer transfer geometry, kernel/logical block-size ratio), so the patch anchors no longer match and each one silently skips. They were dead weight producing misleading "skipped" build-log lines; drop them. Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> Signed-off-by: yihou <yihou@amd.com>
….1 PD)
v0.25.1's moriio_layout.py get_layer_transfer_geometry derives the MLA per-block
transfer size and stride from the KV tensor shape (block_size*latent*element_size)
instead of the authoritative spec.page_size_bytes. For block-scaled fp8 MLA caches
this is wrong in two independent ways:
- DeepSeek-V4 fp8_ds_mla is 576-byte-aligned/padded, so page_size_bytes exceeds
the shape-derived bytes; the dropped tail carries the UE8M0 per-block scale, so
decode dequantizes with a stale scale -> right-structure/wrong-fact output.
- GLM-5.1 lays the cache out per kernel block of size 1 while the scheduler pages
at block_size 16, so the shape-derived length is 16x too small -> only 1/16 of
each block's KV is moved -> total garbage.
Both cross-node MoRIIO PD paths produced garbage while prefill-direct was correct;
Mooncake was correct on the same nodes because it already transfers the full page.
patch_moriio_pagelen.py makes the MLA branch use spec.page_size_bytes for block_len
and page_size_bytes // element_size for block_stride, mirroring Mooncake. It is
idempotent, self-locating, and a no-op for contiguous fp16/bf16 K/V caches. Applied
by the existing Dockerfile.vllm patch loop.
Verified 2026-07-22 on TP4 2-node MoRIIO PD, temp=0: DeepSeek-V4-Pro and
GLM-5.1-FP8 both France->Paris / China->Beijing, coherent and deterministic.
Co-Authored-By: Claude Opus 4 <noreply@anthropic.com>
Signed-off-by: yihou <yihou@amd.com>
The vLLM PD-disagg suite only exercised Mooncake; the MoRIIO connector — and thus the moriio_layout.py MLA page-length fix — had no automated coverage. Add a single GLM-5.1-FP8 case over MoRIIO as the regression guard for that fix. The disagg adapter now selects the connector per-case: a case opts into MoRIIO via extra_env[INFERA_E2E_KV_CONNECTOR]="MoRIIOConnector", otherwise it defaults to Mooncake, so every existing case's argv/env is byte-for-byte unchanged. On the MoRIIO path the adapter emits the connector's extra_config (host_ip/proxy_ip/ proxy_ping/http/handshake/notify ports + tp_size, the router forges the request_id from the handshake/notify ports it reads from etcd) and sets MORI_IB_GID_INDEX; the selector key is stripped from the worker env. No other model gets a MoRIIO variant, and all disagg cases stay @pytest.mark.slow (skipped by default), so CI is unchanged. GlmMoeDsa (block-scaled fp8 MLA + DSA lightning indexer) is the exact cache layout the page-len bug corrupted; the case's launch args mirror the verified recipe (fp8 KV, aiter MoE, glm45 reasoning parser). Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> Signed-off-by: yihou <yihou@amd.com>
Extends the GLM-on-every-engine coverage: after GLM-5.1 on vLLM PD (MoRIIO), add a single-node SGLang mixed case. GlmMoeDsaForCausalLM routes through SGLang's DeepSeek MLA+DSA path, which auto-selects attention_backend=dsa / page_size=64 / tilelang / kv bf16, so the case is intentionally minimal (--reasoning-parser glm45 + --mem-fraction-static 0.85, SGLANG_USE_AITER=1) — the DSv4-specific flags would fight that auto-config. The existing SglangAdapter already passes these through, so no adapter change is needed. Verified 2026-07-23 single-node mix, temp=0: France->Paris, China->Beijing, 2+2->4. Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> Signed-off-by: yihou <yihou@amd.com>
Completes GLM coverage across all three engines. ATOM loads GlmMoeDsaForCausalLM (model_type glm_moe_dsa) natively — it allocates the MLA chunked-prefill workspaces for the GlmMoeDsa MLA+DSA path. The case is intentionally minimal: fp8 KV + --trust-remote-code, with HSA_NO_SCRATCH_RECLAIM=1. No --method mtp: GLM ships no MTP/nextn draft weights, and gfx950 plain decode is correct (the gfx942 broken-plain-decode bug that forced speculative decode on DSv4 does not reproduce). The existing AtomAdapter passes these through, so no adapter change is needed. Verified 2026-07-23 single-node, temp=0 (thinking disabled): France->Paris, China->Beijing, 2+2->4. Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> Signed-off-by: yihou <yihou@amd.com>
Fills the last engine/scenario gap: GLM-5.1 was already covered on vLLM cross-node PD (MoRIIO) and SGLang/ATOM single-node mix; this adds vLLM single-node PD-mix. vLLM v0.25.1 serves GlmMoeDsa via the DeepSeek MLA path; the case uses fp8 KV + aiter (via env) + the glm45 reasoning parser. Single-node mix uses no kv-transfer connector (that is the pd_disag path), so it does not exercise the MoRIIO page-len fix — it is a straight brings-up-and-is-correct guard. The existing VllmAdapter passes these flags through, so no adapter change is needed. Verified 2026-07-23 single-node mix, temp=0 (thinking disabled): France->Paris, China->Beijing, 2+2->4. Co-Authored-By: Claude Opus 4 <noreply@anthropic.com> Signed-off-by: yihou <yihou@amd.com>
dorado269
force-pushed
the
yihou.dev.vllm.glm
branch
from
July 27, 2026 03:24
6ff49e2 to
2df2fed
Compare
jiejingzhangamd
approved these changes
Jul 28, 2026
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.
Description
Fix the vLLM v0.25.1 MoRIIO PD-disaggregation corruption on block-scaled fp8 MLA models (DeepSeek-V4 / GLM-5.1), and add automated GLM-5.1 e2e coverage across all three engines (vLLM, SGLang, ATOM) so it can't regress.
Type of change
Changes
Remove three no-op MoRIIO patches —
patch_moriio_dsa_write.py,patch_moriio_hetero.py,patch_vllm_moriio_blocksize.pytargeted the pre-v0.25.1 connector. v0.25.1 refactored it intomoriio_layout.py, which natively covers all three concerns (full-layerwait_for_save, per-layer transfer geometry, kernel/logical block-size ratio), so their anchors no longer match and each silently skips. Dropped to remove dead weight + misleading "skipped" build-log lines.Fix the MLA page-length transfer geometry — the real bug.
get_layer_transfer_geometryderived the MLA per-block transfer size/stride from tensor shape (block_size*latent*element_size) instead of the authoritativespec.page_size_bytes. Two symptoms, one defect:fp8_ds_mlais 576-byte-aligned/padded →page_size_bytes> shape bytes; the dropped tail carries the UE8M0 per-block scale → decode dequantizes with a stale scale → right-structure/wrong-fact output ("France"→"a good idea").patch_moriio_pagelen.pyusesspec.page_size_bytesforblock_lenandpage_size_bytes // element_sizeforblock_stride, mirroring Mooncake (correct on the same nodes). Idempotent, self-locating, no-op for contiguous fp16/bf16 K/V.Add GLM-5.1-FP8 e2e cases across all engines/scenarios — the regression guards. For vLLM PD-disagg, the adapter now selects the KV connector per-case: a case opts into MoRIIO via
extra_env[INFERA_E2E_KV_CONNECTOR]="MoRIIOConnector", otherwise it defaults to Mooncake so every existing case is byte-for-byte unchanged. Only GLM declares the sentinel; no other model gets a MoRIIO variant. vLLM/SGLang/ATOM also get single-node PD-mixed GLM cases via their existing adapters (no adapter change). All e2e cases stay@pytest.mark.slow(skipped by default) → CI is unchanged.Engine / case coverage (GLM-5.1-FP8)
GLM-5.1-FP8 was brought up and verified correct at temp=0 (France→Paris, China→Beijing, 2+2→4) on every engine/scenario, and each is wired into the suite:
pd_disag/vllm→GLM-5.1-FP8-tp4pd_mixed/vllm→GLM-5.1-FP8-tp4pd_mixed/sglang→GLM-5.1-FP8-tp4pd_mixed/atom→GLM-5.1-FP8-tp4Existing (unchanged) transfer-engine coverage for context: vLLM
pd_disagalso runs Mooncake (default) onQwen3-0.6B-tp2; the MoRIIO path is new and GLM-only.Notes per engine:
GlmMoeDsaForCausalLMauto-selects the DSA attention path (attention_backend=dsa, tilelang); the case is minimal (--reasoning-parser glm45,SGLANG_USE_AITER=1) — forcing DSv4 flags would fight the auto-config.GlmMoeDsaForCausalLMnatively (MLA chunked-prefill workspaces); minimal--kv_cache_dtype fp8 --trust-remote-code. No MTP — GLM ships no MTP/nextn draft weights and gfx950 plain decode is correct (the gfx942 broken-plain-decode bug does not reproduce).Verification
docker build -f deploy/docker/Dockerfile.vllmsucceeds; build log showsmoriio-pagelen: patched(not skipped) and the three removed patches never appear; in-imagemoriio_layout.pycontains the fix marker.pytest --collect-only, offline):pd_disag/vllm→Qwen3-0.6B-tp2(existing, unchanged) +GLM-5.1-FP8-tp4(new MoRIIO). Adapter unit assertions confirm the Mooncake path argv/env is byte-for-byte unchanged and the MoRIIO path emits the correct extra_config with the sentinel stripped from worker env.pd_mixed/vllm→ existing 3 ids unchanged +GLM-5.1-FP8-tp4.pd_mixed/sglang→ existing 4 ids unchanged +GLM-5.1-FP8-tp4.pd_mixed/atom→ existing 4 ids unchanged +GLM-5.1-FP8-tp4.Test plan
pytest --collect-onlylists the new cases; existing ids / default Mooncake path unchangedINFERA_E2E_NODES=<p>,<d> tests/run_tests.sh e2e vllm disag→ GLM MoRIIO correctness probe PASS in-suite🤖 Generated with Claude Code