diff --git a/deploy/docker/Dockerfile.sglang b/deploy/docker/Dockerfile.sglang index bc5d714ba..bd546d5a1 100644 --- a/deploy/docker/Dockerfile.sglang +++ b/deploy/docker/Dockerfile.sglang @@ -3,10 +3,24 @@ # the command. # # ⚠️ The base tag is PINNED and must stay pinned: the GLM-5.2 DSA patches near -# the end of this file are context diffs applied at --fuzz=0 against this exact +# the end of this file are context diffs applied at --fuzz=0 against a specific # sglang release. Bumping the base fails the build there rather than silently # mis-applying. Build with APPLY_SGLANG_DSA_PATCHES=0 if you need a newer base. -ARG SGLANG_BASE_IMAGE=lmsysorg/sglang:v0.5.17-rocm720-mi35x +# +# Moved v0.5.17 -> v0.5.18 for the GLM-5.3 series. GLM-5.3 and GLM-5.3-MXFP4 are +# `glm_moe_dsa` / `GlmMoeDsaForCausalLM`, field-for-field identical to GLM-5.2 +# except `transformers_version`, so the released engine already serves them +# through glm4_moe.py and this image needs no new model source. What the bump +# cost was one hunk: on v0.5.18 the DSA set still applies except for +# draft_cuda_graph_dp_vote's min()-reduce in dp_attn.py, which upstream renamed +# out from under it -- carried by an anchor port, see +# scripts/apply_sglang_dsa_patches.sh. +# +# GLM-5.3-**Flash** is NOT served by this image: it is `glm5_next`, which no +# published sglang release carries at all. That needs the build-time source +# overlay in Dockerfile.sglang.glm53, and deliberately stays out of this file so +# the default build path for Kimi-K3 and GLM-5.2 never fetches unreleased source. +ARG SGLANG_BASE_IMAGE=lmsysorg/sglang:v0.5.18-rocm720-mi35x FROM ${SGLANG_BASE_IMAGE} WORKDIR /opt/infera @@ -113,10 +127,19 @@ RUN if [ "${APPLY_SGLANG_DSA_PATCHES}" = "1" ]; then \ # prefill.py / mooncake/conn.py, which no DSA diff does. A script here exits non-zero # when its anchors drifted, i.e. when the fix did not go in — an engine image that # silently corrupts long prompts is worse than a failed build. +# Two of these anchors moved between v0.5.17 and v0.5.18 — the typing import in +# mooncake/conn.py grew `Set`, and serving_responses.py's `background=` argument +# gained `and not request.stream`. Neither fix landed upstream, so the patches +# are still needed; only their anchors are stale. reanchor_...glm53.sh retargets +# them into a COPY, leaving the originals untouched for the v0.5.17-based +# gfx942 image. It fails loudly if an old literal is gone, which would mean +# upstream moved again (or took the fix) and this stage should be re-derived. COPY deploy/docker/patches/sglang_disagg/ /tmp/sglang-disagg-patches/ +COPY deploy/docker/scripts/reanchor_sglang_disagg_glm53.sh /tmp/reanchor.sh RUN set -eu; \ - for f in /tmp/sglang-disagg-patches/*.py; do echo "[sglang-patch] $f"; python "$f"; done; \ - rm -rf /tmp/sglang-disagg-patches + bash /tmp/reanchor.sh /tmp/sglang-disagg-patches /tmp/sglang-disagg-v0518; \ + for f in /tmp/sglang-disagg-v0518/*.py; do echo "[sglang-patch] $f"; python "$f"; done; \ + rm -rf /tmp/sglang-disagg-patches /tmp/sglang-disagg-v0518 /tmp/reanchor.sh # ---- sglang Responses API patches (not PD-specific) ------------------------ # REQUIRED FOR ANY Codex-driven workload against a model with a custom chat diff --git a/deploy/docker/Dockerfile.sglang.glm53 b/deploy/docker/Dockerfile.sglang.glm53 new file mode 100644 index 000000000..7b6d23586 --- /dev/null +++ b/deploy/docker/Dockerfile.sglang.glm53 @@ -0,0 +1,354 @@ +# Infera SGLang engine image for GLM-5.3-Flash (glm5_next) on gfx950 / MI355X. +# +# WHY A SEPARATE FILE FROM Dockerfile.sglang: three things differ, and each one +# would break the other image if folded in. +# +# 1. Base moves v0.5.17 -> v0.5.18-rocm720-mi35x. No published sglang release +# carries GLM-5.3-Flash yet, so the base alone still cannot load it: the +# model is sglang PR #36507 and its ROCm gfx942/gfx950 enablement is the +# stacked PR #36607. #36507 is still OPEN against main; #36607 merged into +# #36507's branch on 2026-08-28, so neither is in main and no release +# carries the model. This image checks out #36607's head over +# the base's editable sglang tree, which is exactly how AMD ran the +# published GSM8K validation (gfx950 97.65%, 1288/1319). +# 2. The GLM-5.2 DSA patch set is skipped. Those four are --fuzz=0 diffs cut +# against v0.5.17 for glm_moe_dsa PD + DP-attention + EAGLE MTP. This +# recipe runs neither DP-attention nor MTP (AMD MTP is unvalidated for +# GLM-5.3 and does not start under PD at the current upstream cut), and +# the diffs cannot apply to this tree anyway. +# 3. Two patches/sglang_disagg/ anchors moved. They are re-anchored into a +# copy by scripts/reanchor_sglang_disagg_glm53.sh -- the originals stay +# untouched so the v0.5.17 Kimi-K3 image keeps building. +# +# Build from the REPO ROOT: +# docker build -f deploy/docker/Dockerfile.sglang.glm53 -t /infera:glm53- . +# +# If docker.io is unreachable from the build host ("failed to fetch oauth token: +# Post https://auth.docker.io/token: EOF"), pull the base through the cluster's +# Harbor pull-through cache instead -- same image, no Dockerfile edit: +# --build-arg SGLANG_BASE_IMAGE=harbor.crusoe.primus-safe.amd.com/proxy/lmsysorg/sglang:v0.5.18-rocm720-mi35x +# +# --------------------------------------------------------------------------- +# WHEN TO DELETE THIS FILE +# +# This whole file exists for one reason: glm5_next is the only model here that +# NO published sglang release carries. Every other model Infera serves already +# exists in the released engine and needs at most a bug fix from patches/ -- +# that is why Kimi-K3, MiniMax-M2 and DSv4 each added a couple of files under +# patches/ and no Dockerfile of their own, while this one needs a whole recipe. +# +# CHECK BEFORE TOUCHING ANYTHING BELOW. One command, no API token. The control +# probe is not optional -- without it a network failure reads as a 404: +# +# for ref in main v0.5.18; do +# for f in glm5_next glm4_moe; do # glm4_moe = control, must be 200 +# curl -s -o /dev/null -w "$ref $f %{http_code}\n" \ +# "https://raw.githubusercontent.com/sgl-project/sglang/$ref/python/sglang/srt/models/$f.py" +# done +# done +# +# Verified 2026-08-31: glm5_next is 404 on main, v0.5.18 and v0.5.17 while the +# control is 200 on all three; `Glm5Next` appears 0 times in main's +# model_config.py; v0.5.18 is still the newest tag; and lmsysorg/sglang has no +# mi35x image past v0.5.18-rocm720-mi35x (2026-08-21). Both PRs still open. +# +# ONCE glm5_next IS 200 ON A RELEASE TAG: delete this file. Build GLM-5.3 from +# Dockerfile.sglang with `--build-arg SGLANG_BASE_IMAGE=` plus +# `--build-arg APPLY_SGLANG_DSA_PATCHES=0`, and delete +# scripts/reanchor_sglang_disagg_glm53.sh and scripts/verify_glm53_overlay.py +# with it. Two things to re-derive on that tree first, because both were settled +# for THIS tree only: whether the two patches/sglang_disagg/ anchors still need +# re-anchoring (they may have landed upstream), and the two hicache ROCm gates +# behind the patch_hicache_rocm_staged_write_back.py exclusion below. +# --------------------------------------------------------------------------- +ARG SGLANG_BASE_IMAGE=lmsysorg/sglang:v0.5.18-rocm720-mi35x +FROM ${SGLANG_BASE_IMAGE} + +# ---- GLM-5.3-Flash source overlay ------------------------------------------ +# The base installs sglang editable from /sgl-workspace/sglang/python, and that +# directory is a full (non-shallow) clone of sgl-project/sglang, so the overlay +# is a fetch + checkout rather than a file copy. +# +# THE PIN IS DELIBERATE, and it is #36607's HEAD -- not its first commit. +# +# #36607 ("[AMD] Enable GLM-5.3-Flash on gfx942 and gfx950") is stacked on +# #36507's branch, not on main, and it MERGED 2026-08-28 (merge commit +# aa8c950a). Merged means its head c821c425 is frozen: pinning there is +# reproducible without being stale, which is the usual reason to pin early. +# +# Pinning EARLIER on that PR is the trap. Its first commit is 9e692c92, and +# five of the nine commits after it are the ones this image exists for: +# 654df43c [AMD] Support mixed Quark MXFP4 and block-FP8 loading +# bd1cc98b [AMD] Enable GLM shared-expert fusion on gfx95 +# 91c66b87 [AMD] Enable fused k-pool top-k on HIP +# 77a46694 [AMD] Enable AITER mHC for GLM on gfx95 +# Content, same commands across both refs: glm5_next.py 1834 -> 1942 lines, +# aiter references 3 -> 12, quantization/quark/quark.py 1103 -> 1172 lines. So +# 9e692c92 builds an engine that (a) silently takes the pre-mHC path, which +# upstream measured at 4.3-5.4x slower and which no log calls out, and (b) +# cannot load a Quark MXFP4 checkpoint at all. +# +# Do NOT "update" this to pull/36507/head. That branch carried the AMD work +# between 2026-08-28 and 2026-08-30, then LOST it in its 08-31 rebase ("fix: +# resolve CI regressions after GLM-5.3 Flash rebase"): as of 09-01 its head +# c767511e is back to aiter 3 / quark.py 1103. Tracking either PR by NAME is +# unreproducible anyway -- two builds of "the same Dockerfile" would ship two +# different engines. +# +# Known-good fallback if c821c425 misbehaves: 7fa1924c (2026-08-29, taken off +# #36507 while it still carried the AMD work; same quark.py 1172, aiter 11). +# That one has actually served this model on this host. Conversely, if a +# NON-MXFP4 feature misbehaves, read #36507's current content before assuming a +# bug -- it carries fixes this frozen branch does not. +# +# Bump the ref only against a new validation run. +# +# SGLANG_GLM53_PR only locates the ref cheaply, and the build below treats it +# that way: it fetches the PR ref, and if the pinned sha is still not present +# afterwards it fetches the sha directly (github serves fetch-by-sha, +# allowAnySHA1InWant). Both halves are load-bearing. A merged, closed or +# renumbered PR makes the first fetch fail outright; a force-push past the pin +# makes it *succeed* and still not contain the sha -- which is the exact +# reproducibility hazard this pin exists for, so it must not be a build failure +# either. Keep the two consistent when you bump, but the sha is the contract. +# +# Both PRs touch python/sglang/** only -- no C++ sgl-kernel rebuild. The new DSA +# kernels under python/sglang/kernels/jit/ are JIT-compiled at first use. +ARG SGLANG_GLM53_REF=c821c425c31b0e6c8151324b60fbc2857c39eaef +ARG SGLANG_GLM53_PR=36607 +# The ROCm base leaves this clone DIRTY: it swaps the HIP build metadata into +# place (python/pyproject_other.toml -> python/pyproject.toml, and the same for +# python/sglang/kernels/aot/), so those show as 2 modified + 2 deleted tracked +# files and a plain `git checkout` aborts with "local changes would be +# overwritten". Neither PR 36507 nor 36607 touches any of them -- both are +# confined to python/sglang/{srt,kernels/ops} plus tests -- so the safe move is +# to save that build metadata, force the checkout, and put it back. Force alone +# would silently restore the CUDA-flavoured pyproject and break any later +# `pip install` against this tree. +# +# `git diff HEAD`, not `git diff`: the latter compares the worktree to the +# INDEX, so a base image that staged its metadata swap (`git add`) lists +# nothing here, saves nothing, and loses the HIP pyproject to the forced +# checkout below -- while `git status --porcelain` on the line above still +# prints it, so the diagnostic and the logic would disagree about what is +# dirty. Today's base does not stage it; nothing makes that a promise. +COPY deploy/docker/scripts/verify_glm53_overlay.py /usr/local/bin/verify-glm53-overlay +RUN set -eux; \ + cd /sgl-workspace/sglang; \ + mkdir -p /tmp/base-meta/files; \ + git status --porcelain | tee /tmp/base-meta/status.txt; \ + git diff HEAD --name-only --diff-filter=M > /tmp/base-meta/modified.txt; \ + git diff HEAD --name-only --diff-filter=D > /tmp/base-meta/deleted.txt; \ + while read -r f; do mkdir -p "/tmp/base-meta/files/$(dirname "$f")"; cp "$f" "/tmp/base-meta/files/$f"; done < /tmp/base-meta/modified.txt; \ + git fetch --no-tags origin "pull/${SGLANG_GLM53_PR}/head" \ + || echo "[overlay] pull/${SGLANG_GLM53_PR}/head did not fetch; falling back to the sha"; \ + git cat-file -e "${SGLANG_GLM53_REF}^{commit}" 2>/dev/null \ + || git fetch --no-tags origin "${SGLANG_GLM53_REF}"; \ + git checkout --detach --force "${SGLANG_GLM53_REF}"; \ + while read -r f; do mkdir -p "$(dirname "$f")"; cp "/tmp/base-meta/files/$f" "$f"; done < /tmp/base-meta/modified.txt; \ + while read -r f; do rm -f "$f"; done < /tmp/base-meta/deleted.txt; \ + rm -rf /tmp/base-meta; \ + find python -name __pycache__ -type d -prune -exec rm -rf {} +; \ + test -f python/sglang/srt/models/glm5_next.py; \ + python3 /usr/local/bin/verify-glm53-overlay + +# ---- Mooncake: pinned upstream, no private source patches ------------------- +# Unchanged from Dockerfile.sglang -- same script, same upstream ref. The build +# is independent of the sglang Python tree above. +ARG BUILD_MOONCAKE=1 +ARG MOONCAKE_GIT_REF=faae8dd4a6309c3ecd47e0721a83b0250d686fa2 +COPY deploy/docker/scripts/build_mooncake_sglang.sh /tmp/build_mooncake_sglang.sh +RUN if [ "${BUILD_MOONCAKE}" = "1" ]; then \ + MOONCAKE_GIT_REF="${MOONCAKE_GIT_REF}" bash /tmp/build_mooncake_sglang.sh; \ + else \ + echo "BUILD_MOONCAKE=0 — leaving base Mooncake as-is"; \ + fi \ + && rm -f /tmp/build_mooncake_sglang.sh + +# ---- assert cross-host routing in the .so this image will LOAD -------------- +ARG REQUIRE_MOONCAKE_CROSS_HOST_ROUTING=1 +RUN set -eu; \ + if [ "${REQUIRE_MOONCAKE_CROSS_HOST_ROUTING}" != "1" ]; then \ + echo "skipping Mooncake cross-host routing assertion"; \ + else \ + so_dir="$(python3 -c 'import mooncake, os; print(os.path.dirname(mooncake.__file__))')"; \ + so="$(ls "$so_dir"/engine*.so | head -1)"; \ + syms="$(strings "$so")"; \ + printf '%s\n' "$syms" | grep -q "MC_DISABLE_HIP" || { \ + echo "ERROR: $so lacks upstream cross-host HIP locality routing" >&2; \ + exit 1; }; \ + echo "MOONCAKE_CROSS_HOST_ROUTING=present ($so)"; \ + fi + +# ---- libionic (host ionic ABI match) ----------------------------------------- +# Getting this wrong is SILENT. libibverbs prints one warning per device to +# stderr, ibv_get_device_list then returns 0 HCAs, mooncake logs "No RDMA +# devices found" + "Topology discovery complete. Found 0 HCAs." and serves KV +# over TcpTransport instead. Every pod stays Ready, every request succeeds, +# TTFT is just quietly bad. Verified in exactly that state on 2026-08-28. +# +# PICK BY THE HOST, NOT BY VERSION NUMBER. Newer is not higher-ABI: +# +# 54.0-149.g3304be71 abi 4..4 pool line 1.117.1-a-63 +# 54.0-187-1 abi 1..1 pool line 1.117.5-a-77 +# 54.0-192-1 abi 1..1 pool line 1.125.0-a-187 +# 54.0-197-1 abi 1..1 pool line 1.117.5-a-147 +# +# (measured off the ELF -- see deploy/docker/scripts/print_ionic_abi.py) +# +# To choose, read two things off a target node: +# cat /sys/class/infiniband/ionic_0/fw_ver -> e.g. 1.117.1-a-63 +# cat /sys/class/infiniband_verbs/uverbs1/abi_version -> e.g. 4 +# and take the deb from the pool directory whose name matches fw_ver. The +# repo path IS the pairing; it is not a "latest wins" archive. +# +# Default below targets the gfx950 fleet this image is built for +# (crsuse2-m2m-*, ionic_rdma 25.08.4.004, fw 1.117.1-a-63, uverbs abi 4). +# The MI355X/ionic-26.03 fleet that Dockerfile.sglang was written for is the +# OTHER case -- it wants abi 1, i.e. 54.0-187. Do not copy one into the other; +# that is precisely the regression this block replaces, and it cost a full +# benchmark run of TCP-transported KV before anyone noticed. +ARG INSTALL_LIBIONIC=1 +ARG LIBIONIC_DEB_URL=https://repo.radeon.com/amdainic/pensando/ubuntu/1.117.1-a-63/pool/main/r/rdma-core/libionic1_54.0-149.g3304be71_amd64.deb +# Build fails unless the provider that ends up installed declares this range. +# Set to 1 when building for the ionic-26.03 fleet; set empty to skip the gate. +# +# The gate and the download tolerance below are one decision, not two. A failed +# download leaves the base image's provider in place, which is the wrong ABI -- +# that is the whole reason the gate exists -- so "tolerate the download and then +# enforce the ABI" can only ever end in a hard failure two lines later, with the +# tolerance message printed above it to misdirect the reader. Runtime injection +# is a real workflow (an offline builder mounts the .deb on the node instead), +# but it is a deliberate choice, and the way to make it is to clear +# LIBIONIC_REQUIRE_ABI. So: the gate set means the download must succeed; the +# gate cleared means it may fail and the build carries on. +ARG LIBIONIC_REQUIRE_ABI=4 +COPY deploy/docker/scripts/print_ionic_abi.py /usr/local/bin/print-ionic-abi +RUN set -eu; \ + if [ "${INSTALL_LIBIONIC}" = "1" ]; then \ + if curl -4 -fsSL --retry 3 -o /tmp/libionic.deb "${LIBIONIC_DEB_URL}" \ + && dpkg -i --force-downgrade /tmp/libionic.deb \ + && ldconfig \ + && echo "libionic: $(dpkg -l | awk '/libionic1/{print $3}')" \ + && rm -f /tmp/libionic.deb; then \ + :; \ + elif [ -n "${LIBIONIC_REQUIRE_ABI}" ]; then \ + echo "ERROR: could not install ${LIBIONIC_DEB_URL} (offline build?)," >&2; \ + echo " and LIBIONIC_REQUIRE_ABI=${LIBIONIC_REQUIRE_ABI} demands that provider." >&2; \ + echo " To build without it and inject the .deb on the node at runtime," >&2; \ + echo " rebuild with --build-arg LIBIONIC_REQUIRE_ABI= (empty)." >&2; \ + exit 1; \ + else \ + echo "[libionic] install failed and no ABI is required — inject at runtime instead"; \ + fi; \ + else \ + echo "INSTALL_LIBIONIC=0 — using the base image's provider as-is"; \ + fi; \ + if [ -n "${LIBIONIC_REQUIRE_ABI}" ]; then \ + python3 /usr/local/bin/print-ionic-abi > /tmp/ionic-abi.txt || { \ + echo "ERROR: could not read the ionic provider's ABI range off the ELF." >&2; \ + echo " Unverified means unshipped: this gate exists because the wrong" >&2; \ + echo " provider fails open to TCP with nothing in any log." >&2; \ + exit 1; }; \ + cat /tmp/ionic-abi.txt; \ + grep -q "= ${LIBIONIC_REQUIRE_ABI}\.\.${LIBIONIC_REQUIRE_ABI}\$" /tmp/ionic-abi.txt || { \ + echo "ERROR: installed ionic provider does not declare abi ${LIBIONIC_REQUIRE_ABI}..${LIBIONIC_REQUIRE_ABI}." >&2; \ + echo " RDMA would silently fall back to TCP on the target nodes." >&2; \ + echo " Fix LIBIONIC_DEB_URL, or set LIBIONIC_REQUIRE_ABI to what those nodes export." >&2; \ + exit 1; }; \ + rm -f /tmp/ionic-abi.txt; \ + else \ + python3 /usr/local/bin/print-ionic-abi || \ + echo "[libionic] could not read the provider ABI; no gate set, continuing"; \ + fi + +# ---- infera ------------------------------------------------------------------ +WORKDIR /opt/infera +COPY pyproject.toml README.md ./ +COPY infera ./infera +RUN pip install --no-cache-dir ".[sglang]" "setuptools>=83.0.0" + +# ---- sglang PD patch (mooncake early-send KV wait event) -------------------- +# REQUIRED FOR CORRECTNESS: without it, chunked prefill hands every non-final +# chunk to the decode leg while the forward writing those pages is still +# running, so prompts longer than one chunk come back partially wrong with +# nothing in any log. Verified still absent at SGLANG_GLM53_REF (conn.py has no +# wait_event/synchronize() call), so this is not yet upstream. +# +# patch_responses_pd_bootstrap.py rides along: it plumbs bootstrap_host/port/room +# through /v1/responses so a PD pair accepts the request at all. Needed for any +# Codex-style driver; harmless otherwise. +# +# Both are re-anchored first; the helper exits non-zero if an anchor is gone. +COPY deploy/docker/patches/sglang_disagg/ /tmp/sglang-disagg-patches/ +COPY deploy/docker/scripts/reanchor_sglang_disagg_glm53.sh /tmp/reanchor.sh +RUN set -eu; \ + bash /tmp/reanchor.sh /tmp/sglang-disagg-patches /tmp/sglang-disagg-glm53; \ + for f in /tmp/sglang-disagg-glm53/*.py; do echo "[sglang-patch] $f"; python "$f"; done; \ + rm -rf /tmp/sglang-disagg-patches /tmp/sglang-disagg-glm53 /tmp/reanchor.sh + +# ---- sglang Responses API patch (not PD-specific) --------------------------- +# Anchors verified intact at SGLANG_GLM53_REF; applies unmodified. +COPY deploy/docker/patches/sglang_responses/ /tmp/sglang-responses-patches/ +RUN set -eu; \ + for f in /tmp/sglang-responses-patches/*.py; do echo "[sglang-patch] $f"; python "$f"; done; \ + rm -rf /tmp/sglang-responses-patches + +# ---- sglang ROCm hicache host allocator ------------------------------------- +# Only patch_hicache_rocm_host_alloc.py runs here; its anchors are intact. +# +# patch_hicache_rocm_staged_write_back.py is deliberately EXCLUDED. It resyncs +# two gates that disagreed about the staged write-back JIT on ROCm, and at +# SGLANG_GLM53_REF both gates in mem_cache/pool_host/mla.py already read +# `_is_cuda or _is_hip`, i.e. they agree. Its own precondition check cannot run +# either: `class DSAIndexerPoolHost` has moved out of mem_cache/memory_pool_host.py, +# so the script exits 1 rather than apply a gate it cannot verify -- correct +# behaviour, and the reason it is not in the loop. +# +# This is moot for the recipe below, which runs GPU-only. RE-DERIVE BOTH GATES +# BEFORE ENABLING --enable-hierarchical-cache OR kvd ON THIS IMAGE. +COPY deploy/docker/patches/sglang_rocm/patch_hicache_rocm_host_alloc.py /tmp/patch_hicache_rocm_host_alloc.py +RUN set -eu; \ + echo "[sglang-patch] patch_hicache_rocm_host_alloc.py"; \ + python /tmp/patch_hicache_rocm_host_alloc.py; \ + rm -f /tmp/patch_hicache_rocm_host_alloc.py + +# ---- Rust router (multi-core data plane; --router-backend rust) ------------- +# Unchanged from Dockerfile.sglang, including the same-RUN toolchain removal: +# image scanning reads the flattened filesystem, so a compiler deleted in a +# later layer is still reported. Keep the two assertions at the end in sync with +# Dockerfile.sglang -- the second one catches rustup's unguarded +# `. "$HOME/.cargo/env"` left behind in a shell rc, which the PATH check misses. +COPY rust ./rust +RUN set -eu; \ + command -v cc >/dev/null || { apt-get update && apt-get install -y --no-install-recommends build-essential && rm -rf /var/lib/apt/lists/*; }; \ + export LIBCLANG_PATH="$(dirname "$(find /opt/rocm* /usr/lib /usr/lib64 -name 'libclang.so*' 2>/dev/null | head -1)")"; \ + case "$LIBCLANG_PATH" in ""|".") apt-get update && apt-get install -y --no-install-recommends libclang-dev && rm -rf /var/lib/apt/lists/* && export LIBCLANG_PATH="$(dirname "$(find /usr/lib -name 'libclang.so*' 2>/dev/null | head -1)")";; esac; \ + command -v cargo >/dev/null || { curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --profile minimal; . "$HOME/.cargo/env"; }; \ + export PATH="$HOME/.cargo/bin:$PATH"; \ + export LIBCLANG_PATH="${LIBCLANG_PATH:-/opt/rocm/llvm/lib}"; \ + ( cd rust && cargo build --release --bin infera-router ); \ + cp rust/target/release/infera-router /usr/local/bin/infera-router; \ + rm -rf rust "$HOME/.cargo" "$HOME/.rustup" /usr/local/cargo /usr/local/rustup; \ + rm -f /usr/local/bin/cargo /usr/local/bin/rustc /usr/local/bin/rustup /usr/local/bin/rustdoc; \ + sed -i '/\.cargo\/env/d' "$HOME/.profile" "$HOME/.bashrc" 2>/dev/null || true; \ + ! command -v cargo >/dev/null || { echo "cargo still on PATH after cleanup: $(command -v cargo)" >&2; exit 1; }; \ + ! bash -lc true 2>&1 | grep -q cargo || { echo "a login shell still references cargo" >&2; exit 1; }; \ + /usr/local/bin/infera-router --help >/dev/null 2>&1 || { echo "infera-router unusable after cleanup" >&2; exit 1; } + +# ---- final assertion: this image can actually resolve GLM-5.3-Flash --------- +# Cheap, and it turns "the overlay got clobbered by a later layer" into a build +# failure instead of a CrashLoopBackOff 20 minutes into a weight load. +# +# Deliberately NOT a ModelRegistry.get_supported_archs() check: registry.py +# imports every model module with strict=False and logs+swallows failures, so on +# a GPU-less builder (aiter -> rocminfo) glm5_next would be reported "not +# registered" for a reason that has nothing to do with this image. The verifier +# checks the same contract -- module present, class defined, exported via +# EntryClass, compiles -- statically, and only then tries the real import. +RUN set -eu; python3 /usr/local/bin/verify-glm53-overlay + +COPY deploy/docker/scripts/infera_inject_host_ionic.sh /usr/local/bin/infera-inject-host-ionic + +ENTRYPOINT ["/usr/local/bin/infera-inject-host-ionic"] +CMD ["/bin/bash"] diff --git a/deploy/docker/patch.upstream.status.md b/deploy/docker/patch.upstream.status.md index 4ae904116..8df604bed 100644 --- a/deploy/docker/patch.upstream.status.md +++ b/deploy/docker/patch.upstream.status.md @@ -137,6 +137,66 @@ so expect upstream to close this differently than we did. > V4 stack runs on this branch, so that gate would be untested — the script's SCOPE > section records it for whoever gets there. +## sglang carried — `patches/sglang_carried/` (**baked by nothing**) + +Held, reviewable, and deliberately not built. No Dockerfile copies or runs this +directory, which is why these scripts are not in `sglang_rocm/` beside their +siblings: every other patch directory is consumed by an unconditional +`for f in .../*.py; do python "$f"; done`, so a file placed there is wired into +every image that copies it, with no per-file opt-out. +`patches/sglang_carried/README.md` carries the rule and the exit criteria. + +| patch | fixes | upstream issue | upstream PR | ours? | PR state | +|---|---|---|---|---|---| +| `sglang_carried/patch_glm_moe_gate_bias_fp32.py` | `MoEGate.__init__` allocates the MoE `e_score_correction_bias` as **bf16** whenever a quant_config is present and `_use_aiter`, and `biased_grouped_topk_gpu` casts it down again at the aiter call. GLM's bias is a narrow band at a large offset, so bf16 cannot hold it: measured on the checkpoints themselves, GLM-5.3 / -MXFP4 collapse **238 distinct fp32 values to 8** and GLM-5.3-Flash-MXFP4 **282 to 11**, which reorders `noaux_tc` top-k routing | none found | [sglang#37133](https://github.com/sgl-project/sglang/pull/37133) `[GLM-5.2] Keep GlmMoeDsa MoE e_score_correction_bias in fp32` (`xiaobochen-amd`) — same defect, **narrower gate**: see below | no (`xiaobochen-amd`) | OPEN | + +**Status: carried, not applied.** Verified as text and as predicate logic only — +applies to copies of `v0.5.18`, `c821c425` and the live source extracted from +`infera/engine-sglang:v0518-glm53` (result byte-identical to the patched +`v0.5.18` tree), idempotent on re-run, and both drift cases exit 1 having written +neither file. **It has never been executed on a GPU, and no accuracy or +throughput delta has been measured.** Do not read the row above as validation. + +Not wired because the defect is a routing perturbation, not a crash: the server +starts, answers, and reports plausible numbers either way. Rebuilding the engine +image while an alignment campaign is in flight would make our arm the only one +running the fixed router and destroy the like-for-like property those ratios rest +on. The exit criterion is a hardware run plus a measurement of the thing it +claims to change. + +> **Why the gate is not upstream's.** #37133 gates on +> `any("GlmMoeDsa" in arch for arch in config.architectures)`. That predicate is +> False for the whole Flash family and would ship a fix that does not fix it: +> `Glm5NextForConditionalGeneration.__init__` does +> `self.config = config.text_config` and passes the **text** config down, and +> GLM-5.3-Flash-MXFP4's `text_config` has `architectures: None` — there is no arch +> string on the object `MoEGate` receives, so widening the string list cannot +> help. Flash imports `DeepseekV2MoE` (hence the same `MoEGate`) and has the same +> collapse. +> +> This script's `_moe_gate_bias_wants_fp32()` is a union of three tests: +> `moe_router_dtype == "float32"` (the model author's own declaration — carried by +> all four GLM-5.3 configs and GLM-5.2-FP8, and **sglang has zero references to +> the field**, which is the actual root cause); `model_type` (`glm_moe_dsa` / +> `glm5_next*`, needed because GLM-5.2-MXFP4 and GLM-5.1-FP8 predate that field, +> and it survives the NextN rewrite in `configs/model_config.py:623`); and +> upstream's arch test, kept so this stays a superset of #37133 and becomes a +> no-op once it lands. Checked against all 38 checkpoints on this host: 8/8 GLM +> keep fp32, 30/30 non-GLM byte-identical. +> +> **Both edits or neither, and that is not stylistic.** aiter's launcher +> (`csrc/kernels/topk_softmax_kernels_group.cu:1156`) dispatches on +> `gating_output.dtype()` and then `reinterpret_cast`s the bias pointer to that +> same `scalar_t` **without checking the bias tensor's own dtype**. An fp32 bias +> under a bf16 gating tensor is neither an error nor a cast — it reads fp32 bytes +> as bf16. So applying the `deepseek_v2.py` half alone is strictly worse than the +> defect, and the script plans both files before writing either. +> +> **Drop this patch** when a base sglang keeps the GLM bias fp32 on both sides; +> the replacement text is then already present and the script reports "already +> present" and no-ops. If only #37133's form lands, keep it — the Flash half would +> go with it. + ## Mooncake C++ — `patches/mooncake_cpp/` SGLang now builds Mooncake `faae8dd4` directly and carries no private Mooncake diff --git a/deploy/docker/patches/sglang_carried/README.md b/deploy/docker/patches/sglang_carried/README.md new file mode 100644 index 000000000..21b3d2dfb --- /dev/null +++ b/deploy/docker/patches/sglang_carried/README.md @@ -0,0 +1,53 @@ +# `patches/sglang_carried/` — patches we hold, and deliberately do not build + +**No Dockerfile copies or runs anything in this directory.** That is the whole +point of it, and it is the reason these scripts do not live beside their +siblings in `patches/sglang_rocm/`. + +Every other patch directory is consumed the same way: + +```dockerfile +COPY deploy/docker/patches/sglang_rocm/ /tmp/sglang-rocm-patches/ +RUN set -eu; \ + for f in /tmp/sglang-rocm-patches/*.py; do echo "[sglang-patch] $f"; python "$f"; done; \ +``` + +The glob is `*.py` and the loop is unconditional, so **dropping a file into one +of those directories wires it into every image that copies the directory** — +`Dockerfile.sglang` and `Dockerfile.sglang.gfx942` for `sglang_rocm/`. There is +no per-file opt-out. A patch that is ready to read but not ready to build +therefore cannot be stored there without changing two images. + +This directory is the opt-out. Nothing globs it, so a script here is carried, +reviewable, and inert. + +## When a patch belongs here + +When all three hold: + +- the defect is established well enough to write the fix down, +- the fix has **not** been executed on hardware, or its effect has not been + measured, and +- wiring it would perturb something in flight — a benchmark campaign, an + alignment comparison, a packup whose numbers are already reported. + +The last one is the common case, and it is why "carried" is a status and not a +euphemism for "unfinished". A silently changed image invalidates every number +measured before and after it; an unwired script costs nothing and loses nothing. + +## Moving one out + +Move the file into the directory of the image that should carry it +(`sglang_rocm/`, `sglang_dsa/`, …) and update its row in +`../../patch.upstream.status.md` from *carried, not applied* to the image list +it is now baked by. The glob picks it up from there — no Dockerfile edit is +needed, which is exactly why the move must be deliberate. + +Before moving one out, it needs what it did not have when it was parked: a run +on hardware, and a measurement of the thing it claims to change. + +## Contents + +| script | what it fixes | why it is not wired | +|---|---|---| +| `patch_glm_moe_gate_bias_fp32.py` | GLM MoE `e_score_correction_bias` is allocated bf16 under aiter+quant and downcast again at the aiter router boundary; GLM's bias band collapses from 238 distinct fp32 values to 8 in bf16 (Flash: 282 → 11), reordering `noaux_tc` top-k routing | verified as text and as predicate logic only — **never executed on a GPU**, and no accuracy or throughput delta measured. Wiring it would rebuild the engine image mid-campaign and confound an alignment comparison whose validity rests on both arms running the same code | diff --git a/deploy/docker/patches/sglang_carried/patch_glm_moe_gate_bias_fp32.py b/deploy/docker/patches/sglang_carried/patch_glm_moe_gate_bias_fp32.py new file mode 100644 index 000000000..7e7711655 --- /dev/null +++ b/deploy/docker/patches/sglang_carried/patch_glm_moe_gate_bias_fp32.py @@ -0,0 +1,380 @@ +#!/usr/bin/env python3 +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +"""GLM MoE gate: keep `e_score_correction_bias` in fp32 end to end. + +WHAT: on ROCm with aiter, `MoEGate.__init__` allocates the MoE +`e_score_correction_bias` parameter as **bf16** whenever a quant_config is +present, and `biased_grouped_topk_gpu` casts it down again at the aiter call. +GLM's bias is a narrow band at a large offset, so bf16 cannot represent it: the +router ends up choosing among hundreds of experts using a bias with single-digit +distinct levels. This patch keeps the parameter fp32 for GLM MoE models and, at +the aiter boundary, promotes the gating logits to fp32 instead of demoting the +bias. + +WHY: measured on this host, reading the checkpoints' own safetensors -- + + checkpoint tensor range fp32 bf16 + GLM-5.3-MXFP4 L10 mlp.gate.e_score_correction_bias 6.817 - 7.063 238 8 + GLM-5.3 L10 (byte-identical values) 6.817 - 7.063 238 8 + GLM-5.3-Flash-MXFP4 L10 [288 experts] 6.179 - 6.564 281 10 + GLM-5.3-Flash-MXFP4 L11 [288 experts] 6.772 - 7.167 282 11 + + bf16's ULP at 7.0 is 0.03125 and the whole spread is ~0.25, so 238 distinct + fp32 biases collapse into 8 bf16 bins (Flash: 282 -> 11). `noaux_tc` selects + experts by `sigmoid(logits) + correction_bias`, so this is a routing + perturbation. Upstream (sgl-project/sglang#37133) measured **98.50 % of tokens + select a different top-8 expert set** on GLM-5.2, and states plainly that + accuracy benchmarks cannot resolve it (GSM8K 0.941 -> 0.947, GPQA-D 0.8333 -> + 0.8182, both inside the error bars). Those are UPSTREAM's numbers on GLM-5.2 / + MI355X. We have measured the dtype collapse; we have NOT measured a quality + delta, and nothing here should be read as claiming one. + + This is a correctness fix and it COSTS a little throughput: upstream reports a + flat +4.5 us per gating call, <= 1.2 % of a decode step. Do not expect tok/s. + + The trigger is live in both our images, verified link by link: + 1. `SGLANG_USE_AITER=1` is baked into the image environment, so + `_use_aiter = get_bool_env_var("SGLANG_USE_AITER") and _is_hip` is True. + 2. `"quark"` is in the downcast tuple and we launch `--quantization quark`. + 3. `DeepseekV2MoE` forwards the model-level quant_config into `MoEGate` + unchanged, so `quant_config is not None` holds. + 4. On HIP the flashinfer and `_is_cuda` arms of `biased_grouped_topk_gpu` + are skipped, so we land in the `elif _use_aiter:` arm that casts down. + +HOW: two edits that MUST land together (see "BOTH OR NEITHER" below). + + A. `models/deepseek_v2.py` -- add `_moe_gate_bias_wants_fp32(config)` and use + it to skip the bf16 downcast in `MoEGate.__init__`. + B. `layers/moe/topk.py` -- in `biased_grouped_topk_gpu`'s aiter arm, when the + bias arrives fp32, pass it through and promote `gating_output` to fp32; + otherwise keep today's downcast byte-for-byte. + +BOTH OR NEITHER -- why edit B is not optional polish. aiter's kernel launcher +is: + + VLLM_DISPATCH_FLOATING_TYPES_rmTorch(gating_output.dtype(), ..., [&] { + ... reinterpret_cast(gating_output.data_ptr()), + reinterpret_cast(correction_bias.data_ptr()), ... + + (`csrc/kernels/topk_softmax_kernels_group.cu`, read in our own image.) It + dispatches on the GATING dtype only and then **reinterpret_casts the bias to + that same scalar_t with no check of the bias tensor's own dtype**. Handing it + an fp32 bias alongside a bf16 gating tensor is not an error and not a cast -- + it reads fp32 bytes as bf16 and returns silent garbage. So applying A without + B is strictly WORSE than the bug it fixes. That is the concrete reason this + script is all-or-nothing rather than two independent edits. + + The converse direction is safe: `AITER_DTYPE_fp32` is in the dispatch list, so + promoting the gating tensor is a supported instantiation. + +WHY THE GATE IS NOT UPSTREAM'S. #37133 gates on `any("GlmMoeDsa" in arch for +arch in config.architectures)`. That predicate is False for the Flash family and +would leave it broken: + + * `Glm5NextForConditionalGeneration.__init__` does `self.config = + config.text_config` and passes the TEXT config down, so the object `MoEGate` + receives is `text_config` -- and GLM-5.3-Flash-MXFP4's `text_config` has + **`architectures: None`**. There is no arch string to match; widening the + string list does not help. + * Flash imports `DeepseekV2MoE as Glm5NextMoE`, i.e. the very same `MoEGate`, + and has the same collapse (282 -> 11 above). + + So the predicate here is a union of three tests, cheapest first: + + 1. `moe_router_dtype == "float32"` -- the model author's own declaration. + Present on GLM-5.3, GLM-5.3-MXFP4, GLM-5.3-Flash, GLM-5.3-Flash-MXFP4 and + GLM-5.2-FP8-fixed. sglang v0.5.18 ignores the field entirely (zero + references under `python/sglang/srt/`), which is the actual root cause. + 2. `model_type` -- `glm_moe_dsa` (5.1/5.2/5.3 big) or `glm5_next*` (Flash, + whose text config reports `glm5_next_text`). Needed because GLM-5.2-MXFP4 + and GLM-5.1-FP8 predate the `moe_router_dtype` field. `model_type` + survives the NextN rewrite in `configs/model_config.py`, so draft heads + are covered without a substring hack. + 3. upstream's arch-string test, kept so this stays a superset of #37133 and + is a no-op once that lands. + + BLAST RADIUS, measured rather than assumed: of 38 checkpoints on this host, + exactly 5 carry `moe_router_dtype: "float32"` and all 5 are GLM 5.2/5.3. Every + DeepSeek-V3/V4/R1, Kimi-K3, Qwen2.5/3/3.5, MiniMax-M3, Hy3, MiMo, Llama and + Mistral config has it unset and a non-GLM `model_type`, so all of them take + the unchanged branch and stay byte-identical. + +CONTEXT + upstream sgl-project/sglang#37133 "[GLM-5.2] Keep GlmMoeDsa MoE + e_score_correction_bias in fp32" (xiaobochen-amd), OPEN against + main, not merged. Its CI is red only at `pr-gate` / + `*-finish` aggregators -- the missing-`run-ci`-label failure -- + so no AMD job has actually run, and the AMD runners are + `linux-mi300-*` (gfx942), not our gfx950. Treat it as unvalidated + on our hardware. + bases The `MoEGate` dtype block is textually IDENTICAL on v0.5.18 and + c821c425, so edit A has one anchor. `topk.py` drifted -- v0.5.18 + has `correction_bias.to(dtype=gating_output.dtype)` inline where + c821c425 hoisted a `bias` local and a `scaling` local -- so edit + B carries one alternative per base and requires that EXACTLY ONE + of them match. + why a script Upstream's diff FAILS `git apply --check` on v0.5.18 (the + `topk.py` hunk) and its arch predicate is wrong for Flash. Both + bases are pinned and neither can move: c821c425 is a + merged-and-frozen ref whose upstream branch was reverted + wholesale on 2026-09-01, so there is no newer ref to bump to. + Anchoring on source text serves both bases from one file. + NOT covered `fused_topk()` has a second `aiter_biased_grouped_topk` call + with the same downcast. It is the ungrouped path and + `topk_method == "noaux_tc"` never reaches it, so it is left + alone deliberately -- same as upstream. + precedent `router_dtype` is an established transformers config field + (switch_transformers, nllb_moe) defaulting to `"float32"` and + honoured as `getattr(torch, config.router_dtype)`. GLM's + `moe_router_dtype` is the same idea under a namespaced name. + +Self-locating and idempotent. All edits or none, across BOTH files: an anchor +that is missing, ambiguous, or matches in more than one variant writes NOTHING +and exits 1 -- because a half-applied fix here is not a degraded fix, it is a +silently mis-typed kernel argument. +""" + +from __future__ import annotations + +import importlib.util +import os +import sys +from pathlib import Path + +_TAG = "[glm-moe-gate-bias-fp32]" + +# -------------------------------------------------------------------------- +# Edit A -- models/deepseek_v2.py +# -------------------------------------------------------------------------- + +_A_REL = "models/deepseek_v2.py" + +# The helper. Inserted immediately above MoEGate, which is where the only caller +# is; `class MoEGate(nn.Module):` occurs exactly once on both bases. +_A_HELPER_ANCHOR = "class MoEGate(nn.Module):\n" + +_A_HELPER = '''def _moe_gate_bias_wants_fp32(config) -> bool: + """True when this model's MoE gate bias must stay fp32. + + GLM stores `e_score_correction_bias` as a narrow band at a large offset + (~6.2-7.2), where bf16's ULP of 0.03125 collapses hundreds of distinct + biases into single digits and reorders `noaux_tc` top-k routing. + + Three tests, cheapest first, because no single one covers every GLM + checkpoint we serve: + 1. `moe_router_dtype` -- the model author's own declaration. Carried by + GLM-5.3 / -MXFP4 / -Flash / -Flash-MXFP4 and GLM-5.2-FP8. + 2. `model_type` -- GLM-5.1/5.2 predate that field, and the Flash family + hands MoEGate its `text_config`, whose `architectures` is None, so the + arch test below cannot see it. `model_type` is always present and + survives the NextN rewrite in configs/model_config.py. + 3. architectures -- upstream sgl-project/sglang#37133's predicate, kept so + this stays a superset of it and becomes a no-op once it lands. + """ + router_dtype = getattr(config, "moe_router_dtype", None) + if isinstance(router_dtype, str) and router_dtype.lower() in ("float32", "fp32"): + return True + model_type = getattr(config, "model_type", None) or "" + if model_type.startswith("glm_moe_dsa") or model_type.startswith("glm5_next"): + return True + return any("GlmMoeDsa" in arch for arch in (getattr(config, "architectures", None) or [])) + + +''' + +# The downcast itself. Identical text on v0.5.18 and c821c425. +_A_GATE_ANCHOR = """ if config.topk_method == "noaux_tc" and not is_hash_moe: + correction_bias_dtype = torch.float32 + if quant_config is not None: + if _use_aiter and quant_config.get_name() in ( + "fp8", + "compressed_tensors", + "quark", + ): + correction_bias_dtype = torch.bfloat16 +""" + +_A_GATE_PATCHED = """ if config.topk_method == "noaux_tc" and not is_hash_moe: + correction_bias_dtype = torch.float32 + # GLM53_BIASFP32: GLM's bias sits in a ~0.25-wide band around +7, + # where bf16 steps by 0.03125 -- 238 distinct biases become 8 (Flash: + # 282 -> 11), reordering top-k. HF stores it fp32; keep it fp32. + # The matching promote lives at the aiter boundary in + # layers/moe/topk.py and is NOT optional: aiter reinterpret_casts the + # bias to the GATING tensor's dtype without checking it. + if quant_config is not None and not _moe_gate_bias_wants_fp32(config): + if _use_aiter and quant_config.get_name() in ( + "fp8", + "compressed_tensors", + "quark", + ): + correction_bias_dtype = torch.bfloat16 +""" + +# -------------------------------------------------------------------------- +# Edit B -- layers/moe/topk.py, biased_grouped_topk_gpu's aiter arm +# -------------------------------------------------------------------------- + +_B_REL = "layers/moe/topk.py" + +_B_PREAMBLE = """ # GLM53_BIASFP32: do NOT re-downcast an fp32 bias here. aiter dispatches + # on gating_output.dtype and then reinterpret_casts the bias pointer to + # that same scalar_t with no check of its own dtype, so an fp32 bias + # under a bf16 gating tensor is read as garbage rather than rejected. + # Promote the gating logits instead. Gated on the BIAS dtype, so a bias + # that already arrives bf16 is byte-identical to before. + if correction_bias.dtype == torch.float32: + _glm_gating = gating_output.to(torch.float32) + _glm_bias = correction_bias + else: + _glm_gating = gating_output + _glm_bias = {fallback} +""" + +# One variant per base. EXACTLY ONE must match; two matches means the file is +# not what either base says it is. +_B_VARIANTS: list[tuple[str, str, str]] = [ + ( + "v0.5.18 (inline downcast)", + """ topk_ids = torch.empty((token, topk), dtype=torch.int32, device=device) + aiter_biased_grouped_topk( + gating_output, + correction_bias.to(dtype=gating_output.dtype), + topk_weights, + topk_ids, + num_expert_group, +""", + """ topk_ids = torch.empty((token, topk), dtype=torch.int32, device=device) +""" + + _B_PREAMBLE.format(fallback="correction_bias.to(dtype=gating_output.dtype)") + + """ aiter_biased_grouped_topk( + _glm_gating, + _glm_bias, + topk_weights, + topk_ids, + num_expert_group, +""", + ), + ( + "c821c425 (hoisted `bias` local)", + """ topk_ids = torch.empty((token, topk), dtype=torch.int32, device=device) + aiter_biased_grouped_topk( + gating_output, + bias, + topk_weights, + topk_ids, + num_expert_group, +""", + """ topk_ids = torch.empty((token, topk), dtype=torch.int32, device=device) +""" + + _B_PREAMBLE.format(fallback="bias") + + """ aiter_biased_grouped_topk( + _glm_gating, + _glm_bias, + topk_weights, + topk_ids, + num_expert_group, +""", + ), +] + +# Present iff each edit landed. Used for idempotency and for the build-time +# bytecode check, the same way patch 01 greps for `_p1v2_trim`. +_A_MARKER = "not _moe_gate_bias_wants_fp32(config)" +_B_MARKER = "_glm_gating = gating_output.to(torch.float32)" + + +def _srt_dir() -> Path | None: + """Locate `sglang/srt`, preferring the interpreter's own sglang.""" + spec = importlib.util.find_spec("sglang") + if spec and spec.origin: + d = Path(spec.origin).parent / "srt" + if d.is_dir(): + return d + root = os.environ.get("SGLANG_DIR", "/sgl-workspace/sglang") + d = Path(root) / "python" / "sglang" / "srt" + return d if d.is_dir() else None + + +def _fail(msg: str) -> int: + print(f"{_TAG} {msg}", file=sys.stderr) + print(f"{_TAG} sglang drifted — re-anchor the patch, nothing written", file=sys.stderr) + return 1 + + +def _plan_a(path: Path) -> str | int: + """Return the new text for deepseek_v2.py, or an exit code.""" + src = path.read_text() + if _A_MARKER in src: + return src + + n = src.count(_A_GATE_ANCHOR) + if n != 1: + return _fail( + f"{_A_REL}: MoEGate correction-bias dtype block matched {n} times (want 1)" + ) + n = src.count(_A_HELPER_ANCHOR) + if n != 1: + return _fail(f"{_A_REL}: 'class MoEGate(nn.Module):' matched {n} times (want 1)") + + out = src.replace(_A_HELPER_ANCHOR, _A_HELPER + _A_HELPER_ANCHOR, 1) + return out.replace(_A_GATE_ANCHOR, _A_GATE_PATCHED, 1) + + +def _plan_b(path: Path) -> str | int: + """Return the new text for topk.py, or an exit code.""" + src = path.read_text() + if _B_MARKER in src: + return src + + hits = [(name, old, new) for name, old, new in _B_VARIANTS if src.count(old) == 1] + ambiguous = [name for name, old, _ in _B_VARIANTS if src.count(old) > 1] + if ambiguous: + return _fail(f"{_B_REL}: anchor is ambiguous for variant(s) {ambiguous}") + if len(hits) != 1: + names = [name for name, _, _ in _B_VARIANTS] + return _fail( + f"{_B_REL}: {len(hits)} of {len(names)} aiter-call variants matched " + f"(want exactly 1); tried {names}" + ) + + name, old, new = hits[0] + print(f"{_TAG} {_B_REL}: matched variant {name}") + return src.replace(old, new, 1) + + +def main() -> int: + srt = _srt_dir() + if srt is None: + print(f"{_TAG} sglang not importable — skipping") + return 0 + + paths = {_A_REL: srt / _A_REL, _B_REL: srt / _B_REL} + for rel, p in paths.items(): + if not p.is_file(): + return _fail(f"{p} is missing — sglang layout changed") + + # Plan both files before writing either: edit A without edit B hands aiter a + # mis-typed bias pointer, which is worse than the defect. + planned = {} + for rel, p in paths.items(): + result = _plan_a(p) if rel == _A_REL else _plan_b(p) + if isinstance(result, int): + return result + planned[rel] = result + + changed = [rel for rel, text in planned.items() if text != paths[rel].read_text()] + if not changed: + print(f"{_TAG} already present — skipping") + return 0 + + for rel in changed: + paths[rel].write_text(planned[rel]) + print(f"{_TAG} patched {paths[rel]}") + print(f"{_TAG} GLM MoE gate bias now stays fp32 through the aiter router") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/deploy/docker/patches/sglang_dsa/patch_draft_cuda_graph_dp_vote_v0518.py b/deploy/docker/patches/sglang_dsa/patch_draft_cuda_graph_dp_vote_v0518.py new file mode 100644 index 000000000..59700df8c --- /dev/null +++ b/deploy/docker/patches/sglang_dsa/patch_draft_cuda_graph_dp_vote_v0518.py @@ -0,0 +1,151 @@ +#!/usr/bin/env python3 +"""Patch 04's one v0.5.18-drifted edit in dp_attn.py, carried by anchor. + +`draft_cuda_graph_dp_vote.diff` carries seven files. On the v0.5.18 base six +files apply clean, and `dp_attn.py` applies 6 of its 7 hunks -- only hunk 4, the +min()-reduce, rejects. It rejects for a boring reason: upstream turned the +per-field device reads into one D2H copy and renamed the tensor with it, so +every context line in that hunk is stale. + + v0.5.17 v0.5.18 + tp0_info[:, 6].min().item() tp0_info_cpu[:, 6].min() + +Nothing about the mechanism changed, so this is an anchor script rather than a +re-cut diff -- it edits by unique source text and therefore survives that kind +of rename, the same reasoning that makes patch 01 a script. + +WHY THIS ONE HUNK IS THE WHOLE PATCH + Losing it does NOT leave patch 04 partly applied in a way that degrades + gracefully. The other six edits declare the dataclass field, contribute the + rank-local answer to the all-gather, and consume the result -- but hunk 4 is + the only place the gathered column is min()-reduced back onto + `self.can_run_draft_cuda_graph`. Without it that attribute keeps the value + the constructor put there, which is THIS RANK's answer. The vote silently + never happens, every consumer downstream reads a rank-local bool, and the DP + group splits across graph-replay and eager exactly as it did unpatched. + + That is the failure mode the apply script's header warns about: "an inert 04 + looks exactly like a working one until load." It is also why this cannot be + left to the `dp_attn.py:can_run_draft_cuda_graph` marker -- that identifier + is present after six hunks, so the marker passes on an inert patch. The + assertions below are the real gate. + +WHAT (unchanged from the diff -- see its header for the full analysis) + `EagleDraftWorker.draft()` picks graph-replay vs eager PER RANK, and two of + the guard's four terms are rank-dependent by construction. The two paths do + not issue the same host-side collective sequence, so the DP group diverges + and deadlocks on the first routed request. Fix: carry the choice as one more + int64 slot in the MLP-sync all-gather the scheduler already performs, + min()-reduced, so any rank needing eager takes the whole group eager. + +SCOPE + Wired into the `full` arm only, which is the v0.5.18 gfx950 image + (`Dockerfile.sglang`). The `indexer` arm (gfx942 / v0.5.16) does not carry + 04 at all and substitutes it at runtime with + `--json-model-override-args '{"index_share_for_mtp_iteration":false}'`; its + unused sibling port is `patch_draft_cuda_graph_dp_vote_v0516.py`, which + carries all seven edits because on that base the whole file rejects. + + DELETE THIS FILE when the diff is re-cut against a newer base, or when 04 + lands upstream. It is a one-line bridge, not a second source of truth. + +All-or-nothing: the anchor must match exactly once and all six sibling edits +must already be in place, or nothing is written and this exits non-zero -- so a +drifted base fails the build instead of shipping a patch that is present, +inert, and indistinguishable from a working one. +""" + +from __future__ import annotations + +import os +import sys + +REL = "python/sglang/srt/managers/scheduler_components/dp_attn.py" + +# The edit. Anchor is rewritten to itself plus the new line, so this is a pure +# insertion, and it is placed beside the two gates it is reduced like. +ANCHOR = " self.can_run_prefill_cuda_graph = bool(tp0_info_cpu[:, 6].min())\n" +INSERT = ( + " self.can_run_draft_cuda_graph = bool(tp0_info_cpu[:, 7].min())\n" +) + +# Proof that hunk 4 specifically landed. Distinct from the bare identifier, +# which six hunks already satisfy. +MARKER = "self.can_run_draft_cuda_graph = bool(tp0_info_cpu[:, 7]" + +# The six edits the diff itself applies. If any is missing this script is being +# run somewhere it does not belong -- inserting the reduce alone would then read +# a column nothing writes. +SIBLINGS: list[tuple[str, str]] = [ + ("dataclass field", " can_run_draft_cuda_graph: bool\n"), + ("all-gather slot", " int(self.can_run_draft_cuda_graph),\n"), + ("inactive-rank slot", " 1, # can_run_draft_cuda_graph\n"), + ( + "publish onto batch", + " batch.can_run_dp_draft_cuda_graph = mlp_sync_info.can_run_draft_cuda_graph\n", + ), + ("rank-local answer", " can_run_draft_cuda_graph = not (\n"), + ( + "constructor arg", + " can_run_draft_cuda_graph=can_run_draft_cuda_graph,\n", + ), +] + +# Slot 7 needs no width edit -- the all-gather derives its width from the tensor +# itself (`info_width = local_info_tensor.numel()`). These two readers are the +# ones that could have hard-coded the old width; assert they survive. +SLICE_READERS = ("tp0_info_cpu[:, 4:6]", "tp0_info_cpu[:, 5]") + + +def main() -> int: + root = os.environ.get("SGLANG_DIR", "/sgl-workspace/sglang") + path = os.path.join(root, REL) + if not os.path.isfile(path): + print(f"[draft-dp-vote-v0518] MISSING {path}", file=sys.stderr) + return 1 + + with open(path, encoding="utf-8") as f: + src = f.read() + + if MARKER in src: + print("[draft-dp-vote-v0518] already present — skipping") + return 0 + + for name, text in SIBLINGS: + if text not in src: + print( + f"[draft-dp-vote-v0518] sibling edit missing ({name}) — " + "draft_cuda_graph_dp_vote.diff did not apply here, refusing to write", + file=sys.stderr, + ) + return 1 + + n = src.count(ANCHOR) + if n != 1: + print( + f"[draft-dp-vote-v0518] anchor matched {n} times (want 1) — base drifted, " + "refusing to write", + file=sys.stderr, + ) + return 1 + + out = src.replace(ANCHOR, ANCHOR + INSERT, 1) + + for slice_expr in SLICE_READERS: + if slice_expr not in out: + print( + f"[draft-dp-vote-v0518] expected {slice_expr} to survive — " + "check slot layout", + file=sys.stderr, + ) + return 1 + + with open(path, "w", encoding="utf-8") as f: + f.write(out) + + print(f"[draft-dp-vote-v0518] applied the min()-reduce to {REL}") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/deploy/docker/scripts/apply_sglang_dsa_patches.sh b/deploy/docker/scripts/apply_sglang_dsa_patches.sh index 75fae8334..720269004 100755 --- a/deploy/docker/scripts/apply_sglang_dsa_patches.sh +++ b/deploy/docker/scripts/apply_sglang_dsa_patches.sh @@ -7,24 +7,30 @@ # # TWO ARMS, because the set is not uniformly portable across our engine bases: # -# DSA_PATCH_SET=full (default; Dockerfile.sglang, mi35x / v0.5.17) -# patch 01 + dsa_dp_sync + dsa_page_table_rows + draft_cuda_graph_dp_vote. -# The three diffs are `--fuzz=0` against that one release, so this arm -# only works there. +# DSA_PATCH_SET=full (default; Dockerfile.sglang, mi35x / v0.5.18) +# patch 01 + dsa_dp_sync + dsa_page_table_rows + draft_cuda_graph_dp_vote, +# plus one anchor port. The three diffs were cut `--fuzz=0` against +# v0.5.17; on the v0.5.18 base 02a and 02b still apply byte-for-byte (at +# an offset, which is not fuzz), and 04 applies everywhere except one +# hunk of dp_attn.py -- see EXPECT_REJECT below. # DSA_PATCH_SET=indexer (Dockerfile.sglang.gfx942, mi30x / v0.5.16) # patch 01 only. 02b is substituted at RUNTIME by # `--json-model-override-args '{"index_share_for_mtp_iteration":false}'` # -- see patches/sglang_dsa/README.md; the gfx942 recipe MUST pass it. # -# THE TWO EMPTY TABLES BELOW ARE DELIBERATE. `EXPECT_REJECT` scopes tolerance -# for a diff expected to reject one named file; `PORT_SCRIPTS` lists the anchor -# script that carries that file instead. No arm needs either today, so any -# rejection fails the build. One port script is kept unused: -# patches/sglang_dsa/patch_draft_cuda_graph_dp_vote_v0516.py carries 04's -# dp_attn.py to v0.5.16 by anchor -- six of that diff's seven files apply there, -# and dp_attn.py rejects all 7 of its hunks because the two gates beside them -# were renamed. Verified in-image, so 04 need not be re-cut if it is ever -# wanted on that base. +# TWO PORT SCRIPTS EXIST, AND ONLY ONE IS WIRED IN. `EXPECT_REJECT` scopes +# tolerance for a diff expected to reject one named file; `PORT_SCRIPTS` lists +# the anchor script that carries that file instead. Both are populated for the +# `full` arm and empty for `indexer`, so on `indexer` any rejection still fails +# the build. +# +# ..._v0518.py WIRED into `full`. Carries the single hunk of 04's dp_attn.py +# that v0.5.18 rejects. See EXPECT_REJECT below. +# ..._v0516.py UNUSED. Carries all seven of 04's dp_attn.py edits to +# v0.5.16, where the whole file rejects because the two gates +# beside them were renamed. The `indexer` arm does not carry 04 +# at all, so nothing runs it; it is kept, verified in-image, so +# 04 need not be re-cut if it is ever wanted on that base. # # THE `indexer` ARM IS NARROW ON PURPOSE. 02a and 04 can look necessary there # after a GPU fault that is really the driver: an out-of-support host-driver / @@ -69,12 +75,27 @@ case "$DSA_PATCH_SET" in esac # Per-arm escape hatch for a diff that is expected to reject ONE named file -# because an anchor script carries it instead. Both are deliberately EMPTY: no -# arm needs this today. The mechanism stays because it is the only safe way to -# express that tolerance -- scoped to a named file, so any OTHER rejection still -# fails the build. An unset expectation waves nothing through; see the loop. +# because an anchor script carries it instead. Tolerance is scoped to a named +# file, so any OTHER rejection still fails the build, and an unset expectation +# waves nothing through; see the loop. +# +# The `full` arm needs it since the base moved v0.5.17 -> v0.5.18. 04's +# dp_attn.py applies 6 of its 7 hunks there; hunk 4 -- the min()-reduce -- fails +# because upstream collapsed the per-field device reads into one D2H copy and +# renamed the tensor with it (`tp0_info[:, N].min().item()` -> +# `tp0_info_cpu[:, N].min()`). Losing that one hunk does NOT degrade +# gracefully: it is the only place the gathered column is reduced back onto +# `self.can_run_draft_cuda_graph`, so without it every rank reads its own answer +# and the vote silently never happens. The port script re-asserts all six +# sibling edits before writing, which the `dp_attn.py:can_run_draft_cuda_graph` +# marker below cannot -- that identifier is present after six hunks, so the +# marker passes on an inert patch. declare -A EXPECT_REJECT=() PORT_SCRIPTS=() +if [ "$DSA_PATCH_SET" = "full" ]; then + EXPECT_REJECT[draft_cuda_graph_dp_vote.diff]="python/sglang/srt/managers/scheduler_components/dp_attn.py" + PORT_SCRIPTS=(patch_draft_cuda_graph_dp_vote_v0518.py) +fi # module basename : identifier that must be present in the compiled bytecode # diff --git a/deploy/docker/scripts/print_ionic_abi.py b/deploy/docker/scripts/print_ionic_abi.py new file mode 100644 index 000000000..57b24c441 --- /dev/null +++ b/deploy/docker/scripts/print_ionic_abi.py @@ -0,0 +1,92 @@ +#!/usr/bin/env python3 +"""Print the kernel-uverbs ABI range the installed ionic provider declares. + +Why this exists: a wrong libionic fails *silently*. libibverbs prints one +warning per device to stderr, ibv_get_device_list then returns 0 HCAs, and +mooncake quietly serves KV over TCP -- every pod stays Ready, every request +succeeds, and the only symptom is that PD is slow. So the build records the +declared range in the build log, and refuses to produce an image whose +provider cannot be inspected at all. + +The range lives in the provider's `struct verbs_device_ops`: + + static const struct verbs_device_ops ionic_dev_ops = { + .name = "ionic", <- const char * + .match_min_abi_ver = N, <- uint32 + .match_max_abi_ver = N, <- uint32 + ... + +so: find the address of the "ionic" string in .rodata, find the 8-byte +pointer to it in the file, and read the two uint32 that follow. + +Measured on 2026-08-28, all four versions currently in repo.radeon.com: + + 54.0-149.g3304be71 abi 4..4 (pool line 1.117.1-a-63) + 54.0-187-1 abi 1..1 (pool line 1.117.5-a-77) + 54.0-192-1 abi 1..1 (pool line 1.125.0-a-187) + 54.0-197-1 abi 1..1 (pool line 1.117.5-a-147) + +Note the ordering: newer is NOT higher-ABI. Pick by the host, never by +version number. See the Dockerfile block for how. +""" + +import glob +import re +import struct +import subprocess +import sys + +LIBDIR = "/usr/lib/x86_64-linux-gnu" + + +def readelf(*args): + return subprocess.run( + ["readelf", *args], capture_output=True, text=True, errors="replace" + ).stdout + + +def declared_abi(path): + # Address of the bare "ionic" string in .rodata. + off = None + for line in readelf("-p", ".rodata", "-W", path).splitlines(): + m = re.match(r"\s*\[\s*([0-9a-f]+)\]\s+ionic$", line) + if m: + off = int(m.group(1), 16) + if off is None: + return None + m = re.search(r"\.rodata\s+\S+\s+([0-9a-f]+)\s+([0-9a-f]+)", readelf("-S", "-W", path)) + if not m: + return None + va = int(m.group(1), 16) + off + + data = open(path, "rb").read() + needle = struct.pack("/abi_version") + return rc + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/deploy/docker/scripts/reanchor_sglang_disagg_glm53.sh b/deploy/docker/scripts/reanchor_sglang_disagg_glm53.sh new file mode 100755 index 000000000..c44a2131d --- /dev/null +++ b/deploy/docker/scripts/reanchor_sglang_disagg_glm53.sh @@ -0,0 +1,82 @@ +#!/usr/bin/env bash +# Re-anchor the two patches/sglang_disagg/ scripts onto a POST-v0.5.17 sglang +# tree. Both consumers need it and both need the SAME two edits: +# +# Dockerfile.sglang stock v0.5.18 (glm_moe_dsa family) +# Dockerfile.sglang.glm53 PR #36607 head (glm5_next family), stacked on #36507 +# +# The name still says glm53 because that is where it was first needed; it is not +# GLM-5.3-specific. Verified 2026-09-01 inside lmsysorg/sglang:v0.5.18-rocm720-mi35x +# that both drifted anchors below read EXACTLY as they do on the GLM-5.3 branch, +# that disaggregation/common/utils.py still has the OLD typing import (so the +# conn.py edit stays per-file rather than a global sed), and that conn.py still +# has no wait_event/synchronize() call — i.e. the fix is still needed there too. +# +# WHY: both scripts are literal-anchor rewriters that exit 1 when an anchor +# drifted, which is correct — an image whose PD path silently corrupts long +# prompts should not ship. Between the v0.5.17 base and these trees, exactly two +# anchor LINES moved, and neither fix moved upstream: +# +# 1. mooncake/conn.py `from typing import List, Optional, Tuple, Union` +# -> `from typing import List, Optional, Set, Tuple, Union` +# Only conn.py grew `Set`; disaggregation/common/utils.py did NOT, so this +# is a per-file edit, not a global sed. conn.py at the GLM-5.3 head still +# has no wait_event/synchronize() call, i.e. the fix is still needed. +# +# 2. openai/serving_responses.py `background=request.background,` +# -> `background=request.background and not request.stream,` +# The bootstrap trio is still absent at that head, i.e. still needed. +# +# The scripts themselves are copied, not edited in place, so the v0.5.17 Kimi-K3 +# image keeps building from the untouched originals. +# +# Fails loudly if an old literal is gone: that means upstream moved again (or +# took the fix), and the Dockerfile should be re-derived rather than guessing. +set -euo pipefail + +SRC="${1:?usage: reanchor_sglang_disagg_glm53.sh }" +DST="${2:?usage: reanchor_sglang_disagg_glm53.sh }" + +mkdir -p "$DST" +cp "$SRC"/*.py "$DST"/ + +python3 - "$DST" <<'PY' +import sys, pathlib + +dst = pathlib.Path(sys.argv[1]) + +EDITS = [ + ( + "patch_mooncake_early_send_wait_event.py", + # Scoped to the conn.py entry of _EDITS. utils.py keeps the old import. + ''' "disaggregation/mooncake/conn.py": [ + ( + "from typing import List, Optional, Tuple, Union", + "from typing import Any, List, Optional, Tuple, Union",''', + ''' "disaggregation/mooncake/conn.py": [ + ( + "from typing import List, Optional, Set, Tuple, Union", + "from typing import Any, List, Optional, Set, Tuple, Union",''', + 1, + ), + ( + "patch_responses_pd_bootstrap.py", + " background=request.background,\n", + " background=request.background and not request.stream,\n", + 2, + ), +] + +for name, old, new, want in EDITS: + p = dst / name + s = p.read_text() + got = s.count(old) + if got != want: + sys.exit( + f"[reanchor] {name}: found {got}x of the v0.5.17 anchor, want {want}x.\n" + f" The script or upstream moved again — re-derive the anchors\n" + f" against the GLM-5.3 source instead of shipping a guess." + ) + p.write_text(s.replace(old, new)) + print(f"[reanchor] {name}: {want} anchor(s) retargeted to the GLM-5.3 tree") +PY diff --git a/deploy/docker/scripts/verify_glm53_overlay.py b/deploy/docker/scripts/verify_glm53_overlay.py new file mode 100644 index 000000000..38b60131d --- /dev/null +++ b/deploy/docker/scripts/verify_glm53_overlay.py @@ -0,0 +1,135 @@ +#!/usr/bin/env python3 +"""Assert that the GLM-5.3-Flash (glm5_next) overlay actually landed in the image. + +Run at the END of Dockerfile.sglang.glm53, after every later layer has had its +chance to clobber /sgl-workspace/sglang. Cheap, and it turns "a later layer +reinstalled sglang over the overlay" into a build failure instead of a +CrashLoopBackOff twenty minutes into a weight load. + +DELIBERATELY NOT a ModelRegistry.get_supported_archs() check. registry.py +imports every model module with strict=False and logs+swallows the failures, so +on a GPU-less builder (aiter -> rocminfo) glm5_next would be reported "not +registered" for a reason that has nothing to do with this image -- a false +failure that would train people to ignore the check. Instead: verify the same +contract STATICALLY (module present, class defined, exported via EntryClass, +compiles), and only then attempt the real import, treating a GPU-absence error +as a pass. +""" + +import ast +import importlib +import pathlib +import sys +import traceback + +MODULE = "sglang.srt.models.glm5_next" +CLASS = "Glm5NextForConditionalGeneration" +SRC = pathlib.Path("/sgl-workspace/sglang/python/sglang/srt/models/glm5_next.py") + +# An import failure naming any of these is the builder lacking a GPU, not the +# overlay missing. Keep this list tight: anything not listed is a real failure, +# and a marker that is too broad turns this check into a rubber stamp -- which +# is worse than not having it, because the build then claims the overlay landed. +# +# Two that used to be here and are not: +# "HIP error" matches every runtime error HIP raises, a real broken kernel +# included. The device-absence ones are spelled out instead. +# "torch.cuda" appears in the message of any AttributeError raised while +# touching a torch.cuda symbol -- including one raised BY +# glm5_next.py against a torch this image actually ships, which +# is precisely the defect this script exists to catch. +GPU_ABSENCE = ( + "rocminfo", + "No HIP GPUs are available", + "no CUDA-capable device", + "hipErrorNoDevice", + "no ROCm-capable device", + "Found no NVIDIA driver", + "Torch not compiled with CUDA enabled", +) + + +def _is_gpu_absence(exc): + """True only for a device-absence failure raised *outside* the overlay. + + Two conditions, both required. The message test walks the `__cause__` / + `__context__` chain, because aiter's rocminfo probe is routinely re-raised + as something whose own `str()` says nothing about GPUs. The frame test is + the sharper half: whatever the message says, if the innermost frame is in + glm5_next.py then the overlay is what broke, and a builder without a GPU is + not the explanation. The usual true-negative -- glm5_next.py's own + `import aiter` -- puts glm5_next.py in the traceback but never at the + bottom of it. + """ + seen, cur, blobs = set(), exc, [] + while cur is not None and id(cur) not in seen: + seen.add(id(cur)) + blobs.append(f"{type(cur).__name__}: {cur}") + cur = cur.__cause__ or cur.__context__ + joined = "\n".join(blobs).lower() + if not any(m.lower() in joined for m in GPU_ABSENCE): + return False + + innermost = None + cur, seen = exc, set() + while cur is not None and id(cur) not in seen: + seen.add(id(cur)) + frames = traceback.extract_tb(cur.__traceback__) + if frames: + innermost = frames[-1].filename + cur = cur.__cause__ or cur.__context__ + if innermost and pathlib.Path(innermost).resolve() == SRC.resolve(): + return False + return True + + +def fail(msg): + sys.exit(f"[verify-glm53-overlay] FAIL: {msg}") + + +def main(): + if not SRC.is_file(): + fail( + f"{SRC} is absent -- the overlay did not land, or a later layer " + f"reinstalled sglang over it" + ) + + text = SRC.read_text() + try: + tree = ast.parse(text, filename=str(SRC)) + except SyntaxError as e: + fail(f"{SRC} does not parse: {e}") + + classes = {n.name for n in tree.body if isinstance(n, ast.ClassDef)} + if CLASS not in classes: + fail( + f"{SRC} defines no `class {CLASS}` (found {len(classes)} classes) " + f"-- wrong commit checked out?" + ) + + if "EntryClass" not in text: + fail(f"{SRC} has no EntryClass export; sglang's loader would never reach {CLASS}") + + compile(text, str(SRC), "exec") # bytecode-level, not just AST + print(f"[verify-glm53-overlay] static: {CLASS} defined and exported in {SRC}") + + # Best effort. A GPU-absence error on the builder is expected and passes. + try: + mod = importlib.import_module(MODULE) + except Exception as e: # noqa: BLE001 - see GPU_ABSENCE + blob = f"{type(e).__name__}: {e}" + if _is_gpu_absence(e): + print( + f"[verify-glm53-overlay] import skipped, builder has no GPU " + f"({blob.splitlines()[0][:120]})" + ) + return + fail(f"importing {MODULE} raised {blob}") + + if not hasattr(mod, CLASS): + fail(f"{MODULE} imported but does not expose {CLASS}") + print(f"[verify-glm53-overlay] import OK: {MODULE}.{CLASS}") + + +if __name__ == "__main__": + main() diff --git a/examples/sglang_1p1d_glm5.2/engine/leg.sh b/examples/sglang_1p1d_glm5.2/engine/leg.sh index b7fc1eb4f..6dfe23c4d 100755 --- a/examples/sglang_1p1d_glm5.2/engine/leg.sh +++ b/examples/sglang_1p1d_glm5.2/engine/leg.sh @@ -57,8 +57,31 @@ CHUNK="${CHUNK:-65536}" # pool, or a 5-20x slower TCP fallback. See cluster/README.md section 3. export MOONCAKE_DISABLE_HIP_DMABUF="${MOONCAKE_DISABLE_HIP_DMABUF:-1}" export MC_GID_INDEX="${MC_GID_INDEX:?MC_GID_INDEX is required — read it off the preflight report}" -export MC_DISABLE_HIP_TRANSPORT=1 +# MC_DISABLE_HIP_TRANSPORT AND MC_ENABLE_HIP_TRANSPORT ARE DEAD NAMES. The +# shipped mooncake build does not contain either string: +# exact matches in mooncake/engine.cpython-310-x86_64-linux-gnu.so +# MC_DISABLE_HIP 1 <- the real knob +# MC_DISABLE_HIP_TRANSPORT 0 +# MC_ENABLE_HIP_TRANSPORT 0 +# MC_USE_HIP_IPC 1 <- a second gate, uninvestigated +# Measured, not inferred: a leg launched with MC_DISABLE_HIP_TRANSPORT=1 present +# in /proc//environ still logged "HIP transport installed for intra-node GPU +# P2P" 4x per leg. So these two lines have never had an effect, in either the +# two-node or the single-node shape. +# +# Kept, made overridable, and NOT called a fix: it changes nothing either way, +# and it is here only so the pair reads together with the live knob below. +# INFERRED: whoever wrote the original hardcode probably meant to disable hip and +# used a name that does not exist -- which would explain why no reason for the +# line could be established. There was no effect to reason about. +export MC_DISABLE_HIP_TRANSPORT="${MC_DISABLE_HIP_TRANSPORT:-1}" unset MC_ENABLE_HIP_TRANSPORT +# The one that works. Forwarded only when set, so unset == today's behaviour and +# the two-node path does not move. Set it to 1 to actually disable the hip +# transport -- which is the only way to run a meaningful hip-on/hip-off A/B. +# Verify it took: "HIP transport installed for intra-node GPU P2P" must go from +# 4 per leg to 0. If it stays at 4, the gate is elsewhere (MC_USE_HIP_IPC?) -- +# stop rather than hunting for a third name. # MC_MS_FILTERS pins mooncake to named device(s). Required in the dma-buf mode so a non-ODP # rail is never picked (it would pin and double the KV pool); harmless to leave unset when a # peer-mem module is loaded and every rail can carry KV. @@ -143,7 +166,12 @@ CAR_ARGS=() # --enable-cache-report populates usage.prompt_tokens_details.cached_tokens. Without it every # client-side cache-hit metric reads 0 and a prefix-reuse target cannot be checked at all. -EXTRA_ARGS=(--enable-cache-report) +# EXTRA_ENGINE_ARGS is the cluster wrapper's escape hatch for flags this script +# does not model -- e.g. --disable-shared-experts-fusion, which the GLM-5.3 +# wrappers pass. It was previously defined by those wrappers and read here by +# nobody, so the flag never reached the engine. Unquoted on purpose: several +# flags in one string must word-split. Empty by default, so unchanged when unset. +EXTRA_ARGS=(--enable-cache-report ${EXTRA_ENGINE_ARGS:-}) log "$ROLE on $MY_IP:$PORT — tp=$TP dpa=$DPA mtp=$MTP kvaware=$KVAWARE kvd=$KVD gmu=$GMU chunk=$CHUNK ctx=$CTX nic=$NIC ib=$RDMA_IB_DEVICES" @@ -152,7 +180,8 @@ log "$ROLE on $MY_IP:$PORT — tp=$TP dpa=$DPA mtp=$MTP kvaware=$KVAWARE kvd=$KV docker exec -d "$CTR" env \ HIP_VISIBLE_DEVICES="$GPUS" \ MOONCAKE_DISABLE_HIP_DMABUF="$MOONCAKE_DISABLE_HIP_DMABUF" \ - MC_GID_INDEX="$MC_GID_INDEX" MC_DISABLE_HIP_TRANSPORT=1 \ + MC_GID_INDEX="$MC_GID_INDEX" MC_DISABLE_HIP_TRANSPORT="$MC_DISABLE_HIP_TRANSPORT" \ + ${MC_DISABLE_HIP:+MC_DISABLE_HIP="$MC_DISABLE_HIP"} \ ${MC_MS_FILTERS:+MC_MS_FILTERS="$MC_MS_FILTERS"} \ ${MC_MS_FILTERS:+MC_MS_AUTO_DISC="${MC_MS_AUTO_DISC:-0}"} \ ${RDMAV_FORK_SAFE:+RDMAV_FORK_SAFE="$RDMAV_FORK_SAFE"} \ diff --git a/examples/sglang_1p1d_glm5.2/engine/up.sh b/examples/sglang_1p1d_glm5.2/engine/up.sh index a4c4edae0..d1338e585 100755 --- a/examples/sglang_1p1d_glm5.2/engine/up.sh +++ b/examples/sglang_1p1d_glm5.2/engine/up.sh @@ -35,7 +35,9 @@ ${MC_MS_AUTO_DISC:+MC_MS_AUTO_DISC=$MC_MS_AUTO_DISC} \ ${RDMAV_FORK_SAFE:+RDMAV_FORK_SAFE=$RDMAV_FORK_SAFE} \ ${HOST_RDMA_LIB:+HOST_RDMA_LIB=$HOST_RDMA_LIB} \ ${ENTRYPOINT_KEEP:+ENTRYPOINT_KEEP=$ENTRYPOINT_KEEP} \ -${GMU_PREFILL:+GMU_PREFILL=$GMU_PREFILL} ${GMU_DECODE:+GMU_DECODE=$GMU_DECODE}" +${GMU_PREFILL:+GMU_PREFILL=$GMU_PREFILL} ${GMU_DECODE:+GMU_DECODE=$GMU_DECODE} \ +${EXTRA_ENGINE_ARGS:+EXTRA_ENGINE_ARGS=\"$EXTRA_ENGINE_ARGS\"} \ +${MC_DISABLE_HIP:+MC_DISABLE_HIP=$MC_DISABLE_HIP}" log "=== 1/4 containers ===" for h in "$PREFILL_NODE" "$DECODE_NODE"; do @@ -64,11 +66,35 @@ if [ "${DECODE_KVD:-0}" = "1" ]; then start_kvd "$DECODE_NODE"; fi log "=== 3/4 legs ===" # Launch both legs before waiting on either: they load ~400 GB of weights concurrently, and # serialising the waits doubles the bring-up for no reason. +# GPUS is forwarded PER LEG with :+ (inject only when set), NOT :- with a +# default. A default would push GPUS into the two-node path, where leg.sh's own +# `seq 0..TP-1` is correct and a hardcoded list would be wrong at any TP != that +# list's length. Unset on both sides here == today's behaviour exactly. +# +# Until this existed, a single-node pair silently put BOTH legs on the same +# cards: the wrapper's PREFILL_GPUS/DECODE_GPUS were read by nobody, leg.sh:26 +# fell back to `seq 0..TP-1`, and the second leg's weights landed on top of the +# first's. Measured: GPUs 0-3 at 263.8 GB each with 4-7 at 0.3 GB, and the +# prefill leg then died on "Loaded weights leave no GPU memory for the KV cache", +# which reads as a GMU tuning problem and is not one. +# +# The KV-event ports are forwarded PER LEG. They default to leg.sh's own values, +# so the two-node shape is unchanged -- there the legs are on different hosts and +# cannot collide. On a SINGLE-NODE pair they share one network namespace, and +# without distinct values the second leg dies at bind with "port_base at N is not +# available". Nothing a wrapper exports reaches leg.sh except through here: +# `on()` runs a fresh remote shell. on "$PREFILL_NODE" "$COMMON_ENV ROLE=prefill MY_IP=$PREFILL_IP PORT=$PREFILL_PORT \ DPA=${PREFILL_DPA:-0} MTP=${PREFILL_MTP:-0} KVD=${PREFILL_KVD:-1} \ + KV_PUB_PORT=${PREFILL_KV_PUB_PORT:-5557} KV_SNAP_PORT=${PREFILL_KV_SNAP_PORT:-8801} \ + MC_DISABLE_HIP_TRANSPORT=${MC_DISABLE_HIP_TRANSPORT:-1} \ + ${PREFILL_GPUS:+GPUS=$PREFILL_GPUS} \ bash $KIT_DIR/engine/leg.sh" on "$DECODE_NODE" "$COMMON_ENV ROLE=decode MY_IP=$DECODE_IP PORT=$DECODE_PORT \ DPA=${DECODE_DPA:-1} MTP=${DECODE_MTP:-1} KVD=${DECODE_KVD:-0} \ + KV_PUB_PORT=${DECODE_KV_PUB_PORT:-5557} KV_SNAP_PORT=${DECODE_KV_SNAP_PORT:-8801} \ + MC_DISABLE_HIP_TRANSPORT=${MC_DISABLE_HIP_TRANSPORT:-1} \ + ${DECODE_GPUS:+GPUS=$DECODE_GPUS} \ bash $KIT_DIR/engine/leg.sh" # Poll /health from INSIDE each node's container. Never curl a PD leg's port from another diff --git a/examples/sglang_1p1d_glm5.3/README.md b/examples/sglang_1p1d_glm5.3/README.md new file mode 100644 index 000000000..25b540eeb --- /dev/null +++ b/examples/sglang_1p1d_glm5.3/README.md @@ -0,0 +1,498 @@ +# GLM-5.3 (big) — SGLang 1P1D + +Prefill/decode-disaggregated deployment for **GLM-5.3** and **GLM-5.3-MXFP4**, +in two shapes: the usual **two-node** pair, and a **single-node** pair that +splits one 8-GPU box into TP4 prefill + TP4 decode. + +## This kit does not fork the GLM-5.2 kit, and that is deliberate + +GLM-5.3 (big) is `glm_moe_dsa` / `GlmMoeDsaForCausalLM`. Its `config.json` is +identical to GLM-5.2's field for field except `transformers_version` — same +hidden size, same layer count, same expert count, same attention. So the engine +recipe is not *similar* to the GLM-5.2 one, it **is** the GLM-5.2 one. + +[`examples/sglang_1p1d_glm5.2/`](../sglang_1p1d_glm5.2/) already carries that +recipe in `engine/leg.sh`, validated end to end on two clusters and both RDMA +fabric types. Copying those ~600 lines here to change a model path would create +a second source of truth that drifts the first time either is fixed. So this kit +ships **wrappers only** and points `KIT_DIR` at the GLM-5.2 kit, exactly as that +kit's own `cluster/*.sh` files do. + +If you find yourself editing an engine script to serve GLM-5.3, something is +wrong — say so, it is a bug in this arrangement. + +**GLM-5.3-Flash is NOT covered by these wrappers**, but the question of whether +PD is *possible* for it now has an answer — see below. For Flash today, use +[`sglang_mix_glm5.3`](../sglang_mix_glm5.3/) (aggregated). + +### Flash PD: feasible, OUT OF SCOPE, never run + +**Deliberately not pursued.** What follows establishes that the shape is +*possible* — it is a source read, not a validation, and no Flash PD deployment +has ever been brought up. It is recorded so that whoever picks this up starts +from the code reads rather than repeating them, and so that "we did not do this" +is not mistaken for "this does not work". + +The two checks at the end of this section are what that person should run first. + +The obvious reason to expect Flash PD to be broken is that `glm5_next` keeps +**two** pools — the paged KV pool *and* a KDA recurrent-state pool (logged as +`mamba usage`) — while PD hands off after prefill. If the KDA state were not +transferred, the decode leg would start its linear-attention layers from a zero +state, and that fails as **subtly wrong output**, not as a crash. + +**The premise does not hold.** sglang's PD path is not KV-only: it carries a +generic *state component* mechanism, `StateType.MAMBA` is one of its members, and +`glm5_next` resolves onto exactly the pool that mechanism reads. Verified link by +link against the pinned ref `c821c425`: + +| link | where | +|---|---| +| `glm5_next` recognised as recurrent-state | `hybrid_arch.py:114` | +| folded into the generic predicate | `hybrid_arch.py:126` `mambaish_config()` | +| declared to keep SSM state | `kv_cache_builder.py:100` `uses_ssm_state()` | +| KDA state modelled as a mamba2 cache | `configs/glm5_next.py:264` | +| engine builds a hybrid req pool | `kv_cache_configurator.py:896` | +| pool exposes state buffers | `memory_pool.py:1436` `get_state_buf_infos` | +| **PD registers them as transferable** | `disaggregation/utils.py` — duck-typed on `get_state_buf_infos()` → `append_state_component(..., StateType.MAMBA, ...)` | +| prefill builds the payload | `prefill.py:1183` `_mamba_payload()` | +| decode has the matching pool | `decode.py:223` `HybridMambaDecodeReqToTokenPool` | +| transport is generic over components | `mooncake/conn.py:1241` iterates `state_types` | + +**No guard refuses it** — grepping `NotImplementedError|not supported` across +`srt/disaggregation/*.py` for mamba/hybrid/linear/kda/glm5 returns nothing. The +guards that exist are narrow and unrelated. + +**So the worry relocates rather than disappearing: from "not transferred" to +"transferred, but unverified".** Every link above is a code read; nobody has run +it, and whether the KDA conv/ssm buffers survive the mooncake round-trip +bit-exact on gfx950 is untested. That is still the failure mode that produces +subtly wrong output with nothing logged. + +**Two checks, in order, before believing any Flash PD result:** + +1. *Ten minutes, at bring-up.* Confirm at runtime that `StateType.MAMBA` actually + lands in `kv_args.state_types` for `glm5_next` — one log line at + `append_state_component`. That converts the table above from "the code says it + should" into "it did". +2. *The decisive one.* Send the same prompt to a Flash **MIX** deployment and a + Flash **PD** deployment, greedy, and **diff the token sequences**. Intact state + → identical; zeroed decode-leg state → divergence, with nothing logged. The + reference side already exists: `flash-mxfp4` MIX is validated end to end. + +One INFERRED caveat: `mori/conn.py:1041` refuses loudly when `state_types` is +empty; the mooncake path iterates `state_types` without an equivalent check *at +that site*, so an empty list there could be a silent no-op. Whether mooncake +checks elsewhere is UNKNOWN. + +Upstream has three OPEN items in this exact area — **#36651** (adding PD state +transfer for another Flash-class model, so the wiring is not automatic +everywhere), **#37276** (PD + mamba + speculative decoding, which is why running +Flash PD with MTP off sidesteps a known bug), and **#33457** (hybrid-linear KV +transfer under prefill pipeline parallelism). The machinery is live and being +repaired, not absent. + +## Contents + +| path | what | +|---|---| +| [`cluster.2node.sh`](cluster.2node.sh) | two-node pair — the validated shape. Fill in and run | +| [`cluster.singlenode.sh`](cluster.singlenode.sh) | one 8-GPU node split TP4 + TP4 | +| everything else | comes from [`../sglang_1p1d_glm5.2/`](../sglang_1p1d_glm5.2/) — `common.sh`, `engine/*.sh`, `preflight_rdma.sh` | + +## Validation status + +Stated plainly, because the honest answer is short. + +| shape | status | +|---|---| +| the **deployment shape** (1P1D + mooncake + DPA + MTP + kvd + kv-aware) | validated for **GLM-5.2** on two clusters, both fabric types | +| **GLM-5.3 weights** through that shape, two-node | **not yet run.** Same architecture, so expected to work; expectation is not evidence | +| **single-node** TP4+TP4 | **not yet run, and it carries a real unknown** — see below | + +## The single-node path: HIP IPC over XGMI, not loopback RDMA + +**An earlier version of this section said the same-host KV handoff is a loopback +RDMA transfer and that the risk is silent slowness. Both were wrong**, and the +correction changes what you check. Established by reading the mooncake tree and +build cache inside the shipped image: + +The pinned mooncake commit is `01d1eb2a` (2026-07-01), *"[TE] Support rdma+hip +multi-protocol segments for single-node disaggregation (#2682)"* — literally the +single-node disaggregation commit, whose own message reports validation of +single-node 1P1D on MI355X over the rdma+hip path. + +The image builds with `USE_HIP=ON`, `ENABLE_MULTI_PROTOCOL=ON`. On init, +`auto_discover` installs `rdma` (HCAs present, `MC_FORCE_TCP` unset) and then +**composes** `hip` on top — the local segment advertises `"rdma,hip"`. +Registration fans out to every installed transport, so device KV gets both a HIP +IPC buffer and an RDMA buffer, while host aux buffers land on rdma only. +`MultiTransport::selectTransport` then routes **per request** by fixed priority +`hip 4 > cxl 3 > rdma 2 > tcp 1`, so for KV **hip wins**: `hipIpcGetMemHandle` on +the exporter, `hipIpcOpenMemHandle` on the importer, `hipMemcpyAsync` over +enabled peer access. **GPU-to-GPU across XGMI, no NIC in the path.** + +### Everything in this section is single-node-specific, and that is by design + +**Two-node PD never uses hip, regardless of any setting.** `selectTransport` +calls `isHipReachableTarget()` and skips hip buffers whenever the target is on +another host. The in-source rationale, at the build commit: + +> *"This makes the intra-node fast path (hip) and the cross-node path (rdma) work +> automatically from a single multi-protocol segment, without requiring the +> operator to set `MC_DISABLE_HIP`."* + +So the multi-protocol segment is not a configuration problem to be solved — it +resolves itself by target. **Every hip question in this kit is a single-node +question**, and none of it applies to `cluster.2node.sh`. + +### Installation is unconditional; *selection* is what the knob controls + +`transfer_engine_impl.cpp:402-414` installs hip under a bare `#ifdef USE_HIP` +with no runtime condition. The gate is one stage later, at +`multi_transport.cpp:489`: + +```cpp +if (p == "hip") return std::getenv("MC_DISABLE_HIP") ? 0 : 4; +if (p == "rdma") return 2; +``` + +`MC_DISABLE_HIP` demotes hip from priority 4 to 0, so rdma wins for the device KV +pool — which is registered under both. **hip stays installed and stops being +used.** + +**This is why the obvious check is useless.** `HIP transport installed for +intra-node GPU P2P` is an **install-time** log, and the variable gates +**selection**. It reads 4/4 whether hip is carrying KV or demoted to zero, in +every state, forever. Verifying a hip-off arm requires `MC_DISABLE_HIP=1` present +in `/proc//environ` on both legs **plus** the source read above — there is +no log line that will confirm it, and treating the non-flip as evidence of +anything is how a correct hip-off deployment got discarded unmeasured. + +Two related names are **absent from the binary entirely** and do nothing: +`MC_DISABLE_HIP_TRANSPORT` (which `leg.sh:60` sets) and +`MC_ENABLE_HIP_TRANSPORT`. So: **two dead names, one live name whose effect is +invisible to the natural check.** + +One trap that makes the config lie: sglang passes `protocol="rdma"` into +`engine.initialize()` (`MOONCAKE_PROTOCOL` defaults to `"rdma"`). On this build +that argument **does not choose the transport** — outside the EFA/CXI paths it +only feeds `initMemoryAllocator()`. Setting `MOONCAKE_PROTOCOL` will not disable +hip, and seeing `rdma` in the config does not mean KV moves over RDMA. + +### The real risk: the two legs cannot see each other's GPUs + +`cluster.singlenode.sh` sets `PREFILL_GPUS=0,1,2,3` and `DECODE_GPUS=4,5,6,7`, +applied as `HIP_VISIBLE_DEVICES` (`../sglang_1p1d_glm5.2/engine/leg.sh:159`). The +two legs therefore have **disjoint visible device sets**, each seeing 4 devices +renumbered 0-3, and `setupP2PAccess()` only iterates visible devices — so peer +access is enabled *within* each leg and never between them. + +> **This was not true as originally shipped, and the failure is worth knowing.** +> `up.sh` forwards a fixed list of per-leg variables through `on()` — which runs a +> fresh remote shell — and `GPUS` was not among them. Both legs therefore fell +> through to `leg.sh:26`'s default, `seq 0..TP-1`, and **landed on the same four +> cards**. Fixed by forwarding `${PREFILL_GPUS:+GPUS=$PREFILL_GPUS}` per leg; +> conditional expansion, so an unset variable injects nothing and the two-node +> path is unchanged. +> +> **The way it failed is the instructive part.** The prefill leg died with +> `Loaded weights leave no GPU memory for the KV cache under +> --mem-fraction-static=0.7. Raise --mem-fraction-static above 0.773` — a number +> that is arithmetically correct and diagnostically wrong. Taking the engine's +> advice would have let two legs coexist on four cards and produced a deployment +> that *ran*, with every subsequent number meaningless and nothing saying so. +> **Before trusting any memory error, check that both halves of the box are +> loaded** — `rocm-smi --showmeminfo vram` should show weights on GPUs 0-3 *and* +> 4-7 (~408 GB / TP4 ≈ 102 GB per card for GLM-5.3-MXFP4). That distinguishes a +> tuning problem from a topology problem. +> +> **Do not use `base_gpu_id` for this.** It is tempting and it does not work: +> `HIP_VISIBLE_DEVICES=4,5,6,7` renumbers the decode leg's devices to 0-3, so +> `base_gpu_id` is an index into the *visible* set, not the physical one. It reads +> `0` on both legs when the split is broken **and** `0` on both legs when it is +> correct — it does not discriminate at all. The VRAM read is more expensive and +> it is the only unambiguous check here. + +**ANSWERED — it works.** Measured on gfx950 / ROCm 7.2 with the shipped engine +image, two processes in one container, `--ipc=host`: + +``` +exporter HIP_VISIBLE_DEVICES=0,1 writes pattern 7,3,9,1,4,1,5,9 to cuda:0 +importer HIP_VISIBLE_DEVICES=2,3 imports the handle, reads back + -> READ BACK: [7, 3, 9, 1, 4, 1, 5, 9] MATCH +``` + +Repeated **across two separate containers** (importer started with `--ipc=host`), +which is the shape PD actually runs — same disjoint split, same pattern: + +``` +CROSS-CONTAINER IMPORT OK, bytes= 1048576 +READ BACK: [7, 3, 9, 1, 4, 1, 5, 9] MATCH +``` + +The importer **cannot see the exporter's physical GPU** and still mapped its +memory and read the correct bytes. The single-container run alone would have +left the container boundary as an untested variable; it is closed. A bare "import succeeded" would not have +proved this — the handle records device index 0 and the importer's own ordinal 0 +is a *different* physical GPU, so the import could plausibly have mapped local +memory instead. The data pattern is what rules that out; **check the bytes, not +the return code**, if you repeat this. + +Caveat on scope: measured with a 4-GPU visible set split 0,1 / 2,3 (physical +4,5 / 6,7 of that host) rather than the 0-3 / 4-7 split this kit uses. Same node, +same XGMI fabric. Strong evidence, not proof, for the exact split. + +Note `torch.cuda.cudart()` does **not** expose `cudaIpcGetMemHandle` in this +build — use PyTorch's storage IPC path (`untyped_storage()._share_cuda_()` / +`torch.UntypedStorage._new_shared_cuda(*info)`), which is what actually carries +HIP IPC handles here. + +Two source reads that failed to answer this before the measurement, recorded so +nobody repeats them: + +Two attempts to close it from source, both negative, recorded so nobody repeats +them: + +- **The ROCm 7.2 header** (`hip_runtime_api.h:2535-2545`) says + `hipIpcOpenMemHandle` *"can attempt to enable peer access between the devices as + if the user called hipDeviceEnablePeerAccess"*, and points at + `hipDeviceCanAccessPeer` to test it. Suggestive, not decisive: + `hipDeviceCanAccessPeer` takes **visible** ordinals, and under disjoint + `HIP_VISIBLE_DEVICES` the importer cannot name the exporter's device at all. The + doc does not say what happens then. +- **Mooncake's own HIP tests do not cover it.** All three harnesses + (`tests/hip_transport_test.cpp`, `mooncake-wheel/tests/test_transfer_on_hip.py`, + `tent/tests/hip_bandwidth_bench.cpp`) are single-process and single-device, and + `grep -rn HIP_VISIBLE_DEVICES` over the whole repo returns nothing. So the + pinned commit's *"prefill GPU0 / decode GPU1"* validation is not reproducible + from the tree, and its test suite does not exercise two processes with disjoint + visible devices — which is exactly what this kit configures. + +That is the argument for running the probe below rather than reasoning further. + +What *is* established is the shape of each outcome: + +- **If it works:** KV moves over XGMI, and the only positive evidence is the + install line plus the absence of hip errors. +- **If it fails, it fails LOUDLY at transfer time, not silently.** Registration + still succeeds (`hipIpcGetMemHandle` is local), the segment still advertises + `"rdma,hip"`, `selectTransport` still picks hip, and then + `hipIpcOpenMemHandle failed` is logged and the transfer returns + `"device memory not registered"` — surfacing as *"Failed to get kvcache from + prefill instance"*, exactly the pre-fix symptom the pinned commit quotes. + +So **the single-node failure mode is a broken PD, not a slow one** — provided hip +is installed. The silent-slow path exists only if hip is *absent*. + +> **And on this build hip cannot be turned off by the variable anyone would +> reach for.** `leg.sh:60` exports `MC_DISABLE_HIP_TRANSPORT=1` and unsets +> `MC_ENABLE_HIP_TRANSPORT`. **Neither name exists in the shipped mooncake +> binary.** Exact-match against `mooncake/engine.*.so`: +> +> | env name | matches | +> |---|---:| +> | `MC_DISABLE_HIP` | **1** | +> | `MC_USE_HIP_IPC` | 1 | +> | `MC_FORCE_TCP` | 1 | +> | `MC_DISABLE_HIP_TRANSPORT` | **0** | +> | `MC_ENABLE_HIP_TRANSPORT` | **0** | +> +> Confirmed behaviourally as well as by inspection: a run launched with +> `MC_DISABLE_HIP_TRANSPORT=1` — verified present in the process environment via +> `/proc` — still logged `HIP transport installed for intra-node GPU P2P` **4× +> per leg**, identical to a run without it. +> +> Two consequences. **`leg.sh:60` has never had an effect**, in either the +> two-node or single-node path, so it is not evidence that anyone deliberately +> disabled hip — which is likely why no reason for it could be established +> (INFERRED: the author may have intended to disable hip and used a name that +> does not exist). And **any A/B that varies hip must set `MC_DISABLE_HIP`**; +> using the `_TRANSPORT` spelling produces a guaranteed-zero differential that +> reads as a null result rather than as a broken experiment. +> +> Before trusting any such A/B, confirm the discriminator actually flipped: +> `HIP transport installed for intra-node GPU P2P` must go **4/4 → 0/0**. + +If it fails, the fix is a topology change — give both legs all 8 GPUs and split +with `--base-gpu-id` so each process can see its peer's cards — not a mooncake +debug session. + +### Settle it in seconds, before loading any weights + +Two processes with the kit's own disjoint split, exchanging one IPC handle. No +model, no server: + +```bash +# exporter — the prefill leg's GPUs +docker exec -e HIP_VISIBLE_DEVICES=0,1,2,3 python - <<'EOF' +import torch +t = torch.zeros(1<<20, dtype=torch.uint8, device='cuda:0') +h = torch.cuda.cudart().cudaIpcGetMemHandle(t.data_ptr()) +open('/dev/shm/ipc.h','wb').write(bytes(h)); print("exported, holding"); input() +EOF + +# importer — the decode leg's GPUs +docker exec -e HIP_VISIBLE_DEVICES=4,5,6,7 python - <<'EOF' +import torch +torch.zeros(1, device='cuda:0') # init the HIP context first +h = open('/dev/shm/ipc.h','rb').read() +print(torch.cuda.cudart().cudaIpcOpenMemHandle(h, 1)) # 1 = LazyEnablePeerAccess +EOF +``` + +`--ipc=host` is already passed to both containers (`../sglang_1p1d_glm5.2/common.sh:46`), +which HIP IPC across processes requires. + +### What to grep, and the two lines nobody was checking + +| outcome | line | +|---|---| +| HIP transport installed | `HIP transport installed for intra-node GPU P2P` | +| HIP install failed | `Failed to install HIP transport (intra-node GPU P2P unavailable)` | +| RDMA installed | `installTransport, type=rdma` | +| KV not IPC-exportable | `HipTransport: hipIpcGetMemHandle failed` | +| peer's KV not importable | `HipTransport: hipIpcOpenMemHandle failed` | +| two GPUs cannot reach each other | `HipTransport: P2P access not available between device i and device j` | +| TCP forced | `MC_FORCE_TCP is set, using TCP transport only` | +| **TCP fallback (no HCAs)** | **nothing — see below** | + +Two properties of this table matter more than the table: + +1. **The TCP fallback is silent.** TCP is installed with no success log where the + RDMA branch logs `installTransport, type=rdma`. Grep for the *positive* rdma + line and require it; there is no tcp line to find. +2. **No log line says which transport a given transfer used.** `selectTransport` + chooses silently per request, and the two routing `LOG(ERROR)` calls in + `multi_transport.cpp` are commented out in this tree. The install lines tell + you the *capability*, never the *choice*. `MC_LOG_LEVEL=TRACE` adds + per-buffer registration lines — still not per-transfer routing. + +**The existing `MC_FORCE_TCP` / `GID is NULL` checks do not cover this case.** +`MC_FORCE_TCP` is an env var we would have to set ourselves, so counting it only +confirms we did not force TCP by accident. (It does catch one real disaster: if +set, init returns early *before* auto-discover, hip is never installed, and every +KV byte goes over TCP loopback — that genuinely is the 5-20× case.) `GID is NULL` +is per-RDMA-device rail health and is a **cross-host** signal; with the hip path +live the single-node KV transfer never touches a GID, so a count of 0 tells you +nothing about it. + +**So add one line to the single-node smoke: require `HIP transport installed for +intra-node GPU P2P` in BOTH leg logs**, plus zero `hipIpcOpenMemHandle failed`. +Without the first, the segment is `"rdma"` only and KV silently takes loopback +RDMA with nothing raised anywhere. + +What we *have* verified on the reference node, first-hand: + +- 8 ionic RDMA devices on the host, all `PORT_ACTIVE`. +- `ib_peer_mem` **loaded**, so registration **mode A** (bare `ibv_reg_mr` + + peer-mem: nothing pinned, KV pool not duplicated, every rail usable) is the + mode to expect. That is the best of the three. +- Inside the engine image, `libionic 54.0-187-1` (ABI 4) and `ibv_devinfo` + reporting **8 HCAs**. + +**That last check has a trap worth knowing before you repeat it.** The container +must be started with `--device=/dev/infiniband`. Without it `ibv_devinfo` reports +zero HCAs from inside a container on a host that has eight — which is +indistinguishable from a libionic ABI mismatch, and is the same reading that +means "RDMA has silently degraded to TCP". `common.sh` passes it; an ad-hoc +`docker run` will not unless you remember. + +Run `preflight_rdma.sh mode` before the first bring-up either way. Believe its +verdict from inside the container over the host's view: only the vendor provider +libraries the image ships can open a card. + +## Quick start + +```bash +# 1. which registration mode does this fabric support? +IMAGE= bash ../sglang_1p1d_glm5.2/preflight_rdma.sh mode + +# 2. fill in ONE wrapper, then +bash cluster.singlenode.sh up # or cluster.2node.sh +bash cluster.singlenode.sh smoke +bash cluster.singlenode.sh down +``` + +`smoke` is the GLM-5.2 kit's, and it checks each feature with a signal that goes +red when the feature is silently absent — including **`MC_FORCE_TCP` and +`GID is NULL` counts of 0 in both leg logs**, which is the check that catches a +pair that paired successfully and is moving KV over TCP. + +## Notes carried over from the GLM-5.2 kit that still apply + +These cost real debugging time there and the architecture has not changed: + +1. **`--ep-size` and `--enable-dp-attention` are different axes.** Gate both on + one condition and turning DPA off silently collapses the MoE from ep8 to the + TP default, after which no latency delta is attributable to either. +2. **`--chunked-prefill-size` is a GLOBAL budget** that SGLang divides by + `dp_size` only when DP-attention is on. One value serves both modes; + hardcoding the per-rank number in a DPA-off branch cuts it 8x. +3. **Prefill activation OOM is fixed by LOWERING `--mem-fraction-static`**, the + opposite of the decode-side fix. Diagnose by phase: decode retract → raise; + prefill `HSA_STATUS_ERROR_OUT_OF_RESOURCES` at *low* token usage → lower. + Low token usage at the abort is the tell that it was never KV exhaustion. + + **3a. A third form, and the engine's own advice is a trap in it.** On the + single-node shape a leg can abort at startup with + + ``` + ValueError: Loaded weights leave no GPU memory for the KV cache under + --mem-fraction-static=0.7. Raise --mem-fraction-static above 0.773 + ``` + + That number is arithmetically correct and **diagnostically wrong**. It is + computed from the memory actually free at that moment, and the reason there + is none is usually that **the other leg is on the same cards** — measured + once as GPUs 0-3 at 263.8 GB each with 4-7 at 0.3 GB. Raising the fraction + as instructed produces a *working* deployment on the wrong topology: two + legs sharing half the node, every subsequent number meaningless, nothing + logged. + + **Check the cards before you touch the knob.** `rocm-smi --showmeminfo vram` + must show load on *both* halves. Do not use `--showmemuse`'s `VRAM%` (it + does not fall when memory is released) and do not use each leg's + `base_gpu_id` (it is an index into the leg's **visible** set, so + `HIP_VISIBLE_DEVICES=4,5,6,7` renumbers the decode leg to 0-3 and it reads + `base_gpu_id=0` on both legs whether the split is broken or correct). + + Distinguishing the three: **this** form aborts during startup profiling with + weights already loaded and no request served; the classic prefill form + aborts under load at low token usage; the decode form retracts under load at + high token usage. +4. **`SGLANG_OPT_USE_TOPK_V2=0` is mandatory on gfx950.** Without it the model + serves, returns 200s, and returns garbage. +5. **MTP and decode-side radix cache are mutually exclusive upstream**, so + `decode_prefix_len` is always 0 and every turn re-transfers the whole prompt + KV. A prefill-side cache hit saves compute, not bytes — which is why fabric + bandwidth matters on long-prompt agentic workloads even at a high hit rate. +6. **An MTP acceptance length of a steady 4.00 is bad news**, not a good result: + the draft is predicting a repetition loop perfectly. 2-3 is healthy. + +## Two GLM-5.3-specific things to decide before you run + +**MTP.** The GLM-5.2 kit runs EAGLE MTP on the decode leg and that is its +validated configuration. For GLM-5.3 the evidence conflicts: upstream's GLM-5.3 +cookbook says MTP/EAGLE is **disabled on AMD** because the gfx950 draft kernel is +unvalidated, while the OneNexus GLM-5.3-MXFP4 model card runs EAGLE at +`--speculative-num-steps 3` and lists it as validated. Both wrappers here +default `MTP=0`. That is a choice to avoid an unvalidated variable on the first +run, not a finding — resolve it deliberately rather than inheriting it. + +**Shared-experts fusion.** Not a concern for the big MXFP4 checkpoint: its +shared experts are themselves MXFP4 (76 `.weight` / 75 `.weight_scale`, the odd +one being the BF16 MTP layer 78, which is not loaded while MTP is off), so the +precondition for the mismatch does not hold. It *is* a concern for Flash-MXFP4 — +see [`sglang_mix_glm5.3`](../sglang_mix_glm5.3/) and upstream issue #37268. +`glm4_moe.py`'s fusion gate only special-cases `w4afp8` and would fuse under +`quark`, so the wrappers pass `--disable-shared-experts-fusion` as insurance; +upstream #25261 shows this class failing *silently with wrong output* rather +than crashing when the shapes happen to line up. + +## Source + +[`examples/sglang_1p1d_glm5.3/`](.) in [AMD-AGI/Infera](https://github.com/AMD-AGI/Infera) +· [the GLM-5.2 kit this drives](../sglang_1p1d_glm5.2/) +· [aggregated MIX kit for all four GLM-5.3 checkpoints](../sglang_mix_glm5.3/) +· [PD disaggregation concepts](../../manual/features/pd_disaggregation.md) diff --git a/examples/sglang_1p1d_glm5.3/cluster.2node.sh b/examples/sglang_1p1d_glm5.3/cluster.2node.sh new file mode 100755 index 000000000..51da03f0a --- /dev/null +++ b/examples/sglang_1p1d_glm5.3/cluster.2node.sh @@ -0,0 +1,135 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# ============================================================================ +# EDIT THIS FILE. Nothing else needs changing. +# ============================================================================ +# +# GLM-5.3 (big) 1P1D across TWO nodes -- one prefill, one decode, KV over +# mooncake RDMA. This is the shape the GLM-5.2 kit validated end to end on two +# clusters and both fabric types; only the weights differ here, and GLM-5.3 big +# is the same architecture. +# +# NOT YET RUN with GLM-5.3 weights. Same architecture, so it is expected to +# work -- expectation is not evidence, so treat the first run as a bring-up. +# +# It drives the GLM-5.2 kit's engine scripts unchanged -- GLM-5.3 (big) is the +# same architecture, so there is nothing to fork. See the README. +# +# Usage: bash cluster.singlenode.sh up | smoke | bench [conc...] | down +set -euo pipefail +HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" + +# The kit whose engine/ and common.sh actually run. Do not point this at a copy. +export KIT_DIR="${KIT_DIR:-$(cd "$HERE/../sglang_1p1d_glm5.2" && pwd)}" + +# --------------------------------------------------------------------------- +# 1. Nodes -- one prefill, one decode +# --------------------------------------------------------------------------- +# Each node's SSH-reachable name and its DATA-PLANE IP -- not the management +# NIC. The legs advertise these to each other for the KV handoff. +export PREFILL_NODE="${PREFILL_NODE:-}" +export DECODE_NODE="${DECODE_NODE:-}" +export PREFILL_IP="${PREFILL_IP:-}" +export DECODE_IP="${DECODE_IP:-}" + +# --------------------------------------------------------------------------- +# 2. Image and weights (identical on BOTH nodes) +# --------------------------------------------------------------------------- +# The STOCK infera sglang image (deploy/docker/Dockerfile.sglang). GLM-5.3 big +# needs no source overlay -- that is only for the Flash family. +# +# TWO names for one image, and both are needed. preflight_rdma.sh reads IMAGE; +# engine/up.sh and common.sh require INFERA_IMAGE. Exporting only IMAGE makes +# `up` die at its first require_env, before a single container is started. +export IMAGE="${IMAGE:-}" +export INFERA_IMAGE="${INFERA_IMAGE:-$IMAGE}" +# GLM-5.3-MXFP4 or GLM-5.3. Resolve symlinks: where this path crosses an NFS +# mount boundary, bind-mounting the symlink's parent gives the container an +# empty directory, and the failure surfaces much later as an unrelated error. +export MODEL="${MODEL:-}" +export MODEL_MOUNT="${MODEL_MOUNT:-$(dirname "$MODEL")}" # same path on both nodes +export SERVED="${SERVED:-glm-5.3-mxfp4}" + +# --------------------------------------------------------------------------- +# 3. Transport +# --------------------------------------------------------------------------- +# From `preflight_rdma.sh mode`, run on BOTH nodes. Do not guess these. +# MC_GID_INDEX is PER NODE -- two identical machines routinely expose the +# routable GID at different indices, because an empty slot on one shifts +# everything after it. A wrong index fails loudly at init on every DP rank +# (`GID is NULL`). The link-local fe80:: GID is never the answer. Also: a down +# rail must not be listed, and the two nodes can legitimately differ. +export RDMA_IB_DEVICES="${RDMA_IB_DEVICES:-}" +export MC_GID_INDEX="${MC_GID_INDEX:-}" +# preflight_rdma.sh recommends RDMAV_FORK_SAFE=1 in all three of its modes, and +# engine/leg.sh honours it only when it is passed in +# ([ "${RDMAV_FORK_SAFE:-0}" = "1" ]) -- so leaving it unset here silently drops +# that recommendation. It is deliberately NOT defaulted on for the two-node +# shape: that shape was validated end to end WITHOUT it, and a validated path +# should not move on the strength of a recommendation alone. Uncomment to adopt +# it, as its own single-variable round. +# export RDMAV_FORK_SAFE=1 +# Leave MC_MS_FILTERS unset in mode A (peer-mem present). It is required only in +# the dma-buf mode, where KV must be pinned to one ODP-capable card. +# export MC_MS_FILTERS="ionic_0" + +# --------------------------------------------------------------------------- +# 4. Shape +# --------------------------------------------------------------------------- +# A whole node per leg. +export TP="${TP:-8}" +# Left UNSET on purpose, now that engine/up.sh actually forwards these. A whole +# node per leg means leg.sh's own `seq 0..TP-1` is already right and stays right +# if TP changes; pinning a literal 0..7 here would silently contradict TP=4. +# Set them only to place a leg on a specific subset of a node's cards. +export PREFILL_GPUS="${PREFILL_GPUS:-}" +export DECODE_GPUS="${DECODE_GPUS:-}" + +# Separate hosts, so these only need to be free on their own node. +export PREFILL_PORT="${PREFILL_PORT:-30000}" +export DECODE_PORT="${DECODE_PORT:-30001}" +export BOOTSTRAP_PORT="${BOOTSTRAP_PORT:-8998}" +export ROUTER_PORT="${ROUTER_PORT:-8100}" +export ETCD_PORT="${ETCD_PORT:-12379}" + +# --------------------------------------------------------------------------- +# 5. Features -- see the README before changing these two +# --------------------------------------------------------------------------- +# MTP off: upstream's GLM-5.3 cookbook disables EAGLE on AMD while the vendor +# model card runs it at 3 steps. The GLM-5.2 kit runs it ON and that is its +# validated configuration, so this is the one place this wrapper deliberately +# departs from it. Turn it on as its own round, and watch acceptance length -- +# a steady 4.00 is a repetition loop, not a good result. +# EVERY FEATURE KNOB THE KIT READS IS PER LEG. This file previously exported +# single knobs -- MTP, DPA -- which engine/up.sh reads under no name at all, so +# each silently did nothing and fell back to a plausible default. The single +# knobs are kept only as a convenience SEED for the per-leg pair below; the +# per-leg names are what the kit actually consumes. +export MTP="${MTP:-0}" +export PREFILL_MTP="${PREFILL_MTP:-$MTP}" +export DECODE_MTP="${DECODE_MTP:-$MTP}" +# DPA means DECODE-side DP-attention; prefill stays pure TP, as in the GLM-5.2 +# kit. This one was silently CORRECT before -- up.sh defaults happen to be 0/1 -- +# which is worse than visibly broken. +export DPA="${DPA:-1}" +export PREFILL_DPA="${PREFILL_DPA:-0}" +export DECODE_DPA="${DECODE_DPA:-$DPA}" +export KVAWARE="${KVAWARE:-1}" +export PREFILL_KVD="${PREFILL_KVD:-0}" +export DECODE_KVD="${DECODE_KVD:-0}" +# Insurance, not a fix -- this checkpoint's shared experts are themselves MXFP4. +# Kept on because upstream #25261 shows the mismatch failing SILENTLY with wrong +# output when shapes happen to line up. See the README. +export EXTRA_ENGINE_ARGS="${EXTRA_ENGINE_ARGS:---disable-shared-experts-fusion}" + +# Prefill wants activation headroom, decode wants KV pool. Do not equalise them. +export GMU_PREFILL="${GMU_PREFILL:-0.70}" +export GMU_DECODE="${GMU_DECODE:-0.85}" + +for v in PREFILL_IP IMAGE INFERA_IMAGE MODEL RDMA_IB_DEVICES MC_GID_INDEX; do + case "${!v}" in "<"*) echo "edit $(basename "$0"): $v is still a placeholder" >&2; exit 2;; esac +done + +exec bash "$KIT_DIR/engine/${1:?usage: $(basename "$0") up|smoke|bench|down}.sh" "${@:2}" diff --git a/examples/sglang_1p1d_glm5.3/cluster.singlenode.sh b/examples/sglang_1p1d_glm5.3/cluster.singlenode.sh new file mode 100755 index 000000000..5d611d8f6 --- /dev/null +++ b/examples/sglang_1p1d_glm5.3/cluster.singlenode.sh @@ -0,0 +1,184 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# ============================================================================ +# EDIT THIS FILE. Nothing else needs changing. +# ============================================================================ +# +# GLM-5.3 (big) 1P1D on a SINGLE 8-GPU node: TP4 prefill on cards 0-3, TP4 +# decode on cards 4-7, KV moved between them over mooncake. +# +# UNVALIDATED. The two-node shape is validated for GLM-5.2; this one is not, and +# its load-bearing unknown is whether mooncake moves KV between two legs on the +# SAME host, and at what speed. Read the README's "single-node unknown" section +# before trusting any number this produces. `smoke` checks MC_FORCE_TCP and +# GID is NULL in both leg logs precisely because a silent fall back to TCP is +# the failure mode here. +# +# It drives the GLM-5.2 kit's engine scripts unchanged -- GLM-5.3 (big) is the +# same architecture, so there is nothing to fork. See the README. +# +# Usage: bash cluster.singlenode.sh up | smoke | bench [conc...] | down +set -euo pipefail +HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" + +# The kit whose engine/ and common.sh actually run. Do not point this at a copy. +export KIT_DIR="${KIT_DIR:-$(cd "$HERE/../sglang_1p1d_glm5.2" && pwd)}" + +# --------------------------------------------------------------------------- +# 1. Node -- the same host twice +# --------------------------------------------------------------------------- +# Both legs live here, so PREFILL_* and DECODE_* name one machine. up.sh reaches +# each "node" through $SSH_CMD; pointing both at this host is what makes the +# single-node shape work with no change to the engine scripts. +export PREFILL_NODE="${PREFILL_NODE:-$(hostname -s)}" +export DECODE_NODE="$PREFILL_NODE" +# DATA-PLANE IP, not the management NIC -- the legs advertise it to each other. +export PREFILL_IP="${PREFILL_IP:-}" +export DECODE_IP="$PREFILL_IP" + +# --------------------------------------------------------------------------- +# 2. Image and weights +# --------------------------------------------------------------------------- +# The STOCK infera sglang image (deploy/docker/Dockerfile.sglang). GLM-5.3 big +# needs no source overlay -- that is only for the Flash family. +# +# TWO names for one image, and both are needed. preflight_rdma.sh reads IMAGE; +# engine/up.sh and common.sh require INFERA_IMAGE. Exporting only IMAGE makes +# `up` die at its first require_env, before a single container is started. +export IMAGE="${IMAGE:-}" +export INFERA_IMAGE="${INFERA_IMAGE:-$IMAGE}" +# GLM-5.3-MXFP4 or GLM-5.3. Resolve symlinks: where this path crosses an NFS +# mount boundary, bind-mounting the symlink's parent gives the container an +# empty directory, and the failure surfaces much later as an unrelated error. +export MODEL="${MODEL:-}" +export MODEL_MOUNT="${MODEL_MOUNT:-$(dirname "$MODEL")}" +export SERVED="${SERVED:-glm-5.3-mxfp4}" + +# --------------------------------------------------------------------------- +# 3. Transport +# --------------------------------------------------------------------------- +# From `preflight_rdma.sh mode`, run on THIS node. Do not guess these. +# MC_GID_INDEX is per node: the link-local fe80:: GID is never the answer. +# SINGLE-NODE RULE, and it differs from the two-node one: pin ONE device, the +# same on both legs. preflight_rdma.sh states it directly under mode A -- +# "all active rails (cross-node; mooncake pairs by GID subnet). Single-node +# loopback: pin ONE device on both legs instead." Listing all rails here is a +# cross-node recipe applied to a shape that is not cross-node. +# +# This matters even when the hip transport is enabled below: if hip fails to +# install, the fallback is loopback RDMA, and this is what decides whether that +# fallback is correctly configured or not. +export RDMA_IB_DEVICES="${RDMA_IB_DEVICES:-}" +export MC_GID_INDEX="${MC_GID_INDEX:-}" +# preflight_rdma.sh asks for this in ALL THREE modes, and engine/leg.sh honours +# it only when it is passed in ([ "${RDMAV_FORK_SAFE:-0}" = "1" ]). Unset here it +# never reaches the engine, so the preflight's own recommendation would be +# silently dropped. +export RDMAV_FORK_SAFE="${RDMAV_FORK_SAFE:-1}" +# Leave MC_MS_FILTERS unset in mode A (peer-mem present). It is required only in +# the dma-buf mode, where KV must be pinned to one ODP-capable card. +# export MC_MS_FILTERS="ionic_0" + +# THE flag that decides how single-node KV actually moves, and the only place it +# is turned on. engine/leg.sh defaults it to 1 (hip transport OFF), which is +# right for the two-node shape and wrong here: with hip absent the local segment +# advertises "rdma" only, and KV between two legs on ONE host takes LOOPBACK +# RDMA. That path works, raises nothing, and is the silent-slow case -- the +# README's "it fails loudly" only holds when hip is installed. +# +# With hip installed, MultiTransport::selectTransport routes KV by fixed +# priority (hip 4 > cxl 3 > rdma 2 > tcp 1), so hip wins: hipIpcGetMemHandle on +# the exporter, hipIpcOpenMemHandle on the importer, GPU-to-GPU over XGMI with +# no NIC in the path. That is what the pinned mooncake commit 01d1eb2a exists +# for. Measured on this hardware: HIP IPC works across DISJOINT +# HIP_VISIBLE_DEVICES -- an importer that cannot see the exporter's physical GPU +# mapped its memory and read back the correct bytes. +# +# REQUIRED CHECK after bring-up: `HIP transport installed for intra-node GPU +# P2P` must appear in BOTH leg logs, and `hipIpcOpenMemHandle failed` in +# neither. The MC_FORCE_TCP and GID-is-NULL counters do NOT cover this -- there +# is no log line for a TCP fallback, and a same-host hip transfer never touches +# a GID. +export MC_DISABLE_HIP_TRANSPORT="${MC_DISABLE_HIP_TRANSPORT:-0}" +# ...except that name is DEAD. Measured: the shipped mooncake contains +# MC_DISABLE_HIP (1 exact match) and does NOT contain MC_DISABLE_HIP_TRANSPORT +# or MC_ENABLE_HIP_TRANSPORT (0 each); a leg launched with +# MC_DISABLE_HIP_TRANSPORT=1 in /proc//environ still installed hip 4x. +# So hip is ON here whatever the line above says, and the line above is kept +# only because it reads as the pair to this one. +# +# MC_DISABLE_HIP is the live knob and is left UNSET on purpose: hip on is what +# this shape wants. Set MC_DISABLE_HIP=1 for a hip-off A/B, and confirm the +# discriminator actually flipped (4 -> 0 "HIP transport installed" lines per +# leg) BEFORE benchmarking -- otherwise the arm measures nothing and a +# guaranteed-zero differential reads as a null result. +export MC_DISABLE_HIP="${MC_DISABLE_HIP:-}" + +# --------------------------------------------------------------------------- +# 4. Shape +# --------------------------------------------------------------------------- +export TP="${TP:-4}" +export PREFILL_GPUS="${PREFILL_GPUS:-0,1,2,3}" +export DECODE_GPUS="${DECODE_GPUS:-4,5,6,7}" + +# Ports must not collide -- both legs share one host's network namespace. CHECK +# with `ss -lnt`; on a shared node the obvious ones are often taken. +export PREFILL_PORT="${PREFILL_PORT:-30000}" +export DECODE_PORT="${DECODE_PORT:-30001}" +export BOOTSTRAP_PORT="${BOOTSTRAP_PORT:-8998}" +export ROUTER_PORT="${ROUTER_PORT:-8100}" +export ETCD_PORT="${ETCD_PORT:-12379}" + +# The KV-event ports are PER LEG, and on this shape they must differ. In the +# two-node case both legs take engine/leg.sh's defaults (5557/8801) and never +# meet; here they share one network namespace, so the second leg's bind fails +# and the leg never serves -- the same "port_base at N is not available" that +# common.sh's reap() warns about across restarts, happening across legs instead. +# engine/up.sh forwards these per leg. +export PREFILL_KV_PUB_PORT="${PREFILL_KV_PUB_PORT:-5557}" +export PREFILL_KV_SNAP_PORT="${PREFILL_KV_SNAP_PORT:-8801}" +export DECODE_KV_PUB_PORT="${DECODE_KV_PUB_PORT:-5558}" +export DECODE_KV_SNAP_PORT="${DECODE_KV_SNAP_PORT:-8802}" + +# --------------------------------------------------------------------------- +# 5. Features -- see the README before changing these two +# --------------------------------------------------------------------------- +# EVERY FEATURE KNOB THE KIT READS IS PER LEG. This file previously exported +# single knobs -- MTP, DPA, GPUS, EXTRA_ENGINE_ARGS -- which engine/up.sh reads +# under no name at all, so each silently did nothing and fell back to a +# plausible default. MTP=0 still launched the decode leg at mtp=1. Keep the +# single knobs only as a convenience SEED for the per-leg pair below; the +# per-leg names are what the kit actually consumes. +# +# MTP off: upstream's GLM-5.3 cookbook disables EAGLE on AMD while the vendor +# model card runs it at 3 steps. Off avoids an unvalidated variable on a shape +# that is itself unvalidated. Turn it on deliberately, as its own round. +export MTP="${MTP:-0}" +export PREFILL_MTP="${PREFILL_MTP:-$MTP}" +export DECODE_MTP="${DECODE_MTP:-$MTP}" +# DPA means DECODE-side DP-attention; prefill stays pure TP, as in the GLM-5.2 +# kit. This one was silently CORRECT before -- up.sh defaults happen to be 0/1 -- +# which is worse than visibly broken: setting DPA=0 for a single-variable round +# would have produced dp8 anyway with nothing said. +export DPA="${DPA:-1}" +export PREFILL_DPA="${PREFILL_DPA:-0}" +export DECODE_DPA="${DECODE_DPA:-$DPA}" +export KVAWARE="${KVAWARE:-1}" +export PREFILL_KVD="${PREFILL_KVD:-0}" +export DECODE_KVD="${DECODE_KVD:-0}" +# Insurance, not a fix -- this checkpoint's shared experts are themselves MXFP4. +# Kept on because upstream #25261 shows the mismatch failing SILENTLY with wrong +# output when shapes happen to line up. See the README. +export EXTRA_ENGINE_ARGS="${EXTRA_ENGINE_ARGS:---disable-shared-experts-fusion}" + +# Prefill wants activation headroom, decode wants KV pool. Do not equalise them. +export GMU_PREFILL="${GMU_PREFILL:-0.70}" +export GMU_DECODE="${GMU_DECODE:-0.85}" + +for v in PREFILL_IP IMAGE INFERA_IMAGE MODEL RDMA_IB_DEVICES MC_GID_INDEX; do + case "${!v}" in "<"*) echo "edit $(basename "$0"): $v is still a placeholder" >&2; exit 2;; esac +done + +exec bash "$KIT_DIR/engine/${1:?usage: $(basename "$0") up|smoke|bench|down}.sh" "${@:2}" diff --git a/examples/sglang_mix_glm5.3/README.md b/examples/sglang_mix_glm5.3/README.md new file mode 100644 index 000000000..2b4ae4e30 --- /dev/null +++ b/examples/sglang_mix_glm5.3/README.md @@ -0,0 +1,225 @@ +# GLM-5.3 series — SGLang MIX (aggregated) on MI355X + +Runnable deployment kit for the **four GLM-5.3 checkpoints** served the infera +way on a single 8×MI355X (gfx950) node: one aggregated worker, prefix caching +on, fronted by the **infera kv-aware router**. No PD, no RDMA, no second node. + +One file is site-specific. Fill it in and the deployment is three commands. + +```bash +$EDITOR env.sh # the only file you edit +bash engine/up.sh # container -> etcd -> worker -> router +bash engine/smoke.sh # prove each feature is actually live +``` + +## Read this first: GLM-5.3 is two unrelated architectures + +They share a product name and almost nothing else. Which one you are serving +decides **which engine image you need**, and getting it wrong fails at config +load rather than at inference. + +| | `GLM-5.3`, `GLM-5.3-MXFP4` | `GLM-5.3-Flash`, `GLM-5.3-Flash-MXFP4` | +|---|---|---| +| `model_type` | `glm_moe_dsa` | `glm5_next` | +| hidden / layers | 6144 / 78 | 4096 / 45 | +| routed experts | 256 | 288 | +| attention | uniform MLA + DSA | **hybrid**: KDA linear attention + DSA | +| also carries | — | mHC, no-RoPE MLA, compressed indexer k-pool, a vision encoder | +| memory pools | paged KV only | paged KV **plus a KDA state pool** | +| **image** | `Dockerfile.sglang` | **`Dockerfile.sglang.glm53`** | + +**The big pair is GLM-5.2 with different weights.** Their `config.json` is +identical to GLM-5.2's field for field except `transformers_version`, so the +released engine already serves them through `glm4_moe.py`. + +**The Flash pair exists in no released sglang.** `glm5_next.py` is absent from +`lmsysorg/sglang:v0.5.18-rocm720-mi35x` and from the vendor-validated +`lmsysorg/sglang-rocm:v0.5.18-rocm724-mi35x-20260822`. `Dockerfile.sglang.glm53` +brings the engine source in at build time, pinned to a SHA. Point a Flash +variant at the ordinary image and you get: + +``` +ValueError: The checkpoint you are trying to load has model type `glm5_next` +but Transformers does not recognize this architecture. +``` + +That message names `transformers`, which invites the wrong fix. Upgrading +transformers does not help — the branch registers its own `Glm5NextConfig` and +pins the same version. The missing component is sglang. + +`glm5_next` is the **model body**, not an MTP head: `glm5_next.py` exports +`EntryClass = [Glm5NextForConditionalGeneration]` and MTP lives separately in +`glm5_next_nextn.py`. Turning speculative decoding off does not remove the need +for the overlay. + +## Contents + +| path | what | +|---|---| +| [`env.sh`](env.sh) | **the only file you edit** — variant, IP, weights, image, shape, ports | +| `engine/worker.sh` | the real launcher; carries the tuned recipe for all four variants, no site values | +| `engine/up.sh` | container → etcd → worker → router, waiting on health at each step | +| `engine/smoke.sh` | six blocks, each red when a specific feature is *silently* absent | +| `engine/bench.sh` | reference fixed-length sweep via sglang's own `bench_serving` | +| `engine/down.sh` | tear down, then **wait** for VRAM to actually drain | + +## Validation status + +Stated plainly rather than implied. All on 8×MI355X, gfx950, ROCm 7.2, driver +6.14.14, TP4, decode CUDA graphs on unless noted. + +| variant | status | +|---|---| +| `flash-mxfp4` | **validated end to end** — full infera stack, reproduced on two separate nodes, plus a fixed-length sweep | +| `big-fp8` | **validated** — all smoke blocks green, `max_total_num_tokens=1148288` | +| `big-mxfp4` | **validated, with numbers** — all smoke blocks green; AITER FP4 path confirmed dispatching (`torch.float4_e2m1fn_x2`, `per_1x32`) rather than dequantising to BF16; TP8+DPA+MTP fixed-length sweep lands at **0.92 / 1.06 / 0.89 / 1.10 ×** the GLM-5.2 MIX baseline at concurrency 1/8/16/24 | +| `flash-fp8` | **validated, with numbers** — brought up on a third node, 62/62 shards, coherent answers, 8/8 AITER mHC lines; TP4 fixed-length sweep at isl 7400 / osl 320 gives **99.70 tok/s at conc 1** and **456.68 at conc 8** (TPOT 9.20 / 13.50 ms) | +| PD (1P1D) for any variant | **not covered by this kit.** For the big pair the shape is the same as [`sglang_1p1d_glm5.2`](../sglang_1p1d_glm5.2/), which is validated for GLM-5.2 | + +## The one flag you must not drop: `--disable-shared-experts-fusion` + +On `flash-mxfp4` this is load-bearing, not tuning. sglang PR #36607 opened the +gfx950 branch of `glm5_next`'s shared-experts fusion gate +(`glm5_next.py:1414`) **without** carrying the +`quant_blocks_shared_experts_fusion(quant_config)` guard that +`deepseek_v2.py:3069` has. `QuarkConfig.can_fuse_shared_expert()` computes the +correct answer and is never consulted. The checkpoint keeps its shared experts +in BF16 while its routed experts are MXFP4, so the shared expert gets renamed +into routed slot 288 of a packed FusedMoE and weight load dies: + +``` +RuntimeError: The size of tensor a (256) must match the size of tensor b (512) +at non-singleton dimension 1 # fused_moe_triton/layer.py::_load_w2 +``` + +256 is the MXFP4-packed width; 512 is the same tensor unpacked and TP4-sharded. +Upstream issue **#37268** is the identical failure on NVFP4/NVIDIA with the same +accepted workaround. + +Two consequences worth carrying: + +- **The health signal is a line that must be ABSENT.** Grep the worker log for + `Shared experts fusion optimization enabled.` — present means broken. +- **On `flash-fp8` the flag is NOT needed, and this is measured rather than + assumed.** That checkpoint was brought up on a third node with fusion left + **enabled** — `disable_shared_experts_fusion: False` in the resolved args and + `[TP0] Shared experts fusion optimization enabled.` in the log — and it loaded + all 62 shards with no `_load_w2` mismatch and answered correctly. The reason it + is safe was predicted before launch from the checkpoint's own index: **129 + `.weight` / 129 `.weight_scale_inv`**, a strict 1:1 pairing, i.e. uniformly + block-FP8. The MXFP4 precondition — shared experts at a *higher* precision than + the routed experts — is simply absent, so `quant_blocks_shared_experts_fusion` + returns False and fusion is legitimate. + + This is the control arm that makes the story a quantization mismatch rather + than "fusion is broken on gfx950". **The gfx950 fusion path itself works.** What + #36607 shipped unguarded is the *decision* to use it, and that decision is only + wrong when the checkpoint is mixed-precision — which is why the one-line guard + is the right upstream fix rather than reverting the feature. + +- **On `big-mxfp4` the same flag is insurance, not a fix.** That checkpoint's + shared experts are themselves MXFP4, so the precondition is absent. It stays + on by default because upstream #25261 shows this class of mismatch failing + *silently with wrong output* rather than crashing when the shapes happen to + line up. Set `SHARED_EXPERT_FUSION=1` for a clean single-variable perf round. + +## Notes and gotchas + +**1. The AITER mHC lines are the real health check, not the absence of errors.** +On the Flash variants, grep the worker log for two `AITER gfx950 mHC` lines per +rank. Without them the server still starts, still answers correctly, and is +**4.3–5.4× slower**, with nothing in any log calling it out. They are gated on +HIP + gfx95 + `SGLANG_USE_AITER=1`; miss the env and you silently get the slow +path. `smoke.sh` block 5 counts them. + +**2. The Flash family keeps TWO memory pools.** Decode lines must show +`full token usage` *and* `mamba usage`. If `max_running_requests` is clamped at +startup, suspect the KDA state pool first: raise `--mamba-full-memory-ratio` +(default 0.9) or pin `--max-mamba-cache-size`. Do **not** override +`linear_lower_bound` through `--json-model-override-args`. + +**3. DSA flags are `--dsa-*`, not GLM-5.2's `--nsa-*`.** Both spellings exist in +v0.5.18; the `--nsa-*` ones are not what this model wants. + +**4. The DSA-on-ROCm env block is mandatory for the big pair.** Without it the +model serves, returns 200s, and returns garbage — the sparse-attention indexer +takes a path not ported to gfx950. `worker.sh` sets it; `infera.engine.sglang` +also defaults `SGLANG_OPT_USE_TOPK_V2` off on ROCm. + +**5. Do not copy the vendor card's `--cuda-graph-max-bs 2 +--max-running-requests 2`.** Those appear in the published GLM-5.3-MXFP4 recipe +and cap the server at two concurrent requests. That is an accuracy +configuration, not a throughput one. + +**6. Benchmark with `bench_serving`, not a shell loop.** A bash fan-out of +concurrent `curl`s becomes the bottleneck before the engine does — at +concurrency 32 one measured 350 output tok/s while the engine's own log reported +2398 tok/s with an empty queue. + +**7. Resolve the weights symlink yourself.** Where the models path crosses an +NFS mount boundary, bind-mounting the symlink's parent gives the container an +empty directory. The failure appears minutes later as `Unrecognized processing +class`, because `config.json` is the one file that still resolves. `up.sh` binds +`realpath` output. + +**8. MTP/EAGLE is off for both families, and that is unresolved rather than +decided.** Upstream's GLM-5.3 cookbook disables speculative decoding on AMD +because the gfx950 draft kernel is unvalidated; the OneNexus big-MXFP4 card runs +EAGLE at `--speculative-num-steps 3`. Both statements are recorded here; neither +has been tested in this kit. + +**9. `Ctrl-C` on a log tail does not stop anything.** Use `engine/down.sh`, and +check `docker ps`. + +## Reference numbers + +`flash-mxfp4`, TP4 (four GPUs), decode graphs on, `bench_serving` at +ISL 7400 / OSL 320: + +| concurrency | output tok/s | TTFT p50 | +|---:|---:|---:| +| 1 | 111.0 | 255 ms | +| 8 | 561.0 | 1065 ms | +| 16 | 962.6 | 744 ms | +| 24 | 1391.2 | 619 ms | + +These are a reference point on one configuration, not a claim about the family. +In particular they are **not** comparable to GLM-5.2 numbers taken at TP8 with +DP-attention, MTP and kvd enabled: different model, different parallelism, +different feature set. + +## What this kit exposes on the network, and why you should care on a shared node + +These scripts run with `--network=host` and bind on **all interfaces**. That is +deliberate for a single-tenant benchmark box and is **wrong for a shared one** — +and the reference cluster is shared. Three unrelated users logged into one of +these nodes during a single day of this work, and a colleague's job claimed all +eight GPUs of another without notice. + +| what | where | exposed as | +|---|---|---| +| **etcd**, unauthenticated | `up.sh:64-65` | `0.0.0.0:$ETCD_PORT` and peer port | +| router | `up.sh:96` | `0.0.0.0:$ROUTER_PORT` | +| KV event / snapshot sockets | `worker.sh:208` | `tcp://0.0.0.0:$KV_PUB_PORT` | + +**etcd is the one that matters.** It holds worker discovery, it has no auth, and +anyone who can reach the port can mutate the keys that tell the router where to +send traffic. The engine and KV ports are lower stakes but still raw. + +The scripts are left as they were validated rather than quietly re-bound, because +changing a bind changes a recipe that was measured. **If you run this on a shared +or reachable host, change them yourself:** the router and workers are colocated, +so `127.0.0.1` works for etcd and the KV sockets, and only the router port needs +to be reachable — put something that authenticates in front of it. + +**`--trust-remote-code` (`worker.sh:216`) is not removable.** `glm5_next` is not +in any released `transformers`, so the checkpoint's own code must load. The risk +is therefore in the checkpoint directory, not the flag: on the reference cluster +`/apps/data/models` is a shared NFS mount, so **anyone who can write there can +execute code in your container.** If that matters to you, copy the checkpoint to +a directory you own and verify permissions before launching. + +## Source + +[`examples/sglang_mix_glm5.3/`](.) in [AMD-AGI/Infera](https://github.com/AMD-AGI/Infera) +· [GLM-5.2 1P1D kit](../sglang_1p1d_glm5.2/) · [PD disaggregation concepts](../../manual/features/pd_disaggregation.md) diff --git a/examples/sglang_mix_glm5.3/engine/bench.sh b/examples/sglang_mix_glm5.3/engine/bench.sh new file mode 100755 index 000000000..64e114373 --- /dev/null +++ b/examples/sglang_mix_glm5.3/engine/bench.sh @@ -0,0 +1,83 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# Reference fixed-length sweep using sglang's own bench_serving, which ships +# inside the image. +# +# Use this rather than a shell fan-out of curl. A bash loop of N concurrent +# curls becomes the bottleneck well before the engine does: at concurrency 32 +# one measured 350 output tok/s while the engine's own log reported 2398 tok/s +# at #running-req 32 with an empty queue. The client was the limit, not the server. +# +# Three flags are load-bearing and are NOT defaults: +# --random-range-ratio 1.0 pins every prompt to exactly ISL. The default +# draws uniformly and the percentiles then mix +# request sizes; a fixed-length sweep wants a +# delta, not a distribution. +# --temperature 1.0 --top-p 0.95 the checkpoint's own generation_config +# defaults, deliberately NOT greedy. At temperature +# 0 this reasoning model falls into repetition on a +# long prompt and the run reads like corruption. +# --num-prompts 10 x conc enough requests per arm to reach steady state. +# +# KNOWN LIMITATION — long OSL produces degenerate output. Measured, not feared. +# +# The default OSL of 320 is clean. Raise it and the generations collapse into +# repetition loops, at a rate that grows with length. Measured on this stack with +# a 10-gram-repeated->=5x check, on both GLM-5.3 and GLM-5.2: +# +# osl 320 0-1 % of requests +# osl 3300 40-65 % +# osl 17000 96-100 % (worst: one 10-gram repeated 16,990 times) +# +# The tok/s arithmetic is NOT affected -- degenerate requests were measured to +# decode at the same speed as clean ones (+0.2 % at concurrency 1). What is wrong +# is *what the tokens are*. So a long-OSL number here is a valid throughput +# measurement of the engine generating repetitive text, and must not be quoted as +# a quality result. +# +# Cause is NOT established, and one obvious candidate is already ELIMINATED. +# +# It is tempting to blame missing conversational framing, because the args dump +# shows `apply_chat_template=False`. **That reading is wrong.** On this backend +# the flag is parsed and never consumed -- `grep -n apply_chat_template` in +# sglang's `benchmark/serving.py` returns one hit, an assignment that fires only +# for the image/mmmu datasets. And `--backend sglang-oai-chat` posts `messages` +# to `/v1/chat/completions`, which applies the model's chat template server-side +# by definition. **The template has been applied all along.** A flag that is +# parsed but unread reads exactly like a setting. +# +# What remains: the prompt CONTENT is still synthetic (seeded ShareGPT filler, +# repeated and truncated to reach the target ISL), which is a different thing +# from missing framing. And the ROCm silent greedy fallback in EAGLE verify +# (`eagle_utils.py:726`, `_is_hip` in an `or` with `is_all_greedy`, so +# temperature and top_p never reach token selection -- three open upstream PRs, +# #31214 / #32922 / #37134, none merged). Note every arm measured as degenerate +# ran MTP ON, so that path was live in all of them. +# +# If you need long-OSL numbers you can defend on quality, score the saved +# generations for n-gram repetition (`--output-details` saves `generated_texts`) +# rather than trusting the summary. Adding `--apply-chat-template` will NOT help +# -- see above, it does nothing on this backend. +# ACCEPTANCE LENGTH DOES NOT DETECT THIS: an aggregate of 2.96 with 1.25 % at the +# 4.00 ceiling was measured while 54 % of the very same requests were looping. +# +# Usage: bash engine/bench.sh [conc ...] (default 1 8 16 24) +set -uo pipefail +KIT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +# shellcheck source=../env.sh +source "$KIT/env.sh" +SERVED="${SERVED:-glm5.3-$VARIANT}" +ISL="${ISL:-7400}"; OSL="${OSL:-320}" +CONCS=("$@"); [ ${#CONCS[@]} -eq 0 ] && CONCS=(1 8 16 24) + +for c in "${CONCS[@]}"; do + echo "===== isl=$ISL osl=$OSL conc=$c =====" + docker exec "$CTR" python3 -m sglang.bench_serving --backend sglang-oai-chat \ + --host "$MY_IP" --port "$PORT" --model "$SERVED" \ + --dataset-name random --random-input-len "$ISL" --random-output-len "$OSL" \ + --random-range-ratio 1.0 --max-concurrency "$c" --num-prompts $((c * 10)) \ + --temperature 1.0 --top-p 0.95 2>&1 \ + | grep -E "Successful requests|Benchmark duration|Request throughput|Output token throughput|Total token throughput|Median TTFT|P99 TTFT|Mean TPOT|Median E2E" +done diff --git a/examples/sglang_mix_glm5.3/engine/down.sh b/examples/sglang_mix_glm5.3/engine/down.sh new file mode 100755 index 000000000..f574b04b7 --- /dev/null +++ b/examples/sglang_mix_glm5.3/engine/down.sh @@ -0,0 +1,29 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# Tear down and WAIT for VRAM to actually drain. The wait is the point, not the +# kill: `docker rm -f` returns before the driver has reclaimed the allocations, +# and relaunching too early aborts the next distributed bootstrap with a +# misleading "memory capacity is unbalanced" error that reads like a model or +# config problem. +set -uo pipefail +KIT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +# shellcheck source=../env.sh +source "$KIT/env.sh" + +# Named explicitly, never a pattern -- on a shared node a pattern is how you +# remove somebody else's container. +docker rm -f "$CTR" "${CTR}_etcd" >/dev/null 2>&1 +echo "removed $CTR and ${CTR}_etcd; waiting for VRAM to drain" +for i in $(seq 1 60); do + busy=$(rocm-smi --showmemuse 2>/dev/null | awk -v want=",$GPUS," ' + match($0, /GPU\[([0-9]+)\]/, m) && /VRAM%/ { + split($0, f, ": "); if (index(want, "," m[1] ",") && f[2]+0 > 5) printf "%s ", m[1] }') + [ -z "$busy" ] && { echo "GPUs $GPUS at baseline after $((i * 5))s"; exit 0; } + sleep 5 +done +echo "TIMEOUT: GPU(s) $busy still above baseline." >&2 +echo "If they are not ours, that is expected on a shared node -- check first." >&2 +rocm-smi --showpids 2>/dev/null | sed -n '3,12p' >&2 +exit 1 diff --git a/examples/sglang_mix_glm5.3/engine/smoke.sh b/examples/sglang_mix_glm5.3/engine/smoke.sh new file mode 100755 index 000000000..b07b524ce --- /dev/null +++ b/examples/sglang_mix_glm5.3/engine/smoke.sh @@ -0,0 +1,81 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# Prove the deployment is real, not merely alive. +# +# A green /health says a process is listening. It does not say the AITER fast +# path is dispatching, that both memory pools exist, or that the model is +# producing sense rather than noise. Each block below is chosen because it goes +# RED when a specific feature is silently absent. +set -uo pipefail +KIT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +# shellcheck source=../env.sh +source "$KIT/env.sh" +SERVED="${SERVED:-glm5.3-$VARIANT}" +R="http://$MY_IP:$ROUTER_PORT" + +echo "===== 1. router sees exactly one MIXED worker =====" +curl -s "$R/v1/workers" | python3 -c ' +import json,sys +d=json.load(sys.stdin); w=d if isinstance(d,list) else d.get("workers",d) +w=w if isinstance(w,list) else [] +print(f" workers: {len(w)} (want 1)") +for x in w: print(f" {x.get('"'"'url'"'"',x.get('"'"'worker_id'"'"'))} disagg_mode={x.get('"'"'disagg_mode'"'"')} (want mixed)") +' + +echo "===== 2. model is served under the expected name =====" +curl -s "$R/v1/models" | python3 -c 'import json,sys; print(" ", [m["id"] for m in json.load(sys.stdin)["data"]])' + +echo "===== 3. coherent answer, and reasoning separated =====" +# max_tokens is deliberately generous. GLM-5.3 is a thinking model and the leg +# passes --reasoning-parser glm45, so the chain of thought lands in +# reasoning_content but is billed against the SAME budget. At a small value the +# model spends every token thinking, content comes back empty with +# finish_reason "length", and a healthy deployment reads as a failure. +# +# Garbage or repeated tokens here is NOT a sampling problem. On the big +# variants it is the signature of the DSA-on-ROCm env block not taking effect. +curl -s "$R/v1/chat/completions" -H 'Content-Type: application/json' -d "{ + \"model\": \"$SERVED\", + \"messages\": [{\"role\": \"user\", \"content\": \"In two sentences, explain why prefix caching helps agentic workloads.\"}], + \"max_tokens\": 600, \"temperature\": 0}" | python3 -c ' +import json,sys +d=json.load(sys.stdin); m=d["choices"][0]["message"] +print(" content :", (m.get("content") or "")[:220]) +print(" reasoning_content chars:", len(m.get("reasoning_content") or ""), "(want > 0)") +print(" finish :", d["choices"][0].get("finish_reason"), "| usage:", d["usage"])' + +echo "===== 4. arithmetic, and instruction-following =====" +curl -s "$R/v1/chat/completions" -H 'Content-Type: application/json' -d "{ + \"model\": \"$SERVED\", + \"messages\": [{\"role\": \"user\", \"content\": \"What is 17 * 23? Reply with only the number.\"}], + \"max_tokens\": 512, \"temperature\": 0}" \ + | python3 -c 'import json,sys; print(" answer:", repr(json.load(sys.stdin)["choices"][0]["message"]["content"]), "(want 391)")' + +echo "===== 5. engine-side evidence =====" +docker exec "$CTR" bash -c ' +L=/tmp/glm53_mix.log +# THE health check that matters. Two AITER mHC lines per rank. Their absence +# does not raise: the server answers correctly and is several times slower, +# with nothing anywhere saying so. Flash variants only. +echo " AITER mHC lines : $(grep -c "mHC" $L) (flash: want 2 per rank)" +# Must be ABSENT. Present means shared-experts fusion is on, which on a +# mixed-precision Quark checkpoint mis-loads the BF16 shared expert into a +# packed routed slot. See the note in worker.sh. +echo " fusion-enabled line : $(grep -c "Shared experts fusion optimization enabled" $L) (want 0 on mxfp4)" +# The flash family keeps TWO pools: the paged KV pool and a KDA state pool. +# Both must appear, or the linear-attention path is not what you think it is. +echo " decode lines w/ 2 pools: $(grep -c "full token usage.*mamba usage" $L)" +# The KDA state pool caps concurrency independently of --max-running-requests. +echo " resolved max_running : $(grep -o "max_running_requests is capped to [0-9]*" $L | tail -1)" +echo " memory access fault : $(grep -c "memory access fault" $L) (want 0)" +echo " HIP error : $(grep -c "HIP error" $L) (want 0)" +# torch._dynamo/metrics_context tracebacks are compile-telemetry noise and +# appear in healthy runs; excluded here so a real one is visible. +echo " Traceback (non-dynamo) : $(grep "Traceback" $L | grep -vc "_dynamo") (want 0)" +grep -m1 "max_total_num_tokens" $L | cut -c1-150 +' 2>/dev/null + +echo "===== 6. router policy =====" +docker exec "$CTR" grep -om1 "kv-aware" /tmp/router.log 2>/dev/null | sed 's/^/ policy: /' diff --git a/examples/sglang_mix_glm5.3/engine/up.sh b/examples/sglang_mix_glm5.3/engine/up.sh new file mode 100755 index 000000000..a41cf8943 --- /dev/null +++ b/examples/sglang_mix_glm5.3/engine/up.sh @@ -0,0 +1,106 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# Bring up a single-node MIX (aggregated) GLM-5.3 deployment: +# container -> etcd -> infera worker -> kv-aware router. +# Reads every site value from ../env.sh. Runs ON the node. +set -uo pipefail +KIT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +# shellcheck source=../env.sh +source "$KIT/env.sh" + +case "$MY_IP" in "<"*) echo "edit env.sh: MY_IP is still a placeholder" >&2; exit 2;; esac +case "$IMAGE" in "<"*) echo "edit env.sh: IMAGE is still a placeholder" >&2; exit 2;; esac +[ -d "$MODEL" ] || { echo "MODEL not a directory on this host: $MODEL" >&2; exit 2; } + +# Bind the REALPATH, not the symlink's parent. See the note in env.sh: on the +# reference cluster the models path crosses an NFS mount boundary, and binding +# the wrong side gives the container an empty directory whose failure surfaces +# many minutes later as an unrelated-looking processor error. +MODEL_REAL="$(realpath "$MODEL")" +MOUNT_REAL="$(dirname "$MODEL_REAL")" + +echo "===== 1. teardown (ours only) =====" +# Named explicitly. Never a pattern: on a shared node a pattern is how you +# remove somebody else's container. +docker rm -f "$CTR" "${CTR}_etcd" >/dev/null 2>&1 +sleep 5 + +echo "===== 2. GPUs $GPUS free? =====" +# Advisory, not a kill. If a GPU we want is busy this stops and says so -- it +# never reclaims a process, because on a shared node that process is somebody's +# multi-day job. +busy=$(rocm-smi --showmemuse 2>/dev/null | awk -v want=",$GPUS," ' + match($0, /GPU\[([0-9]+)\]/, m) && /VRAM%/ { + split($0, f, ": "); if (index(want, "," m[1] ",") && f[2]+0 > 5) printf "%s ", m[1] }') +[ -n "$busy" ] && { echo " GPU(s) $busy already in use -- pick a free subset via GPUS=, or wait." >&2; exit 1; } +echo " clear" + +echo "===== 3. container =====" +docker run -d --name "$CTR" --network=host --ipc=host --shm-size=64G \ + --device=/dev/kfd --device=/dev/dri \ + --group-add video --group-add render --cap-add=SYS_PTRACE --cap-add=IPC_LOCK \ + --security-opt seccomp=unconfined --ulimit memlock=-1:-1 \ + -v "$MOUNT_REAL":"$MOUNT_REAL":ro \ + "$IMAGE" sleep infinity >/dev/null || exit 1 +sleep 5 + +echo "===== 4. etcd =====" +# Two etcd traps, both of which have cost time here. +# +# 1. The etcd v3.5.x image has an empty ENTRYPOINT and Cmd=[/usr/local/bin/etcd]; +# passing `etcd` as argv[0] dumps usage and exits 2. Hence --entrypoint. +# 2. The THREE peer flags below must move together. Override only +# --listen-peer-urls and etcd exits 1 with "--initial-cluster has +# default=http://localhost:2380 but missing from +# --initial-advertise-peer-urls". A script that hardcodes the default 2380 +# never sees this -- which is exactly why it bites the first time you land on +# a node where 2380 is taken, and it has bitten two operators here +# independently. +docker run -d --name "${CTR}_etcd" --network=host --entrypoint /usr/local/bin/etcd \ + quay.io/coreos/etcd:v3.5.14 \ + --advertise-client-urls "http://$MY_IP:$ETCD_PORT" \ + --listen-client-urls "http://0.0.0.0:$ETCD_PORT" \ + --listen-peer-urls "http://0.0.0.0:$((ETCD_PORT + 1))" \ + --initial-advertise-peer-urls "http://$MY_IP:$((ETCD_PORT + 1))" \ + --initial-cluster "default=http://$MY_IP:$((ETCD_PORT + 1))" >/dev/null +sleep 5 +curl -sf -m5 "http://$MY_IP:$ETCD_PORT/version" >/dev/null \ + && echo " etcd up" || { echo " ETCD FAILED"; docker logs --tail 5 "${CTR}_etcd"; } + +echo "===== 5. worker =====" +docker cp "$KIT/engine/worker.sh" "$CTR":/worker.sh >/dev/null +docker exec -d "$CTR" env \ + MY_IP="$MY_IP" ETCD_IP="$MY_IP" ETCD_PORT="$ETCD_PORT" \ + MODEL="$MODEL_REAL" VARIANT="$VARIANT" TP="$TP" GPUS="$GPUS" PORT="$PORT" \ + SERVED="${SERVED:-glm5.3-$VARIANT}" CUDA_GRAPH="${CUDA_GRAPH:-1}" \ + bash /worker.sh +echo " launching -> /tmp/glm53_mix.log in $CTR" + +echo "===== 6. wait for /health =====" +# Cold start is minutes, not seconds: several hundred GB of weights, then AITER +# JIT, then graph capture. Silence is not a hang. 650 s observed on a node with +# a cold NFS cache. +for i in $(seq 1 240); do + if docker exec "$CTR" curl -sf -m3 "http://$MY_IP:$PORT/health" >/dev/null 2>&1; then + echo " serving after $((i * 10))s"; break + fi + docker exec "$CTR" pgrep -f infera.engine.sglang >/dev/null 2>&1 || { + echo " worker died -- last lines:"; docker exec "$CTR" tail -25 /tmp/glm53_mix.log; exit 1; } + sleep 10 +done + +echo "===== 7. kv-aware router =====" +docker exec "$CTR" bash -c "printf '%s\n' '#!/bin/bash' \ + 'exec python3 -m infera.server --host 0.0.0.0 --port $ROUTER_PORT \ + --discovery-backend etcd --etcd-endpoint $MY_IP:$ETCD_PORT \ + --request-transport http --kv-event-transport zmq --router-policy kv-aware \ + --router-tokenizer-path $MODEL_REAL' > /run_router.sh && chmod +x /run_router.sh" +docker exec -d "$CTR" bash -c 'nohup /run_router.sh > /tmp/router.log 2>&1' +sleep 20 +docker exec "$CTR" bash -c \ + "curl -sf -m5 http://$MY_IP:$ROUTER_PORT/health >/dev/null && echo ' router healthy' \ + || { echo ' router not ready'; tail -20 /tmp/router.log; }" + +echo "===== up. clients use http://$MY_IP:$ROUTER_PORT =====" diff --git a/examples/sglang_mix_glm5.3/engine/worker.sh b/examples/sglang_mix_glm5.3/engine/worker.sh new file mode 100755 index 000000000..8546dfb8c --- /dev/null +++ b/examples/sglang_mix_glm5.3/engine/worker.sh @@ -0,0 +1,220 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# The real launcher, and the file that carries the tuned recipe. Runs INSIDE +# the container, staged there by up.sh. Every site-specific value arrives as an +# env var from env.sh; there are no addresses or paths in here. +# +# Launches through `python3 -m infera.engine.sglang` rather than +# sglang.launch_server so infera's etcd discovery and kv-aware routing work. +set -u +MY_IP="${MY_IP:?MY_IP=this node data-plane IP}" +MODEL="${MODEL:?MODEL=weights dir inside the container}" +VARIANT="${VARIANT:?VARIANT=flash-mxfp4|flash-fp8|big-mxfp4|big-fp8}" +ETCD_IP="${ETCD_IP:-$MY_IP}" +ETCD_PORT="${ETCD_PORT:-12379}" +PORT="${PORT:-30000}" +TP="${TP:-4}" +GPUS="${GPUS:-$(seq -s, 0 $((TP - 1)))}" +SERVED="${SERVED:-glm5.3-$VARIANT}" +LOG="${LOG:-/tmp/glm53_mix.log}" + +# Decode CUDA graphs. ON by default: measured 12.5 -> 117.9 output tok/s at +# concurrency 1 on flash-mxfp4 and 12-15 -> ~110 on flash-fp8, both ~7.5x, for +# ~33-82 s of capture. The bs list is graph COVERAGE, not a concurrency cap -- a +# decode batch is padded up to the next captured size and anything larger runs +# eager. Sizes above --max-running-requests are dropped at capture time. +# +# --cuda-graph-max-bs is DEPRECATED on v0.5.18 in favour of +# --cuda-graph-max-bs-decode; it still resolves today but will break on a base +# bump. This kit passes --cuda-graph-backend-decode / --cuda-graph-bs-decode, +# which are the current spelling. The vendor recipes still use the old one. +# +# BUDGET THE VRAM AT TP4. Capture cost was measured at 15.4-17.4 GB per rank at +# TP4, against ~1.4 GB in an earlier TP8 bring-up of the same model -- fewer +# ranks means each one's graphs cover a larger shard. Do not carry a TP8 figure +# into a TP4 plan. +# Prefill graphs stay disabled: that is what upstream validated on gfx950, and +# prefill is where the DSA/KDA shape variance lives. +CUDA_GRAPH="${CUDA_GRAPH:-1}" +GRAPH_BS="${GRAPH_BS:-1 2 4 8 16 24 32 48 64 96 128}" +KVAWARE="${KVAWARE:-1}" + +# Cache-hit accounting. OFF by default because it is not free, but you must turn +# it ON for any workload whose result depends on prefix reuse -- an agentic +# replay, or anything passing bench_serving's --cache-report. +# +# Without it the server still answers normally and simply reports nothing: +# `usage.prompt_tokens_details` comes back **null** through the router and the +# engine logs `#cached-token: 0`. That is indistinguishable from a genuine 0 % +# hit rate, so a cache-sensitive benchmark reads as "the cache never worked" +# rather than "the counter was never enabled". Verified on this stack. +CACHE_REPORT="${CACHE_REPORT:-0}" + +# --- common ROCm env -------------------------------------------------------- +# SGLANG_USE_AITER gates the AITER fast paths on gfx950. On the flash variants +# it gates PR #36607's mHC pre/post dispatch specifically, and its absence is +# SILENT: the server starts, answers correctly, and is 4.3-5.4x slower with +# nothing in any log saying so. smoke.sh greps for the mHC lines for that reason. +export SGLANG_USE_AITER=1 +export SAFETENSORS_FAST_GPU=1 HIP_FORCE_DEV_KERNARG=1 HSA_NO_SCRATCH_RECLAIM=1 +export NCCL_IGNORE_CPU_AFFINITY=1 +# Stable block hashes -> stable kv-aware keys across restarts. +export PYTHONHASHSEED=0 +export SGLANG_HOST_IP="$MY_IP" HOST_IP="$MY_IP" +export INFERA_SGLANG_READY_TIMEOUT="${READY_TIMEOUT:-3600}" +NIC=$(ip -o -4 addr show | awk -v ip="$MY_IP" '$4 ~ ("^" ip "/") {print $2; exit}') +[ -n "$NIC" ] && export SGLANG_LOCAL_IP_NIC="$NIC" GLOO_SOCKET_IFNAME="$NIC" + +ARGS=() +case "$VARIANT" in + flash-*) + # --- glm5_next family --------------------------------------------------- + # DSA flags are --dsa-*-backend here. GLM-5.2 used --nsa-*; carrying that + # spelling forward gets unknown-flag errors. + # THE KDA-POOL CLAMP IS REAL AND IT FIRES. This family keeps a second + # memory pool for the linear-attention state, and the scheduler will cap + # concurrency against it regardless of what you ask for here. Observed on + # every rank at TP4: + # max_running_requests is capped to 200 by the mamba state cache + # (max_mamba_cache_size=1000, 5 state slots per request). To raise it: + # increase --mamba-full-memory-ratio or --max-mamba-cache-size, or halve + # the state size with --mamba-ssm-dtype bfloat16. + # Read the RESOLVED value out of the worker log, not out of server_args -- + # server_args records what was requested, and reading it instead is how one + # bring-up concluded the clamp had not fired when it had. + ARGS+=(--dsa-prefill-backend tilelang --dsa-decode-backend tilelang + --kv-cache-dtype "${KV_DTYPE:-bfloat16}" + --context-length "${CTX:-65536}" + --mem-fraction-static "${GMU:-0.80}" + --max-running-requests "${MAX_RUNNING:-32}" + --chunked-prefill-size "${CHUNK:-4096}" + --max-prefill-tokens "${MAX_PREFILL:-16384}" + --mm-feature-transport cpu) + + # --disable-shared-experts-fusion is LOAD-BEARING on mxfp4, not tuning. + # sglang PR #36607 opened the gfx950 branch of glm5_next's fusion gate + # (glm5_next.py:1414) without carrying the + # quant_blocks_shared_experts_fusion(quant_config) guard that + # deepseek_v2.py:3069 has, and QuarkConfig.can_fuse_shared_expert() -- which + # computes the right answer -- is never consulted. The checkpoint's BF16 + # shared expert is then renamed into routed slot 288 of an MXFP4-packed + # FusedMoE and weight load dies with + # RuntimeError: The size of tensor a (256) must match the size of tensor b + # (512) at non-singleton dimension 1 + # in fused_moe_triton/layer.py::_load_w2. Upstream issue #37268 is the same + # bug on NVFP4/NVIDIA and documents the same workaround. + # Set SHARED_EXPERT_FUSION=1 to re-enable, e.g. on a uniformly-quantized + # checkpoint where fusion is both correct and profitable. + SHARED_EXPERT_FUSION="${SHARED_EXPERT_FUSION:-0}" + [ "$SHARED_EXPERT_FUSION" = "0" ] && ARGS+=(--disable-shared-experts-fusion) + + if [ "$VARIANT" = "flash-mxfp4" ]; then + # Quark MXFP4 (fp4 E2M1, 1x32 block scales). --quantization is explicit + # per the vendor model card; the aiter runner dispatches native FP4 MoE + # kernels (torch.float4_e2m1fn_x2). With `triton` the checkpoint is + # dequantized to BF16 GEMMs -- it still serves, and it is much slower. + ARGS+=(--quantization "${QUANT:-quark}" --moe-runner-backend "${MOE_RUNNER:-aiter}") + # Vendor-set for this checkpoint and absent from the FP8 recipe. Not noise. + export SGLANG_OPT_DEEPGEMM_HC_PRENORM=0 + else + # FP8 original: config.json already carries the quantization, so no + # --quantization flag. + ARGS+=(--moe-runner-backend "${MOE_RUNNER:-triton}") + fi + ;; + + big-*) + # --- glm_moe_dsa family: the GLM-5.2 code path -------------------------- + # MANDATORY on gfx950. Without this env block the model serves, returns + # 200s, and returns GARBAGE, because the sparse-attention indexer takes a + # path not ported to this architecture. infera.engine.sglang already + # defaults SGLANG_OPT_USE_TOPK_V2 off on ROCm + # (infera/engine/rocm_dsa_env.py); it is repeated so that a bare + # launch_server run of this same recipe behaves identically. + export SGLANG_ROCM_FUSED_DECODE_MLA=0 SGLANG_OPT_USE_TILELANG_INDEXER=1 + export SGLANG_OPT_USE_TOPK_V2=0 SGLANG_OPT_USE_JIT_NORM=0 + + # READING THE STARTUP LINE UNDER DP-ATTENTION: the engine prints PER-RANK + # values while /get_server_info reports the GLOBAL ones. Ask for + # --max-running-requests 256 --chunked-prefill-size 65536 at dp8 and the + # startup line says 32 and 8192 -- that is 256/8 and 65536/8, a DIVISION, + # not a clamp, even though the line reads exactly like one. Confirm against + # /get_server_info before investigating a cap that is not there. + # + # --ep-size is emitted unconditionally and OUTSIDE any DP-attention branch: + # expert parallelism and attention parallelism are different axes, and + # gating both on one condition silently collapses the MoE whenever DPA is + # off, after which no latency delta is attributable to either. + # --max-running-requests is passed EXPLICITLY rather than left to the + # engine's memory-derived default. The default is not wrong, but it is + # derived from whatever VRAM happens to be free, so two runs of the same + # recipe on differently-loaded nodes silently get different admission + # limits -- and a benchmark then measures the limit, not the engine. + ARGS+=(--ep-size "${EP_SIZE:-$TP}" + --dsa-prefill-backend tilelang --dsa-decode-backend tilelang + --kv-cache-dtype "${KV_DTYPE:-fp8_e4m3}" + --context-length "${CTX:-262144}" + --max-running-requests "${MAX_RUNNING:-32}" + --mem-fraction-static "${GMU:-0.80}" + --chunked-prefill-size "${CHUNK:-65536}") + + # The aiter custom all-reduce kernel deadlocks on this architecture under + # speculative verify. Disabled independently of MTP so that any "MTP on vs + # off" comparison stays a one-variable comparison. + ARGS+=(--disable-custom-all-reduce) + + if [ "$VARIANT" = "big-mxfp4" ]; then + # Quantization is AUTO-DETECTED from config.json; the vendor card states + # no --quantization flag is required. + ARGS+=(--moe-runner-backend "${MOE_RUNNER:-aiter}") + # Insurance rather than a fix: glm4_moe.py:1174's fusion gate only + # special-cases w4afp8 and would fuse under quark, but this checkpoint's + # shared experts are themselves MXFP4, so the precondition is absent. + # Kept on because upstream #25261 shows this class of mismatch failing + # SILENTLY with wrong output rather than crashing. Set + # SHARED_EXPERT_FUSION=1 for a clean single-variable performance round. + [ "${SHARED_EXPERT_FUSION:-0}" = "0" ] && ARGS+=(--disable-shared-experts-fusion) + fi + ;; + *) echo "unknown VARIANT: $VARIANT" >&2; exit 2 ;; +esac + +# MTP/EAGLE is deliberately NOT enabled for either family. Upstream's GLM-5.3 +# cookbook disables speculative decoding on AMD because the gfx950 draft kernel +# is unvalidated, while the OneNexus big-MXFP4 card runs EAGLE at 3 steps. That +# contradiction is recorded rather than resolved; do not add --speculative-* +# without re-deriving it. + +# kvd / hierarchical cache is OFF. On gfx950 (xnack-) hicache stores raw host +# data_ptr()s that a GPU kernel dereferences while hipHostRegister maps those +# pages at a different device VA, and the process aborts with +# "Memory access fault by GPU node-N on address ". The fix lives in +# patches/sglang_rocm/; confirm it is in your image before turning either on. + +[ "$CACHE_REPORT" = "1" ] && ARGS+=(--enable-cache-report) + +if [ "$CUDA_GRAPH" = "1" ]; then + ARGS+=(--cuda-graph-backend-decode full --cuda-graph-backend-prefill disabled + --cuda-graph-bs-decode $GRAPH_BS) +else + ARGS+=(--cuda-graph-backend-decode disabled --cuda-graph-backend-prefill disabled) +fi + +INFERA_ARGS=(--advertise-host "$MY_IP" --etcd-endpoint "$ETCD_IP:$ETCD_PORT" + --discovery-backend etcd --request-transport http --kv-event-transport zmq) +if [ "$KVAWARE" = "1" ]; then + INFERA_ARGS+=(--kv-events-bind "tcp://0.0.0.0:${KV_PUB_PORT:-5557}" + --kv-snapshot-port "${KV_SNAP_PORT:-8801}") +else + INFERA_ARGS+=(--no-enable-kv-events) +fi + +echo "[glm53-mix] variant=$VARIANT ip=$MY_IP:$PORT tp=$TP gpus=$GPUS graph=$CUDA_GRAPH -> $LOG" +HIP_VISIBLE_DEVICES="$GPUS" python3 -m infera.engine.sglang \ + --model-path "$MODEL" --served-model-name "$SERVED" --tp-size "$TP" --trust-remote-code \ + --host "$MY_IP" --port "$PORT" \ + --watchdog-timeout 3600 \ + --reasoning-parser glm45 --tool-call-parser glm47 \ + "${ARGS[@]}" "${INFERA_ARGS[@]}" > "$LOG" 2>&1 diff --git a/examples/sglang_mix_glm5.3/env.sh b/examples/sglang_mix_glm5.3/env.sh new file mode 100755 index 000000000..fe3ad7ae0 --- /dev/null +++ b/examples/sglang_mix_glm5.3/env.sh @@ -0,0 +1,71 @@ +#!/usr/bin/env bash +# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. +# SPDX-License-Identifier: MIT +# +# ============================================================================ +# EDIT THIS FILE. Nothing else in this kit needs changing for your site. +# ============================================================================ +# +# Usage: bash engine/up.sh | engine/smoke.sh | engine/bench.sh [conc...] | engine/down.sh +# All of them source this file. + +# --------------------------------------------------------------------------- +# 1. Which checkpoint +# --------------------------------------------------------------------------- +# flash-mxfp4 | flash-fp8 | big-mxfp4 | big-fp8 +# +# THE VARIANT DECIDES WHICH IMAGE YOU NEED, because GLM-5.3 is two unrelated +# architectures sharing a product name: +# +# flash-* model_type glm5_next, GlmMoeDsaForCausalLM's opposite number -- +# hybrid KDA-linear + DSA attention, mHC, natively multimodal. +# Exists in NO released sglang. Needs the image built from +# deploy/docker/Dockerfile.sglang.glm53, which overlays sglang +# PR #36607 at a pinned SHA. +# big-* model_type glm_moe_dsa. Field-for-field identical to GLM-5.2 +# except transformers_version, so the released engine already +# serves it. Needs the ordinary deploy/docker/Dockerfile.sglang +# image. +# +# Crossing them fails at CONFIG LOAD, not at inference: +# ValueError: The checkpoint you are trying to load has model type +# `glm5_next` but Transformers does not recognize this architecture. +# That message names transformers and invites the wrong fix. The missing +# component is sglang, not transformers. +export VARIANT="${VARIANT:-flash-mxfp4}" + +# --------------------------------------------------------------------------- +# 2. Site +# --------------------------------------------------------------------------- +# This node's DATA-PLANE IP -- not the management NIC. Clients and the router +# reach the worker here. +export MY_IP="${MY_IP:-}" + +# Weights. MODEL must be the directory, MODEL_MOUNT its parent. +# +# Resolve symlinks yourself. On the reference cluster /apps/data/models is a +# symlink to /perf_apps/data/models and /perf_apps is a SEPARATE NFS mount, so +# bind-mounting the symlink's parent gives the container an empty directory and +# every path under it dangles. The failure surfaces far downstream as +# "Unrecognized processing class", because config.json is the one file docker +# still resolves. up.sh binds `realpath` output for this reason. +export MODEL="${MODEL:-/apps/data/models/GLM-5.3-Flash-MXFP4}" + +# Engine image. Must match VARIANT -- see section 1. +export IMAGE="${IMAGE:-}" + +# --------------------------------------------------------------------------- +# 3. Shape +# --------------------------------------------------------------------------- +# TP4 is what AMD validated for the Flash MXFP4 checkpoint, and it leaves four +# GPUs free for a second arm on an 8-GPU node. TP8 works too; raise GPUS with it. +export TP="${TP:-4}" +export GPUS="${GPUS:-0,1,2,3}" + +# Ports. CHECK THESE ARE FREE (`ss -lnt`) -- on a shared node they often are +# not. 2379/2380 in particular are frequently held by somebody else's etcd. +export ETCD_PORT="${ETCD_PORT:-12379}" +export PORT="${PORT:-30000}" +export ROUTER_PORT="${ROUTER_PORT:-8100}" + +export CTR="${CTR:-glm53_mix}" diff --git a/tests/e2e/harness/matrix.py b/tests/e2e/harness/matrix.py index 63c3011e8..d0888a26e 100644 --- a/tests/e2e/harness/matrix.py +++ b/tests/e2e/harness/matrix.py @@ -91,6 +91,19 @@ DEEPSEEK_V4_PRO = "deepseek-ai/DeepSeek-V4-Pro" GLM_5_1_FP8 = "zai-org/GLM-5.1-FP8" +# GLM-5.3 is TWO unrelated architectures sharing a product name, and the split +# decides which engine image a case needs: +# * GLM-5.3 / GLM-5.3-MXFP4 model_type glm_moe_dsa, GlmMoeDsaForCausalLM. +# Field-for-field identical to GLM-5.2 except transformers_version, so the +# released engine already serves them via glm4_moe.py -- Dockerfile.sglang. +# * GLM-5.3-Flash / -Flash-MXFP4 model_type glm5_next, hybrid KDA-linear + +# DSA + mHC, natively multimodal. In NO released sglang; needs the build-time +# source overlay in Dockerfile.sglang.glm53. +GLM_5_3 = "zai-org/GLM-5.3" +GLM_5_3_MXFP4 = "OneNexus/GLM-5.3-MXFP4" +GLM_5_3_FLASH = "zai-org/GLM-5.3-Flash" +GLM_5_3_FLASH_MXFP4 = "OneNexus/GLM-5.3-Flash-MXFP4" + EXTRA_ARGS: dict[str, tuple[str, ...]] = {} # default verbatim extra launch args # config.json keys that mark a Mixture-of-Experts model (any present with an diff --git a/tests/e2e/pd_mixed/sglang/matrix.py b/tests/e2e/pd_mixed/sglang/matrix.py index dc7d243ac..eee95683f 100644 --- a/tests/e2e/pd_mixed/sglang/matrix.py +++ b/tests/e2e/pd_mixed/sglang/matrix.py @@ -10,7 +10,17 @@ import pytest -from ...harness.matrix import DEEPSEEK_V4_PRO, GLM_5_1_FP8, GPT_OSS, KIMI_K26_MXFP4, expand_cases +from ...harness.matrix import ( + DEEPSEEK_V4_PRO, + GLM_5_1_FP8, + GLM_5_3, + GLM_5_3_FLASH, + GLM_5_3_FLASH_MXFP4, + GLM_5_3_MXFP4, + GPT_OSS, + KIMI_K26_MXFP4, + expand_cases, +) # [enable, model, tp, ep, dp_attn] (+ optional opts dict). A tuple/list on an axis # enumerates it (e.g. (True, False) runs both). MoE models can exercise ep. @@ -124,6 +134,218 @@ "server_ready_timeout": 1800, }, ], + # ---- GLM-5.3 series ----------------------------------------------------- + # All four are parked (enable=False) ON PURPOSE, and not because they are + # unproven: each recipe below was brought up and smoke-checked on 8xMI355X. + # They stay off because each needs ~300-700 GB of weights pre-staged and 4 + # GPUs for 10+ minutes of cold start, which no CI runner here has. Flipping + # the first field to True is the whole activation step. + # + # TWO IMAGES. The Flash rows need deploy/docker/Dockerfile.sglang.glm53 + # (glm5_next exists in no released sglang; that file overlays sglang PR + # #36607 at c821c425). The two big rows run on the stock + # deploy/docker/Dockerfile.sglang image. Enabling a Flash row against the + # wrong image fails at CONFIG LOAD with "model type `glm5_next` but + # Transformers does not recognize this architecture" -- which names + # transformers and invites the wrong fix; the missing piece is sglang. + # + # resolve_model() maps these ids to /. Pre-staged + # trees here are flat (.../GLM-5.3-Flash-MXFP4), so symlink the vendor + # prefix or the id falls back to a Hub pull of several hundred GB. + [ + False, + GLM_5_3_FLASH_MXFP4, + 4, + False, + False, + { + # --disable-shared-experts-fusion is LOAD-BEARING, not tuning. PR + # #36607 opened the gfx950 branch of glm5_next's fusion gate + # (glm5_next.py:1414) without carrying deepseek_v2.py:3069's + # quant_blocks_shared_experts_fusion() guard, so this checkpoint's + # BF16 shared expert is renamed into routed slot 288 of an + # MXFP4-packed FusedMoE and weight load dies with + # "size of tensor a (256) must match tensor b (512)". + # Upstream #37268 is the same bug on NVFP4/NVIDIA, same workaround. + # Health check: "Shared experts fusion optimization enabled." must + # be ABSENT from the worker log. + "args": [ + "--quantization", + "quark", + "--moe-runner-backend", + "aiter", + "--kv-cache-dtype", + "bfloat16", + "--dsa-prefill-backend", + "tilelang", + "--dsa-decode-backend", + "tilelang", + "--disable-shared-experts-fusion", + "--reasoning-parser", + "glm45", + "--tool-call-parser", + "glm47", + "--mm-feature-transport", + "cpu", + "--context-length", + "65536", + "--mem-fraction-static", + "0.80", + "--max-running-requests", + "32", + "--chunked-prefill-size", + "4096", + "--max-prefill-tokens", + "16384", + ], + # SGLANG_USE_AITER gates #36607's AITER mHC dispatch. Without it the + # server starts, answers correctly, and is 4.3-5.4x slower with + # nothing in any log saying so -- grep the worker log for two + # "AITER gfx950 mHC" lines per rank. + # SGLANG_OPT_DEEPGEMM_HC_PRENORM=0 is vendor-set for this + # checkpoint and is absent from the FP8 recipe; not noise. + "env": {"SGLANG_USE_AITER": "1", "SGLANG_OPT_DEEPGEMM_HC_PRENORM": "0"}, + # 650 s cold start observed on a node with a cold NFS cache. + "server_ready_timeout": 3600, + }, + ], + [ + False, + GLM_5_3_FLASH, + 4, + False, + False, + { + # FP8 original. No --quantization (config.json carries fp8) and the + # triton MoE runner rather than aiter's FP4 path. + # UNVERIFIED at time of writing: whether this checkpoint also needs + # --disable-shared-experts-fusion. It does IFF its shared experts + # are kept at a higher precision than its routed experts -- count + # mlp.shared_experts .weight vs .weight_scale* in the safetensors + # index. Add the flag if load dies in _load_w2/_load_w13 with a + # 2:1 shape mismatch. + "args": [ + "--moe-runner-backend", + "triton", + "--kv-cache-dtype", + "bfloat16", + "--dsa-prefill-backend", + "tilelang", + "--dsa-decode-backend", + "tilelang", + "--reasoning-parser", + "glm45", + "--tool-call-parser", + "glm47", + "--context-length", + "65536", + "--mem-fraction-static", + "0.85", + ], + "env": {"SGLANG_USE_AITER": "1"}, + "server_ready_timeout": 3600, + }, + ], + [ + False, + GLM_5_3_MXFP4, + 4, + True, + False, + { + # glm_moe_dsa -- the GLM-5.2 code path, stock image. Quantization is + # auto-detected from config.json; no --quantization needed. + # The DSA-on-ROCm env block is MANDATORY on gfx950: without it the + # model serves, returns 200s, and returns garbage, because the + # sparse-attention indexer takes a path not ported to this arch. + # infera.engine.sglang defaults SGLANG_OPT_USE_TOPK_V2 off on ROCm + # (infera/engine/rocm_dsa_env.py); it is repeated here so a bare + # launch_server run of this same row behaves identically. + # + # --disable-shared-experts-fusion is insurance rather than a fix: + # glm4_moe.py:1174's gate only special-cases w4afp8 and would fuse + # under quark, but this checkpoint's shared experts are themselves + # MXFP4 (76 .weight / 75 .weight_scale, the odd one being the BF16 + # MTP layer 78, which is not loaded while MTP is off), so the + # precondition is absent. Upstream #25261 shows this class failing + # SILENTLY with wrong output rather than crashing, which is why it + # is defaulted on. Drop it for a clean single-variable perf round. + # + # NOT the vendor card's --cuda-graph-max-bs 2 --max-running-requests + # 2: that is a concurrency-2 accuracy configuration, not a + # throughput one, and must not be copied into a benchmark arm. + "args": [ + "--kv-cache-dtype", + "fp8_e4m3", + "--moe-runner-backend", + "aiter", + "--dsa-prefill-backend", + "tilelang", + "--dsa-decode-backend", + "tilelang", + "--disable-shared-experts-fusion", + "--disable-custom-all-reduce", + "--reasoning-parser", + "glm45", + "--tool-call-parser", + "glm47", + "--context-length", + "262144", + "--mem-fraction-static", + "0.80", + "--chunked-prefill-size", + "65536", + ], + "env": { + "SGLANG_USE_AITER": "1", + "SGLANG_ROCM_FUSED_DECODE_MLA": "0", + "SGLANG_OPT_USE_TILELANG_INDEXER": "1", + "SGLANG_OPT_USE_TOPK_V2": "0", + "SGLANG_OPT_USE_JIT_NORM": "0", + }, + "server_ready_timeout": 3600, + }, + ], + [ + False, + GLM_5_3, + 4, + True, + False, + { + # FP8 original of the big model. Same code path as the MXFP4 row; + # only the weights and the absent quantization flag differ. 704 GB + # at TP4 leaves ~55 GB per GPU for KV at GMU 0.80 -- measured, not + # estimated (max_total_num_tokens=1148288). + "args": [ + "--kv-cache-dtype", + "fp8_e4m3", + "--dsa-prefill-backend", + "tilelang", + "--dsa-decode-backend", + "tilelang", + "--disable-custom-all-reduce", + "--reasoning-parser", + "glm45", + "--tool-call-parser", + "glm47", + "--context-length", + "262144", + "--mem-fraction-static", + "0.80", + "--chunked-prefill-size", + "65536", + ], + "env": { + "SGLANG_USE_AITER": "1", + "SGLANG_ROCM_FUSED_DECODE_MLA": "0", + "SGLANG_OPT_USE_TILELANG_INDEXER": "1", + "SGLANG_OPT_USE_TOPK_V2": "0", + "SGLANG_OPT_USE_JIT_NORM": "0", + }, + "server_ready_timeout": 3600, + }, + ], ]