docs(cookbook): add AMD MI300X/MI325X/MI355X support for GLM-5.2

[andyluo7]

Summary

Adds AMD MI300X / MI325X / MI355X support to the GLM-5.2 cookbook. GLM-5.2 shares the GLM-5.1 / DeepSeek-V3.2 (glm_moe_dsa) architecture, so the validated GLM-5.1 ROCm recipe carries over with only the model path swapped.

Changes

• Deploy panel (docs_new/src/snippets/configs/zai-org/glm-5.2.jsx):
  • mi300x, mi325x, mi355x added to supportedHardware + ROCm dockerImages.
  • Single-node TP8 FP8/BF16 cells using the DSA tilelang backend (--dsa-prefill-backend tilelang --dsa-decode-backend tilelang --chunked-prefill-size 131072 --mem-fraction-static 0.80 --watchdog-timeout 1200).
  • DSA-prefill Context Parallel disabled on AMD (verified on Hopper only).
• Page (docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx): description + an AMD bullet in Configuration Tips.

Notes

• No MTP on AMD (yet). EAGLE/MTP speculative decoding currently does not build for gfx950 (the spec-decode kernel fails to compile), so the AMD cells omit --speculative-* and serve without MTP. This is documented in Config Tips. Once the gfx950 spec-decode kernel lands, MTP cells can be added.
• Validation: GLM-5.2-FP8 served end-to-end on MI355X-class (gfx950) hardware at TP8 via the DSA tilelang path. MI300X/MI325X (gfx942) and BF16 cells mirror the GLM-5.1 recipe and are marked verified: false pending a run.

Test plan

• Render the GLM-5.2 cookbook page; confirm AMD appears in the Deploy panel and generates the expected ROCm sglang serve command.
• Serve zai-org/GLM-5.2-FP8 on MI300X/MI325X to flip those cells to verified: true.

🤖 Generated with Claude Code

---

CI States

Latest PR Test (Base): ✅ Run #28475719904
Latest PR Test (Extra): ❌ Run #28475719698

[gemini-code-assist]

Warning

You have reached your daily quota limit. Please wait up to 24 hours and I will start processing your requests again!

[zijiexia]

As AMD currently not support MTP, can we disable MTP in playground for AMD? Also the three strategies for amd seems identical.

[zijiexia]

Does the bf16 fit in single node mi300x? might need multinodes

[zijiexia]

Can we add disable reason for AMD?

[andyluo7]

Thanks for the review @zijiexia! Addressed all three points in 3530557:

1. Disable MTP for AMD — the Speculative card's mtp-516 and mtp-112 options now carry disable: { hw: ["mi355x","mi325x","mi300x"] } + a disableReason (gfx950 spec-decode kernel doesn't build; DSA nextn draft path is CUDA-only). They render disabled with the reason on AMD.

2. Three AMD strategies were identical — now differentiated via batch-shaping levers (no MTP needed):

   • low-latency: large --chunked-prefill-size 131072, default cuda-graph bs
   • balanced: --chunked-prefill-size 32768 + --cuda-graph-max-bs 128 + --max-running-requests 80
   • high-throughput: --cuda-graph-max-bs 256 + --max-running-requests 256
     Only the MI355X FP8 low-latency cell is benchmarked, so I kept that verified: true and set the inferred balanced/high-throughput cells to verified: false.

3. BF16 single-node fit — you're right, MI300X (1.5 TB HBM) can't hold the ~1.51 TB BF16 weights + KV. Removed the MI300X BF16 single-node cells (kept FP8). MI325X (2 TB) and MI355X (2.3 TB) BF16 single-node cells remain. Updated the cookbook MDX bullet to call this out.

Also extended the CP knob's disableReason to mention the AMD ROCm DSA-CP path isn't validated yet.

[andyluo7]

Heads-up: while validating this recipe I found a base-accuracy bug affecting GLM-5.2-FP8 on gfx950 (MI350X/MI355X) in current SGLang ROCm images, and updated the PR accordingly.

What's wrong: block-FP8 linear layers route to aiter's gemm_a8w8_blockscale_bpreshuffle GEMM on gfx950, which is numerically incorrect there (ROCm 7.2). The error is small per-layer but compounds across all 78 layers and silently corrupts output — in-context reasoning collapses while short factual prompts can still look right.

Evidence (MI350X, TP4, v0.5.13.post1-rocm720-mi35x):

• GSM8K (200Q, 5-shot): 0.000 / 72% invalid stock → 0.925 / 0% invalid with the fix
• Per-layer residual cosine vs vLLM (pos1): decays to 0.10 at the final layer stock → 1.00 with the fix
• The fix is one line — use the plain CK gemm_a8w8_blockscale (what vLLM uses) instead of the bpreshuffle variable: _use_aiter_bpreshuffle_gfx95 = False. It also restores EAGLE/MTP acceptance to 2.65–3.77 (was ~1.3), so MTP actually works on gfx950 once this is fixed.

Scope: gfx950 only (is_gfx95_supported()); gfx942 (MI300X/MI325X) is unaffected (uses the triton blockscale kernel).

