diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index 0e9dc263faca..c53cc230ba2e 100644 --- a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx +++ b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx @@ -1,6 +1,6 @@ --- title: GLM-5.2 -description: "Deploy GLM-5.2 with SGLang — Z.ai's DeepSeek-Sparse-Attention (DSA) Mixture-of-Experts model with MTP speculative decoding and 1M context, on H200, B200, B300, and GB300." +description: "Deploy GLM-5.2 with SGLang — Z.ai's DeepSeek-Sparse-Attention (DSA) Mixture-of-Experts model with MTP speculative decoding and 1M context, on H200, B200, B300, GB300, and AMD MI300X/MI325X/MI355X." tag: NEW --- @@ -113,6 +113,14 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **BF16 weights need more GPUs.** The full-precision build (`zai-org/GLM-5.2`, ~1.5 TB) does not fit a single 8×H200 / 8×B200 / 4×GB300 node. It fits single-node on **8×B300** (TP8, ~2.1 TB HBM) — **verified**; on the smaller GPUs it needs a **multi-node** layout (e.g. 2×8×H200 or 2×8×B200 at TP16, 2×4×GB300 at TP8), and those **multi-node BF16 recipes are still proposed/inferred** (`verified: false`). FP8 is the recommended deployment. Use the same DSA / MTP / chunked-prefill guidance as FP8. On B300, BF16 low-latency matches FP8 (the sm103 FP8 path is not yet optimized), but FP8 wins at the balanced/high-throughput points. - **Chunked-prefill size is regime-dependent.** At long input (8K+) the default `--chunked-prefill-size 2048` is too small and leaves the balanced point prefill-bound (queueing dominates TTFT). Raising it to `--chunked-prefill-size 32768` on the balanced recipe gave roughly **+34–78% output throughput and −39–59% TTFT** on 8×H200 and 8×B200 (8K-in / 1K-out) in our testing. It is **neutral for high-throughput** (decode-bound there) — keep the default. `--max-running-requests` tracks KV capacity, not a tuning free-for-all: ~60–90 concurrent 8K+1K FP8 requests fit on a single 8-GPU node, so pin balanced near `--max-running-requests 80` and let high-throughput run wider. +- **AMD GPUs (MI300X / MI325X / MI355X).** FP8 (`zai-org/GLM-5.2-FP8`) runs single-node at `tp=8` on all three. BF16 (`zai-org/GLM-5.2`, ~1.51 TB) only fits single-node on **MI325X** (2 TB HBM) and **MI355X** (2.3 TB); **MI300X** (1.5 TB) cannot hold the BF16 weights plus KV cache on one node, so use FP8 there (or a multi-node BF16 layout once validated). Use the DSA tilelang backend (`--dsa-prefill-backend tilelang --dsa-decode-backend tilelang`) and add `--chunked-prefill-size 131072` plus `--watchdog-timeout 1200` (20 min for weight loading). FP8 uses about half the memory of BF16 (~89 GB/GPU vs ~175 GB/GPU). GLM-5.2 and DeepSeek-V3.2 share the same model structure; for other DSA / HiSparse tips see the [DeepSeek-V3.2 cookbook](../DeepSeek/DeepSeek-V3_2). + + +**gfx950 block-FP8 accuracy: fixed as of the pinned MI355X image (`v0.5.13.post1-rocm720-mi35x-20260618`).** Earlier SGLang ROCm images miscompiled AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM on gfx950 (ROCm 7.2): the error was small per layer but compounded across all 78 layers and silently corrupted output — in-context reasoning broke (GSM8K ≈ 0) while short factual prompts still looked fine. The root cause was a gfx950/ROCm-7.2 miscompile of the CK kernel (a packed illegal-type FMA that relied on an LLVM coercion pass removed in ROCm 7.2; non-deterministic wrong rows near tile boundaries). This is resolved in the pinned image and newer: GLM-5.2-FP8 on MI350X/MI355X (gfx950) was re-validated at TP4 and TP8 — **GSM8K ≈ 0.96 (0% invalid)** and **15/15 needle-in-haystack retrieval to ~118K tokens**. **MI300X / MI325X (gfx942) were never affected.** If you must run an older image, treat gfx950 FP8 output as unverified. Background: [sgl-project/sglang#28685](https://github.com/sgl-project/sglang/issues/28685) (analysis) and the upstream CK fix [ROCm/rocm-libraries#8639](https://github.com/ROCm/rocm-libraries/pull/8639) (scalar FMA + accumulator anchor; restores correctness and determinism at -O3). + + +- **MTP / EAGLE speculative decoding** is disabled for AMD in the Deploy panel. The block-FP8 accuracy bug that previously degraded it is now fixed (see note above), but MTP on gfx950 still depends on the spec-decode draft kernel, which is not yet validated on this hardware (and at `--speculative-num-steps > 3` hits a separate build issue). Until MTP is validated on gfx950, omit the `--speculative-*` flags and serve without MTP. + ## 3. Advanced Usage ### 3.1 Reasoning diff --git a/docs_new/src/snippets/configs/zai-org/glm-5.2-benchmarks.jsx b/docs_new/src/snippets/configs/zai-org/glm-5.2-benchmarks.jsx index e107bc61fb71..0994b99ac46c 100644 --- a/docs_new/src/snippets/configs/zai-org/glm-5.2-benchmarks.jsx +++ b/docs_new/src/snippets/configs/zai-org/glm-5.2-benchmarks.jsx @@ -169,4 +169,35 @@ export const benchmarks = [ ttft_ms: 6370, tpot_ms: 280, tokens_per_sec_per_gpu: 430 }, ], }, + // ---- MI355X + FP8 ---- gfx950, TP8, DSA tilelang, NO MTP (disabled on AMD). + // Measured on lmsysorg/sglang-rocm:v0.5.13.post1-rocm720-mi35x-20260618, flush-cache every run. + // No spec-decoding, so not directly comparable to the NVIDIA low-latency cells (EAGLE MTP). + { + match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "low-latency", nodes: "single" }, + sglang_version: "0.5.13.post1", + speed: [ + { workload: { dataset: "random", isl: 8192, osl: 1024, max_concurrency: 1 }, + ttft_ms: 634, tpot_ms: 13.56, tokens_per_sec_per_gpu: 9 }, + { workload: { dataset: "random", isl: 8192, osl: 1024, max_concurrency: 16 }, + ttft_ms: 5411, tpot_ms: 23.60, tokens_per_sec_per_gpu: 69 }, + ], + }, + { + match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "balanced", nodes: "single" }, + sglang_version: "0.5.13.post1", + speed: [ + { workload: { dataset: "random", isl: 8192, osl: 1024, max_concurrency: 64 }, + ttft_ms: 19526, tpot_ms: 46.50, tokens_per_sec_per_gpu: 122 }, + { workload: { dataset: "random", isl: 8192, osl: 1024, max_concurrency: 256 }, + ttft_ms: 117866, tpot_ms: 56.12, tokens_per_sec_per_gpu: 116 }, + ], + }, + { + match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "high-throughput", nodes: "single" }, + sglang_version: "0.5.13.post1", + speed: [ + { workload: { dataset: "random", isl: 8192, osl: 1024, max_concurrency: 1024 }, + ttft_ms: 432058, tpot_ms: 106.44, tokens_per_sec_per_gpu: 141 }, + ], + }, ]; diff --git a/docs_new/src/snippets/configs/zai-org/glm-5.2.jsx b/docs_new/src/snippets/configs/zai-org/glm-5.2.jsx index 8d8a8ac257df..0c28384172e2 100644 --- a/docs_new/src/snippets/configs/zai-org/glm-5.2.jsx +++ b/docs_new/src/snippets/configs/zai-org/glm-5.2.jsx @@ -6,6 +6,7 @@ export const config = { supportedHardware: [ "h200", "b200", "gb300", "b300", + "mi355x", "mi325x", "mi300x", ], // Single released checkpoint — no size/mode split. @@ -93,6 +94,9 @@ sgl-eval run aime25 \\ b200: "lmsysorg/sglang:latest", gb300: "lmsysorg/sglang:latest", b300: "lmsysorg/sglang:latest", + mi355x: "lmsysorg/sglang-rocm:v0.5.13.post1-rocm720-mi35x-20260618", + mi325x: "lmsysorg/sglang-rocm:v0.5.13.post1-rocm700-mi30x-20260616", + mi300x: "lmsysorg/sglang-rocm:v0.5.13.post1-rocm700-mi30x-20260616", // NVFP4 needs the dev image with modelopt_fp4 support (per-quant override). "b200|nvfp4": "lmsysorg/sglang:dev-glm52-nvfp4", "b300|nvfp4": "lmsysorg/sglang:dev-glm52-nvfp4", @@ -113,8 +117,8 @@ sgl-eval run aime25 \\ knobs: [ { id: "tp", label: "TP", values: [null, 4, 8] }, { id: "cp", label: "CP (DSA prefill)", values: [null, 1, 2, 4, 8], - disable: { hw: ["b200", "gb300", "b300"] }, - disableReason: "DSA prefill Context Parallel is verified on Hopper (H200); the Blackwell sm100 DSA-CP FP8 rope kernel is not yet adapted." }, + disable: { hw: ["b200", "gb300", "b300", "mi355x", "mi325x", "mi300x"] }, + disableReason: "DSA prefill Context Parallel is verified on Hopper (H200); the Blackwell sm100 DSA-CP FP8 rope kernel is not yet adapted, and the ROCm DSA-CP path is not yet validated on AMD (MI300X/MI325X/MI355X)." }, { id: "dpAttn", label: "DP-Attention", values: [null, false, 4, 8], labels: { "auto": "Auto", "false": "Off" } }, @@ -149,10 +153,14 @@ sgl-eval run aime25 \\ { id: "off", label: "Off (greedy)" }, { id: "mtp-516", label: "EAGLE / MTP 5-1-6 (low-latency)", flags: ["--speculative-algorithm EAGLE", "--speculative-num-steps 5", - "--speculative-eagle-topk 1", "--speculative-num-draft-tokens 6"] }, + "--speculative-eagle-topk 1", "--speculative-num-draft-tokens 6"], + disable: { hw: ["mi355x", "mi325x", "mi300x"] }, + disableReason: "MTP/EAGLE speculative decoding is not yet validated on AMD ROCm (MI300X/MI325X/MI355X): the gfx950 spec-decode draft kernel is not yet validated and at --speculative-num-steps > 3 hits a separate build issue; the DSA nextn draft path is CUDA-only." }, { id: "mtp-112", label: "EAGLE / MTP 1-1-2 (balanced)", flags: ["--speculative-algorithm EAGLE", "--speculative-num-steps 1", - "--speculative-eagle-topk 1", "--speculative-num-draft-tokens 2"] }, + "--speculative-eagle-topk 1", "--speculative-num-draft-tokens 2"], + disable: { hw: ["mi355x", "mi325x", "mi300x"] }, + disableReason: "MTP/EAGLE speculative decoding is not yet validated on AMD ROCm (MI300X/MI325X/MI355X): the gfx950 spec-decode draft kernel is not yet validated and at --speculative-num-steps > 3 hits a separate build issue; the DSA nextn draft path is CUDA-only." }, ], }, @@ -800,5 +808,277 @@ sgl-eval run aime25 \\ "--port {{PORT}}", ], }, + // ==================================================================== + // AMD MI300X / MI325X / MI355X (ROCm) — TP8, DSA tilelang backend. + // No MTP: disabled in the Speculative card for AMD (the gfx950 spec-decode + // draft kernel is not yet validated, and num-steps>3 hits a separate build + // issue). Strategies differ only by batch-shaping levers + // (cuda-graph-max-bs / max-running-requests / chunked-prefill): + // low-latency — large chunked-prefill, default bs. + // balanced — chunked-prefill 32768 + bs128, max-running 80. + // high-throughput — bs256, max-running 256. + // ACCURACY: the earlier gfx950 block-FP8 bpreshuffle miscompile (GSM8K ~0) is + // fixed as of the pinned mi355x image (...-20260618); MI355X FP8 was re-validated + // (GSM8K ~0.96, NIAH 15/15 to ~118K) and all three FP8 strategies are benchmarked + // + marked verified:true (see glm-5.2-benchmarks.jsx). All BF16 and all gfx942 + // (MI325X/MI300X) cells stay verified:false (not yet benchmarked, but correct). + // BF16 (~1.51 TB) only fits single-node on MI325X (2 TB) / MI355X (2.3 TB); + // MI300X (1.5 TB) needs multi-node, so its BF16 cells are omitted. + // ==================================================================== + { + match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "low-latency", nodes: "single" }, + verified: true, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "balanced", nodes: "single" }, + verified: true, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 32768", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 128", + "--max-running-requests 80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "high-throughput", nodes: "single" }, + verified: true, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi355x", variant: "default", quant: "bf16", strategy: "low-latency", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi355x", variant: "default", quant: "bf16", strategy: "balanced", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 32768", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 128", + "--max-running-requests 80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi355x", variant: "default", quant: "bf16", strategy: "high-throughput", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi325x", variant: "default", quant: "fp8", strategy: "low-latency", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi325x", variant: "default", quant: "fp8", strategy: "balanced", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 32768", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 128", + "--max-running-requests 80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi325x", variant: "default", quant: "fp8", strategy: "high-throughput", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi325x", variant: "default", quant: "bf16", strategy: "low-latency", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi325x", variant: "default", quant: "bf16", strategy: "balanced", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 32768", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 128", + "--max-running-requests 80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi325x", variant: "default", quant: "bf16", strategy: "high-throughput", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi300x", variant: "default", quant: "fp8", strategy: "low-latency", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi300x", variant: "default", quant: "fp8", strategy: "balanced", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--chunked-prefill-size 32768", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 128", + "--max-running-requests 80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi300x", variant: "default", quant: "fp8", strategy: "high-throughput", nodes: "single" }, + verified: false, + env: [], + flags: [ + "--model-path {{MODEL_NAME}}", + "--tp 8", + "--dsa-prefill-backend tilelang", + "--dsa-decode-backend tilelang", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, ], };