From 29a66485bcfa3e6880ffcd96fbd7443fed70915d Mon Sep 17 00:00:00 2001 From: andyluo7 Date: Tue, 16 Jun 2026 17:44:47 -0700 Subject: [PATCH 1/9] docs(cookbook): add AMD MI300X/MI325X/MI355X support for GLM-5.2 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 changed. Adds AMD to the GLM-5.2 cookbook: - Deploy panel (glm-5.2.jsx): mi300x/mi325x/mi355x in supportedHardware, ROCm docker images, and single-node TP8 FP8/BF16 cells using the DSA tilelang backend (--dsa-prefill-backend/--dsa-decode-backend tilelang, --chunked-prefill-size 131072, --watchdog-timeout 1200). DSA-prefill CP is disabled on AMD (verified on Hopper only). - No MTP on AMD: EAGLE spec-decode kernel does not yet build for gfx950, so AMD cells omit the --speculative-* flags (documented in Config Tips). - FP8 on MI355X validated end-to-end (serves at TP8 via DSA tilelang). Co-Authored-By: Claude Opus 4 (1M context) --- .../cookbook/autoregressive/GLM/GLM-5.2.mdx | 4 +- .../src/snippets/configs/zai-org/glm-5.2.jsx | 300 +++++++++++++++++- 2 files changed, 302 insertions(+), 2 deletions(-) diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index 19d6a1577a56..081b3be7702a 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, 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, GB300, and AMD MI300X/MI325X/MI355X." tag: NEW --- @@ -100,6 +100,8 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **BF16 weights need more GPUs (unverified).** 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); 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). The BF16 recipes in the panel are **proposed/inferred, not yet benchmarked** (`verified: false`) — FP8 is the recommended deployment. Use the same DSA / MTP / chunked-prefill guidance as FP8. - **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).** Both FP8 (`zai-org/GLM-5.2-FP8`) and BF16 (`zai-org/GLM-5.2`) run single-node at `tp=8`. 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). **MTP / EAGLE speculative decoding is not currently usable on AMD for GLM-5.2** — the spec-decode kernel does not yet build for gfx950 — so omit the `--speculative-*` flags and serve without MTP. The Deploy panel's AMD cells already reflect this (no MTP). 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). + ## 3. Advanced Usage ### 3.1 Reasoning 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 17325716318f..7ea5551fecb6 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. @@ -90,6 +91,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-20260616", + mi325x: "lmsysorg/sglang-rocm:v0.5.13.post1-rocm700-mi30x-20260616", + mi300x: "lmsysorg/sglang-rocm:v0.5.13.post1-rocm700-mi30x-20260616", }, github: { @@ -106,7 +110,7 @@ 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"] }, + 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." }, { id: "dpAttn", label: "DP-Attention", values: [null, false, 4, 8], @@ -628,5 +632,299 @@ sgl-eval run aime25 \\ "--port {{PORT}}", ], }, + + // ==================================================================== + // AMD MI300X / MI325X / MI355X (ROCm) — TP8, DSA tilelang backend. + // No MTP: EAGLE spec-decode kernel does not yet build for gfx950. + // Mirrors the GLM-5.1 ROCm recipe (model path swapped). + // ==================================================================== + { + 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 131072", + "--mem-fraction-static 0.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", + "--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: "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 131072", + "--mem-fraction-static 0.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", + "--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: "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 131072", + "--mem-fraction-static 0.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", + "--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: "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 131072", + "--mem-fraction-static 0.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", + "--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: "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 131072", + "--mem-fraction-static 0.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", + "--chunked-prefill-size 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi300x", 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: "mi300x", 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 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, + { + match: { hw: "mi300x", 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", + "--chunked-prefill-size 131072", + "--mem-fraction-static 0.80", + "--watchdog-timeout 1200", + "--host {{HOST_IP}}", + "--port {{PORT}}", + ], + }, ], }; From 3530557cc2fe53b4ff50b31ff463580cdc74f2fd Mon Sep 17 00:00:00 2001 From: andyluo7 Date: Thu, 18 Jun 2026 10:52:01 -0700 Subject: [PATCH 2/9] =?UTF-8?q?docs(cookbook):=20address=20GLM-5.2=20AMD?= =?UTF-8?q?=20review=20=E2=80=94=20disable=20MTP,=20differentiate=20strate?= =?UTF-8?q?gies,=20drop=20MI300X=20BF16?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - Disable EAGLE/MTP options in the Speculative card for AMD (gfx950 spec kernel unsupported; DSA nextn draft is CUDA-only). - Add AMD to the CP knob's disable reason. - Differentiate the three AMD strategies via batch levers (cuda-graph-max-bs / max-running-requests / chunked-prefill); only the MI355X FP8 low-latency cell stays verified, the rest are inferred. - Remove MI300X BF16 single-node cells — 1.51 TB weights do not fit 1.5 TB HBM; keep MI325X/MI355X BF16 which do fit. Update MDX accordingly. Co-Authored-By: Claude Opus 4 (1M context) --- .../cookbook/autoregressive/GLM/GLM-5.2.mdx | 2 +- .../src/snippets/configs/zai-org/glm-5.2.jsx | 130 ++++++++---------- 2 files changed, 56 insertions(+), 76 deletions(-) diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index 081b3be7702a..c452878327bd 100644 --- a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx +++ b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx @@ -100,7 +100,7 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **BF16 weights need more GPUs (unverified).** 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); 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). The BF16 recipes in the panel are **proposed/inferred, not yet benchmarked** (`verified: false`) — FP8 is the recommended deployment. Use the same DSA / MTP / chunked-prefill guidance as FP8. - **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).** Both FP8 (`zai-org/GLM-5.2-FP8`) and BF16 (`zai-org/GLM-5.2`) run single-node at `tp=8`. 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). **MTP / EAGLE speculative decoding is not currently usable on AMD for GLM-5.2** — the spec-decode kernel does not yet build for gfx950 — so omit the `--speculative-*` flags and serve without MTP. The Deploy panel's AMD cells already reflect this (no MTP). 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). +- **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). **MTP / EAGLE speculative decoding is not currently usable on AMD for GLM-5.2** — the spec-decode kernel does not yet build for gfx950 — so omit the `--speculative-*` flags and serve without MTP. The Deploy panel disables the Speculative card's MTP options on AMD to reflect this. 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). ## 3. Advanced Usage 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 7ea5551fecb6..0e9125547a46 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 @@ -111,7 +111,7 @@ sgl-eval run aime25 \\ { 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", "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." }, + 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" } }, @@ -146,10 +146,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 supported on AMD ROCm (MI300X/MI325X/MI355X): the gfx950 spec-decode kernel does not build and 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 supported on AMD ROCm (MI300X/MI325X/MI355X): the gfx950 spec-decode kernel does not build and the DSA nextn draft path is CUDA-only." }, ], }, @@ -635,8 +639,17 @@ sgl-eval run aime25 \\ // ==================================================================== // AMD MI300X / MI325X / MI355X (ROCm) — TP8, DSA tilelang backend. - // No MTP: EAGLE spec-decode kernel does not yet build for gfx950. - // Mirrors the GLM-5.1 ROCm recipe (model path swapped). + // No MTP: EAGLE/MTP spec-decode is not yet supported on gfx950 (kernel does + // not build; DSA nextn draft path is CUDA-only), so it is disabled in the + // Speculative card for AMD. Strategies differ only by batch-shaping levers + // (cuda-graph-max-bs / max-running-requests / chunked-prefill): + // low-latency — validated config: large chunked-prefill, default bs. + // balanced — chunked-prefill 32768 + bs128, max-running 80. + // high-throughput — bs256, max-running 256. + // Only the MI355X FP8 low-latency cell is benchmarked; the rest are inferred + // from it (verified:false). 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 until a validated multi-node recipe exists. // ==================================================================== { match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "low-latency", nodes: "single" }, @@ -656,15 +669,17 @@ sgl-eval run aime25 \\ }, { match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "balanced", nodes: "single" }, - verified: true, + 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", + "--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}}", @@ -672,15 +687,16 @@ sgl-eval run aime25 \\ }, { match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "high-throughput", nodes: "single" }, - verified: true, + 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", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", "--watchdog-timeout 1200", "--host {{HOST_IP}}", "--port {{PORT}}", @@ -711,8 +727,10 @@ sgl-eval run aime25 \\ "--tp 8", "--dsa-prefill-backend tilelang", "--dsa-decode-backend tilelang", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--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}}", @@ -727,8 +745,9 @@ sgl-eval run aime25 \\ "--tp 8", "--dsa-prefill-backend tilelang", "--dsa-decode-backend tilelang", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", "--watchdog-timeout 1200", "--host {{HOST_IP}}", "--port {{PORT}}", @@ -759,8 +778,10 @@ sgl-eval run aime25 \\ "--tp 8", "--dsa-prefill-backend tilelang", "--dsa-decode-backend tilelang", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--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}}", @@ -775,8 +796,9 @@ sgl-eval run aime25 \\ "--tp 8", "--dsa-prefill-backend tilelang", "--dsa-decode-backend tilelang", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", "--watchdog-timeout 1200", "--host {{HOST_IP}}", "--port {{PORT}}", @@ -807,8 +829,10 @@ sgl-eval run aime25 \\ "--tp 8", "--dsa-prefill-backend tilelang", "--dsa-decode-backend tilelang", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--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}}", @@ -823,8 +847,9 @@ sgl-eval run aime25 \\ "--tp 8", "--dsa-prefill-backend tilelang", "--dsa-decode-backend tilelang", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", "--watchdog-timeout 1200", "--host {{HOST_IP}}", "--port {{PORT}}", @@ -855,8 +880,10 @@ sgl-eval run aime25 \\ "--tp 8", "--dsa-prefill-backend tilelang", "--dsa-decode-backend tilelang", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--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}}", @@ -871,56 +898,9 @@ sgl-eval run aime25 \\ "--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: "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: "mi300x", 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 131072", - "--mem-fraction-static 0.80", - "--watchdog-timeout 1200", - "--host {{HOST_IP}}", - "--port {{PORT}}", - ], - }, - { - match: { hw: "mi300x", 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", - "--chunked-prefill-size 131072", - "--mem-fraction-static 0.80", + "--mem-fraction-static 0.85", + "--cuda-graph-max-bs 256", + "--max-running-requests 256", "--watchdog-timeout 1200", "--host {{HOST_IP}}", "--port {{PORT}}", From 1e4cb7ef86692d76462eacf10053395671b52652 Mon Sep 17 00:00:00 2001 From: andyluo7 Date: Thu, 18 Jun 2026 16:05:42 -0700 Subject: [PATCH 3/9] docs(cookbook): warn GLM-5.2 gfx950 block-FP8 accuracy bug; mark AMD cells unverified A diffuse numerical bug in AMD aiter's gemm_a8w8_blockscale_bpreshuffle GEMM silently corrupts GLM-5.2 block-FP8 output on MI350X/MI355X (gfx950) in current SGLang ROCm images (GSM8K ~0; in-context reasoning broken), while gfx942 (MI300X/MI325X) is unaffected. Add a Warning callout, correct the MTP note (MTP works on gfx950 once the block-FP8 fix is applied), and set all AMD cells to verified:false until a fixed image ships. The deployment configs themselves are correct. Co-Authored-By: Claude Opus 4 (1M context) --- .../cookbook/autoregressive/GLM/GLM-5.2.mdx | 8 ++++++- .../src/snippets/configs/zai-org/glm-5.2.jsx | 23 +++++++++++-------- 2 files changed, 20 insertions(+), 11 deletions(-) diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index c452878327bd..e6dd6fd05364 100644 --- a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx +++ b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx @@ -100,7 +100,13 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **BF16 weights need more GPUs (unverified).** 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); 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). The BF16 recipes in the panel are **proposed/inferred, not yet benchmarked** (`verified: false`) — FP8 is the recommended deployment. Use the same DSA / MTP / chunked-prefill guidance as FP8. - **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). **MTP / EAGLE speculative decoding is not currently usable on AMD for GLM-5.2** — the spec-decode kernel does not yet build for gfx950 — so omit the `--speculative-*` flags and serve without MTP. The Deploy panel disables the Speculative card's MTP options on AMD to reflect this. 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). +- **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). + + +**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** This is an engine bug, not a config issue; the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking upstream in SGLang (the fix is a one-line change to use the plain `gemm_a8w8_blockscale` kernel, matching vLLM, which restores GSM8K ≈ 0.93 and also makes EAGLE/MTP usable on gfx950). Until then, treat gfx950 GLM-5.2 output as **unverified**, or run on gfx942 / use vLLM. + + +- **MTP / EAGLE speculative decoding** is disabled for AMD in the Deploy panel for now: on current gfx950 images it is degraded by the block-FP8 accuracy bug above, and at `--speculative-num-steps > 3` it additionally hits a separate spec-decode kernel build issue. With the block-FP8 fix applied, EAGLE/MTP at `num-steps ≤ 3` works well on gfx950 (accept length ~2.7–3.8). Until the fix ships, omit the `--speculative-*` flags and serve without MTP. ## 3. Advanced Usage 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 0e9125547a46..c1193947caf9 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 @@ -639,21 +639,24 @@ sgl-eval run aime25 \\ // ==================================================================== // AMD MI300X / MI325X / MI355X (ROCm) — TP8, DSA tilelang backend. - // No MTP: EAGLE/MTP spec-decode is not yet supported on gfx950 (kernel does - // not build; DSA nextn draft path is CUDA-only), so it is disabled in the - // Speculative card for AMD. Strategies differ only by batch-shaping levers - // (cuda-graph-max-bs / max-running-requests / chunked-prefill): - // low-latency — validated config: large chunked-prefill, default bs. + // No MTP: disabled in the Speculative card for AMD (current gfx950 images are + // affected by the block-FP8 accuracy bug below, and num-steps>3 hits a + // separate spec-decode kernel 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. - // Only the MI355X FP8 low-latency cell is benchmarked; the rest are inferred - // from it (verified:false). 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 until a validated multi-node recipe exists. + // ACCURACY: all cells are verified:false. On MI350X/MI355X (gfx950), current + // SGLang ROCm images route block-FP8 to aiter gemm_a8w8_blockscale_bpreshuffle, + // which is numerically wrong on gfx950 (compounds across layers; GSM8K ~0) — + // see the Warning in the page body. gfx942 (MI300X/MI325X) is unaffected but + // not yet benchmarked. The configs are correct and will pass once the engine + // fix ships. 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, + verified: false, env: [], flags: [ "--model-path {{MODEL_NAME}}", From fa56770e8dcb6dd3244a9b4eefe59b24fb5dac53 Mon Sep 17 00:00:00 2001 From: andyluo7 Date: Thu, 18 Jun 2026 16:07:54 -0700 Subject: [PATCH 4/9] docs(cookbook): link upstream issue sgl-project/sglang#28685 in gfx950 warning Co-Authored-By: Claude Opus 4 (1M context) --- docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index e6dd6fd05364..071245f025aa 100644 --- a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx +++ b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx @@ -103,7 +103,7 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **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). -**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** This is an engine bug, not a config issue; the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking upstream in SGLang (the fix is a one-line change to use the plain `gemm_a8w8_blockscale` kernel, matching vLLM, which restores GSM8K ≈ 0.93 and also makes EAGLE/MTP usable on gfx950). Until then, treat gfx950 GLM-5.2 output as **unverified**, or run on gfx942 / use vLLM. +**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** This is an engine bug, not a config issue; the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking upstream in [sgl-project/sglang#28685](https://github.com/sgl-project/sglang/issues/28685) (the fix is a one-line change to use the plain `gemm_a8w8_blockscale` kernel, matching vLLM, which restores GSM8K ≈ 0.93 and also makes EAGLE/MTP usable on gfx950). Until then, treat gfx950 GLM-5.2 output as **unverified**, or run on gfx942 / use vLLM. - **MTP / EAGLE speculative decoding** is disabled for AMD in the Deploy panel for now: on current gfx950 images it is degraded by the block-FP8 accuracy bug above, and at `--speculative-num-steps > 3` it additionally hits a separate spec-decode kernel build issue. With the block-FP8 fix applied, EAGLE/MTP at `num-steps ≤ 3` works well on gfx950 (accept length ~2.7–3.8). Until the fix ships, omit the `--speculative-*` flags and serve without MTP. From c070fe142c3a39a260af0faa0b74d0e1574020f2 Mon Sep 17 00:00:00 2001 From: andyluo7 Date: Fri, 19 Jun 2026 14:08:06 -0700 Subject: [PATCH 5/9] docs(cookbook): cite upstream CK fix ROCm/rocm-libraries#8639 in gfx950 warning Point the GLM-5.2 gfx950 block-FP8 accuracy warning at the validated upstream CK kernel fix (scalar FMA + accumulator anchor) instead of the interim plain-kernel workaround. Co-Authored-By: Claude Opus 4 (1M context) --- docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index e74b2691e3ff..fa40abf2870d 100644 --- a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx +++ b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx @@ -103,7 +103,7 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **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). -**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** This is an engine bug, not a config issue; the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking upstream in [sgl-project/sglang#28685](https://github.com/sgl-project/sglang/issues/28685) (the fix is a one-line change to use the plain `gemm_a8w8_blockscale` kernel, matching vLLM, which restores GSM8K ≈ 0.93 and also makes EAGLE/MTP usable on gfx950). Until then, treat gfx950 GLM-5.2 output as **unverified**, or run on gfx942 / use vLLM. +**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** The root cause is a gfx950/ROCm-7.2 miscompile of the CK `gemm_a8w8_blockscale_bpreshuffle` 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 an engine bug, not a config issue — the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking: [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; validated to restore correctness and determinism at -O3). Until that fix ships in a released ROCm/aiter image, treat gfx950 GLM-5.2 output as **unverified** — run on gfx942 (MI300X/MI325X) or use vLLM. - **MTP / EAGLE speculative decoding** is disabled for AMD in the Deploy panel for now: on current gfx950 images it is degraded by the block-FP8 accuracy bug above, and at `--speculative-num-steps > 3` it additionally hits a separate spec-decode kernel build issue. With the block-FP8 fix applied, EAGLE/MTP at `num-steps ≤ 3` works well on gfx950 (accept length ~2.7–3.8). Until the fix ships, omit the `--speculative-*` flags and serve without MTP. From 075566ee4bdb37e630c0440dbe8d1345025e0445 Mon Sep 17 00:00:00 2001 From: zijiexia <37504505+zijiexia@users.noreply.github.com> Date: Fri, 19 Jun 2026 22:24:06 -0700 Subject: [PATCH 6/9] Fix description and update accuracy warning in GLM-5.2 Corrected the description formatting and updated the warning about known accuracy issues on MI350X / MI355X in the GLM-5.2 documentation. --- docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index fa40abf2870d..8363d6cdff1e 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, GB300, and AMD MI300X/MI325X/MI355X." +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 --- @@ -103,7 +103,7 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **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). -**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** The root cause is a gfx950/ROCm-7.2 miscompile of the CK `gemm_a8w8_blockscale_bpreshuffle` 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 an engine bug, not a config issue — the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking: [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; validated to restore correctness and determinism at -O3). Until that fix ships in a released ROCm/aiter image, treat gfx950 GLM-5.2 output as **unverified** — run on gfx942 (MI300X/MI325X) or use vLLM. +**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** The root cause is a gfx950/ROCm-7.2 miscompile of the CK `gemm_a8w8_blockscale_bpreshuffle` 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 an engine bug, not a config issue — the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking: [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; validated to restore correctness and determinism at -O3). Until that fix ships in a released ROCm/aiter image, treat gfx950 GLM-5.2 output as **unverified** — run on gfx942 (MI300X/MI325X). - **MTP / EAGLE speculative decoding** is disabled for AMD in the Deploy panel for now: on current gfx950 images it is degraded by the block-FP8 accuracy bug above, and at `--speculative-num-steps > 3` it additionally hits a separate spec-decode kernel build issue. With the block-FP8 fix applied, EAGLE/MTP at `num-steps ≤ 3` works well on gfx950 (accept length ~2.7–3.8). Until the fix ships, omit the `--speculative-*` flags and serve without MTP. From 05d1c29aacb7294143edbe042e09eee63bafd513 Mon Sep 17 00:00:00 2001 From: andyluo7 Date: Fri, 26 Jun 2026 20:20:37 -0700 Subject: [PATCH 7/9] docs(cookbook): mark GLM-5.2 gfx950 block-FP8 fixed; pin 0618 image, verify MI355X FP8 Re-validated GLM-5.2-FP8 on MI350X/MI355X (gfx950) on image 0618 at TP4 and TP8: GSM8K ~0.96 (0% invalid) and 15/15 NIAH retrieval to ~118K tokens. The block-FP8 bpreshuffle GEMM miscompile (#28685 / ROCm#8639) no longer reproduces on this image. Pin mi355x image to -20260618, flip MI355X FP8 low-latency to verified, and downgrade the gfx950 warning to a historical "fixed" note. Co-Authored-By: Claude Opus 4 (1M context) --- docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx | 8 ++++---- docs_new/src/snippets/configs/zai-org/glm-5.2.jsx | 4 ++-- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx index 8363d6cdff1e..1c6f940234ed 100644 --- a/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx +++ b/docs_new/cookbook/autoregressive/GLM/GLM-5.2.mdx @@ -102,11 +102,11 @@ import { Playground } from "/src/snippets/_playground.jsx"; - **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). - -**Known accuracy bug on MI350X / MI355X (gfx950) in current SGLang ROCm images.** Block-FP8 linear layers are routed to AMD aiter's `gemm_a8w8_blockscale_bpreshuffle` GEMM, which is numerically incorrect on gfx950 (ROCm 7.2). The error is small per layer but compounds across all 78 layers and silently corrupts output — in-context reasoning breaks (GSM8K ≈ 0) while short factual prompts may still look fine. **MI300X / MI325X (gfx942) are not affected.** The root cause is a gfx950/ROCm-7.2 miscompile of the CK `gemm_a8w8_blockscale_bpreshuffle` 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 an engine bug, not a config issue — the deployment flags above are correct and will produce accurate output once a fixed image is available. Tracking: [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; validated to restore correctness and determinism at -O3). Until that fix ships in a released ROCm/aiter image, treat gfx950 GLM-5.2 output as **unverified** — run on gfx942 (MI300X/MI325X). - + +**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 for now: on current gfx950 images it is degraded by the block-FP8 accuracy bug above, and at `--speculative-num-steps > 3` it additionally hits a separate spec-decode kernel build issue. With the block-FP8 fix applied, EAGLE/MTP at `num-steps ≤ 3` works well on gfx950 (accept length ~2.7–3.8). Until the fix ships, omit the `--speculative-*` flags and serve without MTP. +- **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 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 91066c6a0ef9..a433e7967eff 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 @@ -91,7 +91,7 @@ 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-20260616", + 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", }, @@ -656,7 +656,7 @@ sgl-eval run aime25 \\ // ==================================================================== { match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "low-latency", nodes: "single" }, - verified: false, + verified: true, env: [], flags: [ "--model-path {{MODEL_NAME}}", From 29c87c7f6166e968a9b196cdbbece2784b77d557 Mon Sep 17 00:00:00 2001 From: andyluo7 Date: Sat, 27 Jun 2026 10:23:20 -0700 Subject: [PATCH 8/9] docs(cookbook): add MI355X FP8 benchmarks; verify all 3 GLM-5.2 AMD strategies Benchmarked GLM-5.2-FP8 on MI355X (gfx950, TP8, no MTP) on image 0618 at the standard random ISL 8192 / OSL 1024 shape: low-latency (conc 1/16), balanced (conc 64/256), high-throughput (conc 1024). Add the speed numbers to glm-5.2-benchmarks.jsx and flip the balanced + high-throughput FP8 cells to verified:true (low-latency already verified). Co-Authored-By: Claude Opus 4 (1M context) --- .../configs/zai-org/glm-5.2-benchmarks.jsx | 31 +++++++++++++++++++ .../src/snippets/configs/zai-org/glm-5.2.jsx | 16 +++++----- 2 files changed, 39 insertions(+), 8 deletions(-) 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 0324576a68d2..f44c8828b56d 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 @@ -176,4 +176,35 @@ export const benchmarks = [ { match: { hw: "b300", variant: "default", quant: "nvfp4", strategy: "balanced", nodes: "single" } }, { match: { hw: "gb300", variant: "default", quant: "nvfp4", strategy: "low-latency", nodes: "single" } }, { match: { hw: "gb300", variant: "default", quant: "nvfp4", strategy: "balanced", nodes: "single" } }, + // ---- 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 7835d38bcc82..866ff4519165 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 @@ -720,12 +720,12 @@ sgl-eval run aime25 \\ // 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 low-latency is - // re-validated (GSM8K ~0.96, NIAH 15/15 to ~118K) and marked verified:true. The - // balanced/high-throughput, all BF16, and all gfx942 (MI325X/MI300X) cells are - // 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. + // 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" }, @@ -745,7 +745,7 @@ sgl-eval run aime25 \\ }, { match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "balanced", nodes: "single" }, - verified: false, + verified: true, env: [], flags: [ "--model-path {{MODEL_NAME}}", @@ -763,7 +763,7 @@ sgl-eval run aime25 \\ }, { match: { hw: "mi355x", variant: "default", quant: "fp8", strategy: "high-throughput", nodes: "single" }, - verified: false, + verified: true, env: [], flags: [ "--model-path {{MODEL_NAME}}", From db73c2f0e3e14bfe22e3c13bc0fe975890e542f0 Mon Sep 17 00:00:00 2001 From: Zijie Xia Date: Tue, 30 Jun 2026 13:57:45 -0700 Subject: [PATCH 9/9] docs(cookbook): restore GLM-5.2 AMD supportedHardware + MTP/CP disable lost in merge MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The 2026-06-27 merge with main (b1ec59f, "Conflicts: glm-5.2.jsx") dropped two earlier AMD deltas during conflict resolution, leaving the panel non-functional: - supportedHardware no longer listed mi355x/mi325x/mi300x, so the engine never rendered the AMD hardware tabs (buildHardwareGroups filters the catalog by supportedHardware) — all AMD cells, ROCm dockerImages, and MI355X benchmarks were dead. - The Speculative card's two EAGLE/MTP options and the CP (DSA prefill) knob no longer disabled AMD, regressing the addressed review request and contradicting the surviving MDX prose / config comments. Re-add the three mi* ids to supportedHardware and restore the AMD disable blocks (disableReason wording aligned with the current MDX, which uses the "not yet validated / num-steps>3 build issue" framing). Co-Authored-By: Claude Opus 4.8 --- docs_new/src/snippets/configs/zai-org/glm-5.2.jsx | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) 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 866ff4519165..ddac539bad7f 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. @@ -114,8 +115,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" } }, @@ -150,10 +151,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." }, ], },