Changes in this push:

• Added a <Warning> callout documenting the gfx950 accuracy bug (and that MI300X/MI325X are unaffected).
• Corrected the MTP note (MTP is degraded by this same bug on gfx950; works at num-steps ≤ 3 once fixed).
• Set all AMD cells to verified: false until a fixed image ships — the deployment configs themselves are correct.

I'll file a separate upstream issue for the kernel bug and link it here.

[andyluo7]

Filed the upstream kernel bug: #28685 (linked from the cookbook Warning callout).

[andyluo7]

@zijiexia all review points are addressed (MTP disabled for AMD, strategies differentiated, MI300X BF16 removed, CP/AMD disable reasons added), plus the AMD cells are now marked verified: false with a <Warning> documenting the gfx950 block-FP8 accuracy bug (tracked in #28685). Could you take another look when you have a moment? Also happy to hear your preference on whether to merge now (with the warning) or hold until the #28685 fix lands.

[zijiexia]

@zijiexia all review points are addressed (MTP disabled for AMD, strategies differentiated, MI300X BF16 removed, CP/AMD disable reasons added), plus the AMD cells are now marked verified: false with a <Warning> documenting the gfx950 block-FP8 accuracy bug (tracked in #28685). Could you take another look when you have a moment? Also happy to hear your preference on whether to merge now (with the warning) or hold until the #28685 fix lands.

Hi @andyluo7 , Any idea why #28423 works? Isn't they are the same architecture?

[andyluo7]

Update on the gfx950 accuracy bug behind the AMD verified: false cells: the upstream CK kernel fix is now open and validated — ROCm/rocm-libraries#8639 ("Deterministic blockscale bpreshuffle GEMM on gfx950 and avoid the illegal-type pattern").

Root cause confirmed: the CK gemm_a8w8_blockscale_bpreshuffle kernel used a packed 2-lane illegal-type FMA that relied on an LLVM coercion pass removed in ROCm 7.2 → miscompiled at -O2/-O3 (correct at -O1), producing non-deterministic wrong rows near tile boundaries. #8639 replaces it with a scalar FMA + a VGPR accumulator anchor.

I validated #8639 against this repro on MI350X/ROCm 7.2 (rebuilt the aiter bpreshuffle module at stock -O3): per-M-tile wrong rows gone (clean M=32..128), GLM-5.2 shapes cos=1.0000, and 30 repeated launches bitwise-identical. Details in #28685.

I've updated the cookbook <Warning> here to point at #8639 as the tracked fix. No cell changes yet — the warning + verified: false stay until #8639 lands in CK → aiter → a released lmsysorg/sglang-rocm image; once it ships we can re-benchmark gfx950 accuracy, flip the AMD cells to verified: true, and re-evaluate MTP.

[andyluo7]

Good question @zijiexia — I dug into #28423 (DeepSeek-V4 AMD cookbook). The short version: the bug isn't architecture-specific, it's a gfx950 block-FP8 GEMM kernel miscompile, so what matters is the quantization path and the GPU arch, not the model. #28423's verified AMD cells mostly avoid the exact path that breaks:

• MI300X (FP8) — gfx942, unaffected. The bpreshuffle bug is gfx950-only; gfx942 routes block-FP8 through the triton blockscale kernel. So [AMD] Update v4 amd cookbook #28423's MI300X FP8 cells working is fully consistent with my finding.
• MI355X (FP4) — different kernel. DeepSeek-V4's primary verified MI355X cells are quant: "fp4". FP4 goes through the FP4 GEMM path, not gemm_a8w8_blockscale_bpreshuffle, so it never touches the broken kernel.
• MI355X (FP8) — same path, would be affected. [AMD] Update v4 amd cookbook #28423 does have mi355x + fp8 verified: true cells (e.g. deepseek-v4.jsx L1503/1547/1595, L1787/1831/1879). Those route through the same gfx950 bpreshuffle kernel and would hit the same corruption. But [AMD] Update v4 amd cookbook #28423's verification was docs-rendering only ("Built the docs locally with mint dev… renders correctly"; the accuracy/speed checklist boxes are unchecked) — so it was very likely not GSM8K-tested on gfx950 FP8, and the issue would be silent there (short factual prompts look fine; only multi-step reasoning collapses). Might be worth a quick GSM8K spot-check on those V4 MI355X-FP8 cells too.

Why GLM-5.2 can't sidestep it the way V4 does: DeepSeek-V4 ships an FP4 checkpoint, so its recommended AMD path on gfx950 is FP4 (clean kernel). GLM-5.2 only ships FP8 and BF16 — no FP4 — so on MI350X/MI355X it's forced onto exactly the block-FP8 bpreshuffle path that miscompiles. That's why this surfaces sharply here.

This is now root-caused to a ROCm-7.2 toolchain regression and has an upstream CK fix open + validated: ROCm/rocm-libraries#8639 (tracked from #28685). Once it lands in a released lmsysorg/sglang-rocm image, I'll re-benchmark and flip the GLM-5.2 AMD cells to verified: true. Until then the <Warning> + verified: false are the safe state.

[sunxxuns]

SGLANG_FP8_PAGED_MQA_LOGITS_TORCH=0 enable long context correctness.

[andyluo7]

@richardchen @zijiexia — validated both issues on MI350X / MI355X (gfx950), image lmsysorg/sglang-rocm:v0.5.13.post1-rocm720-mi35x-20260618, and pushed the fixes (commit 05d1c29aa).

TL;DR: both bugs are clear on the 0618 image; no config/env change needed.

1. Long-context / DSA indexer (sunxxuns's SGLANG_FP8_PAGED_MQA_LOGITS_TORCH=0 note)

Swept three indexer paths at TP8, each with GSM8K-200 (5-shot) + needle-in-haystack to ~118K tokens:

Config	indexer	TORCH	GSM8K	NIAH
current PR (no indexer env)	deep_gemm (default)	=0	0.970	15/15
TORCH=1	torch ref	=1	0.965	15/15
full V4 aiter bundle	aiter	=1	0.955	15/15

All three retrieve perfectly to 118K tokens. TORCH=0 vs =1 makes no correctness difference — the PR's current default (=0) is already correct, and the V4 aiter indexer bundle is not required for GLM-5.2. So no env change.

Note: my first NIAH run showed 0/15 everywhere — that was a harness artifact (GLM-5.2 is a reasoning model and my max_tokens=16 was eaten by thinking tokens). Re-ran with enable_thinking:false → 15/15.

2. Block-FP8 bpreshuffle GEMM (#28685 / ROCm/rocm-libraries#8639)

Re-ran GSM8K at TP4 (the original repro config) and TP8 on 0618 with bpreshuffle on (default — _use_aiter_bpreshuffle_gfx95 still evaluates True):

• TP8: 0.970 (0% invalid) · TP4: 0.960 (0% invalid)

The catastrophic collapse (GSM8K≈0 / 72% invalid) no longer reproduces — the CK kernel fix has landed in the aiter/CK build inside the 0618 image, so the workaround (_use_aiter_bpreshuffle_gfx95=False) is no longer needed. The earlier TP4=0.000 was on an older image.

Changes in this push

• Pin mi355x image → ...mi35x-20260618.
• Flip MI355X FP8 low-latency to verified: true (the benchmarked cell, correctness re-confirmed).
• Downgrade the gfx950 <Warning> to a historical <Note> ("fixed as of 0618; older images affected").
• MTP note: block-FP8 dependency removed, but MTP stays disabled (its spec-decode kernel is a separate, not-yet-validated path).

I intentionally left MI355X FP8 balanced/high-throughput, all BF16, and all MI325X/MI300X cells at verified: false — those weren't benchmarked in this pass.

Caveat: I verified via GSM8K (short-ctx) + NIAH (long-ctx retrieval), not a full RULER/LongBench reasoning eval. If you want belt-and-suspenders before signing off, that's the remaining gap. Otherwise this should be good to merge once the run-ci labels are added (the red checks are just the label gate, not test failures) and the branch conflict is resolved.

[andyluo7]

Update: AMD MI355X FP8 now tuned for both latency (MTP-5) and throughput (fp8 KV)

Two additions on top of the base AMD recipe, both validated on MI350X (gfx950, TP8, image lmsysorg/sglang-rocm:v0.5.14-rocm720-mi35x-20260630):

1. EAGLE MTP 5-1-6 on the low-latency cell — enabled by #29373 (cuda_runtime.h guard, in v0.5.14) + --disable-overlap-schedule to work around a v0.5.14 spec-v2 ROCm DSA draft-extend bug (#29785). Accept length ~3–5, SGLANG_DSA_ENABLE_MTP_PRECOMPUTE_METADATA=1.

2. fp8 KV cache on all 3 FP8 cells — --kv-cache-dtype fp8_e4m3 + SGLANG_ROCM_FUSED_DECODE_MLA=0 (mirrors the InferenceX GLM-5.1 recipe). +21–46% output throughput, no accuracy change.

Accuracy (full GSM8K, 1319 questions): fp8 KV 0.946 vs bf16 0.945 — statistically identical.

Throughput (random, ISL 8192 / OSL 1024, tok/s/GPU):

Strategy	Conc	bf16 KV	fp8 KV
Low-latency (MTP-5)	1 / 16	25 / 105	29 / 139
Balanced	64 / 256	122 / 120	148 / 151
High-throughput	1024	140	204

Notes:

• Measured on MI350X (~1000 W); a true MI355X (~1400 W) should be somewhat higher.
• Balanced numbers were re-measured with the balanced cell's own config (cgbs128 / mrr80), not the high-throughput config.
• gfx942 (MI300X/MI325X) and the MXFP4 build keep MTP disabled: gfx942 unvalidated; MXFP4's bf16 MTP layer doesn't load into the MXFP4 draft model.

The --disable-overlap-schedule flag can be dropped once #29785 is fixed. Ready for review once the run-ci label is added.