[AMD] Use the triton DSA backend for GLM-5.2 MXFP4 on MI355X (#39106)

docs/cookbook/autoregressive/GLM/GLM-5.2.mdx

@@ -118,7 +118,7 @@ import { Playground } from "/src/snippets/_playground.jsx";
 - **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).
-- **MI355X MXFP4 (gfx950-only).** AMD publishes a Quark-quantized **`amd/GLM-5.2-MXFP4`** build for MI355X. It needs `--trust-remote-code` (Quark's custom quant config) and runs at `tp=4` (the 4-bit MoE weights fit a 4-GPU slice) with `--kv-cache-dtype fp8_e4m3`, the same DSA tilelang backends, `--chunked-prefill-size`, and `--watchdog-timeout` as the FP8/BF16 recipes above. This recipe is carried over from the validated `amd/GLM-5.1-MXFP4` MI355X deployment (same DSA architecture family) and has not yet been benchmarked on GLM-5.2, so the Deploy panel marks it unverified.
+- **MI355X MXFP4 (gfx950-only).** AMD publishes a Quark-quantized **`amd/GLM-5.2-MXFP4`** build for MI355X. It needs `--trust-remote-code` (Quark's custom quant config) and runs at `tp=4` (the 4-bit MoE weights fit a 4-GPU slice) with `--kv-cache-dtype fp8_e4m3`, the DSA triton backends (`--dsa-prefill-backend triton --dsa-decode-backend triton`, also SGLang's ROCm default), and the same `--chunked-prefill-size` and `--watchdog-timeout` as the FP8/BF16 recipes above. At `tp=4` this shape — 16 query heads per rank on an FP8 KV cache — is the one the gfx950 sparse-MLA kernels are tuned for; the `tp=8` recipes above stay on tilelang, which is what their published numbers were measured with. This recipe is carried over from the validated `amd/GLM-5.1-MXFP4` MI355X deployment (same DSA architecture family) and has not yet been benchmarked on GLM-5.2, so the Deploy panel marks it unverified.
 
 <Note>
 **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).

docs/src/snippets/configs/zai-org/glm-5.2.jsx

@@ -1011,6 +1011,9 @@ sgl-eval run aime25 \\
     // mi355x cells. MTP (mtp-314, steps=3) is validated on MI355X gfx950 with
     // this precision — see the mtp-314 cell below. Not yet benchmarked for
     // GLM-5.2 on the base strategies → verified:false.
+    // DSA backend: triton (SGLang's ROCm default). At TP4 this shape hits the
+    // gfx950 FP8 sparse-MLA tuning; the TP8 FP8/BF16 cells above stay on
+    // tilelang, where their published numbers were measured.
     // ====================================================================
     {
       match: { hw: "mi355x", variant: "default", quant: "mxfp4", strategy: "low-latency", nodes: "single" },
@@ -1021,8 +1024,8 @@ sgl-eval run aime25 \\
         "--model-path {{MODEL_NAME}}",
         "--tp 4",
         "--kv-cache-dtype fp8_e4m3",
-        "--dsa-prefill-backend tilelang",
-        "--dsa-decode-backend tilelang",
+        "--dsa-prefill-backend triton",
+        "--dsa-decode-backend triton",
         "--chunked-prefill-size 131072",
         "--mem-fraction-static 0.80",
         "--watchdog-timeout 1200",
@@ -1039,8 +1042,8 @@ sgl-eval run aime25 \\
         "--model-path {{MODEL_NAME}}",
         "--tp 4",
         "--kv-cache-dtype fp8_e4m3",
-        "--dsa-prefill-backend tilelang",
-        "--dsa-decode-backend tilelang",
+        "--dsa-prefill-backend triton",
+        "--dsa-decode-backend triton",
         "--chunked-prefill-size 32768",
         "--mem-fraction-static 0.85",
         "--cuda-graph-max-bs-decode 128",
@@ -1059,8 +1062,8 @@ sgl-eval run aime25 \\
         "--model-path {{MODEL_NAME}}",
         "--tp 4",
         "--kv-cache-dtype fp8_e4m3",
-        "--dsa-prefill-backend tilelang",
-        "--dsa-decode-backend tilelang",
+        "--dsa-prefill-backend triton",
+        "--dsa-decode-backend triton",
         "--mem-fraction-static 0.85",
         "--cuda-graph-max-bs-decode 256",
         "--max-running-requests 256",
@@ -1086,8 +1089,8 @@ sgl-eval run aime25 \\
         "--model-path {{MODEL_NAME}}",
         "--tp 4",
         "--kv-cache-dtype fp8_e4m3",
-        "--dsa-prefill-backend tilelang",
-        "--dsa-decode-backend tilelang",
+        "--dsa-prefill-backend triton",
+        "--dsa-decode-backend triton",
         "--speculative-algorithm EAGLE",
         "--speculative-num-steps 3",
         "--speculative-eagle-topk 1",