Skip to content

Feature/Fix: support GLM for all engines mix and vllm moriio pd & adding all to test - #25

Merged
jiejingzhangamd merged 6 commits into
mainfrom
yihou.dev.vllm.glm
Jul 28, 2026
Merged

Feature/Fix: support GLM for all engines mix and vllm moriio pd & adding all to test#25
jiejingzhangamd merged 6 commits into
mainfrom
yihou.dev.vllm.glm

Conversation

@dorado269

@dorado269 dorado269 commented Jul 23, 2026

Copy link
Copy Markdown
Collaborator

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

  • Documentation change (change only to the documentation, either a fix or a new content)
  • Bug fix (non-breaking change which fixes an issue)
  • New feature (non-breaking change which adds functionality)
  • Breaking change (fix or feature that would cause existing functionality to not work as expected)
  • Infra/Build change
  • Code refactoring

Changes

  1. Remove three no-op MoRIIO patchespatch_moriio_dsa_write.py, patch_moriio_hetero.py, patch_vllm_moriio_blocksize.py targeted the pre-v0.25.1 connector. v0.25.1 refactored it into moriio_layout.py, which natively covers all three concerns (full-layer wait_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.

  2. Fix the MLA page-length transfer geometry — the real bug. get_layer_transfer_geometry derived the MLA per-block transfer size/stride from tensor shape (block_size*latent*element_size) instead of the authoritative spec.page_size_bytes. Two symptoms, one defect:

    • DeepSeek-V4 fp8_ds_mla is 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").
    • GLM-5.1 lays the cache out per kernel block of size 1 while the scheduler pages at block_size 16 → shape-derived length is 16× too small → only 1/16 of each block's KV moved → total garbage ("is is is").

    patch_moriio_pagelen.py uses spec.page_size_bytes for block_len and page_size_bytes // element_size for block_stride, mirroring Mooncake (correct on the same nodes). Idempotent, self-locating, no-op for contiguous fp16/bf16 K/V.

  3. 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:

Engine Scenario KV transfer TP e2e case (test id) Live temp=0 verify
vLLM PD-disaggregated (cross-node) MoRIIO (this PR's fix) 4 pd_disag/vllmGLM-5.1-FP8-tp4 ✅ PASS (2-node)
vLLM PD-mixed (single-node) n/a (in-process) 4 pd_mixed/vllmGLM-5.1-FP8-tp4 ✅ PASS
SGLang PD-mixed (single-node) n/a (in-process) 4 pd_mixed/sglangGLM-5.1-FP8-tp4 ✅ PASS
ATOM PD-mixed (single-node) n/a (in-process) 4 pd_mixed/atomGLM-5.1-FP8-tp4 ✅ PASS

Existing (unchanged) transfer-engine coverage for context: vLLM pd_disag also runs Mooncake (default) on Qwen3-0.6B-tp2; the MoRIIO path is new and GLM-only.

Notes per engine:

  • vLLM/MoRIIO (pd_disag) is the only case exercising this PR's page-len fix (cross-node RDMA KV transfer). Garbage before the fix, correct after. The vLLM pd_mixed case uses no kv-transfer connector, so the fix is a no-op there — it's a straight brings-up-and-is-correct guard.
  • SGLang: GlmMoeDsaForCausalLM auto-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.
  • ATOM: loads GlmMoeDsaForCausalLM natively (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

  • Image build (cluster, MI355X node): docker build -f deploy/docker/Dockerfile.vllm succeeds; build log shows moriio-pagelen: patched (not skipped) and the three removed patches never appear; in-image moriio_layout.py contains the fix marker.
  • Fix correctness (TP4 2-node MoRIIO PD, temp=0): DeepSeek-V4-Pro and GLM-5.1-FP8 both France→Paris / China→Beijing, coherent and deterministic. Mooncake was correct on the same nodes throughout (differential reference).
  • Test wiring (pytest --collect-only, offline):
    • pd_disag/vllmQwen3-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

  • Image builds with the fix baked in via the Dockerfile patch loop
  • pytest --collect-only lists the new cases; existing ids / default Mooncake path unchanged
  • GLM-5.1-FP8 temp=0 correctness verified live on all four engine/scenario combos (vLLM MoRIIO 2-node, vLLM single-node mix, SGLang single-node, ATOM single-node)
  • (optional, needs 2 nodes) INFERA_E2E_NODES=<p>,<d> tests/run_tests.sh e2e vllm disag → GLM MoRIIO correctness probe PASS in-suite

🤖 Generated with Claude Code

@dorado269
dorado269 force-pushed the yihou.dev.vllm.glm branch from 4b214b5 to 77128d0 Compare July 23, 2026 08:35
@dorado269
dorado269 marked this pull request as ready for review July 23, 2026 11:09
@dorado269 dorado269 changed the title Yihou.dev.vllm.glm Feature/Fix: support GLM for all engines mix and vllm moriio pd & adding all to test Jul 23, 2026
@jiejingzhangamd

Copy link
Copy Markdown
Collaborator

we add DOC in release check, so please add sign off in git commit.

dorado269 and others added 6 commits July 27, 2026 03:23
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
dorado269 force-pushed the yihou.dev.vllm.glm branch from 6ff49e2 to 2df2fed Compare July 27, 2026 03:24
@jiejingzhangamd
jiejingzhangamd merged commit 1750e23 into main Jul 28, 2026
74 of 88 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants