GLM-5.2-FP8 on gfx950: re-measure on SGLang 0.5.17, fill the two empty cells - #3
Merged
Merged
Conversation
Re-measures the gfx950 cell on rocm/sgl-dev:v0.5.17-rocm724-mi35x-20260820 (ROCm 7.2.4, aiter d9e5ef7ce). Three of the existing cell statements have expired on that image, two of them being its own declared gaps; the playbook documents what changed and how to check each claim, and the generator turns bench_serving JSON into models.js rows the way gen_cookbook_rows.py does for Kimi-K3. No data change in this commit. The continuity check that makes the rest worth reading: the published recipe carried onto the new image reproduces its own c=16 row to 0.7% (531.7 vs 535.66 output tok/s, TPOT 25.01 vs 25.22), so everything else is a delta measured on one harness rather than two labs.
verify-cookbook.js fails a launch command containing a flag with no glossary entry, and the three new gfx950 recipes use four it does not have yet: --schedule-policy, --num-continuous-decode-steps, --speculative-num-steps and --speculative-eagle-topk. Also rewrites --kv-cache-dtype, which said bf16 is required with the DSA tilelang path. That is true on CUDA and not on ROCm: _check_tilelang_dsa_fp8_kv raises only when not hip, and the measured gfx950 pool roughly doubles on fp8_e4m3.
…P works, fp8 KV legal Replaces the single gfx950 cell with three and closes both gaps it declared. gfx942 is untouched. Rows generated by gen_glm52_mi355x_rows.py; each cell's published launch command is built from the argv.txt of the server that produced its numbers, so a recipe cannot drift from its own table. 1. The two mandatory bpreshuffle patches are retired on ROCm >= 7.2. The cell stated its own exit condition -- 'the CK rewrite ROCm/rocm-libraries#8639 supersedes the disable workaround; not in aiter 7d604afe5'. This image ships aiter d9e5ef7ce and #8639 is in its CK submodule, and sglang 0.5.17 gates rather than disables. A/B on identical argv: GSM8K 0.980 both arms, and the patched arm is slower at every concurrency on both request shapes (-8.8/-8.5/-9.0% at ISL 8192, -7.2/-7.0/-7.0% at 76k), with output-token counts identical to 0.00% in all six pairs. The gotcha is rewritten, not deleted: below ROCm 7.2 the patches are still required. 2. gaps[1] 'MTP not enabled on AMD' -- it is. GlmMoeDsaForCausalLMNextN, draft in the same checkpoint, one flag. Isolated against a control differing by exactly that flag: 2.58-2.76x decode from 8k to 256k for 1.03-1.06x prefill. 3. 'KV-cache dtype must be bfloat16 with DSA tilelang' is a CUDA rule. On ROCm fp8_e4m3 takes the pool 1,654,976 -> 3,194,368 tokens at +0.0pp GSM8K and -0.38% accept length. 4. gaps[0] 'balanced / high-throughput' -- both cells filled. Two things the cells say out loud rather than bury. The balanced recipe ships --mem-fraction-static 0.85, not 0.88: at 0.88 it aborts the process on a 131k prefill with HSA_STATUS_ERROR_OUT_OF_RESOURCES at a KV token usage of 0.04, and 0.85 clears 262k for 5.5% of pool and no measurable speed. And the benchmark rows are --dataset-name random, which this repo already established overstates accept length; degeneracy_probe.py run against GLM-5.2 returns 9 unique tokens of 512 with a top 8-gram repeated 27 times, and accept length is 3.994 on random against 2.969 on ShareGPT and 3.555 on GSM8K. The rows stay for comparability with the cell they replace; the cell says to size from real text. Three honest gaps replace the two closed ones: long-context accuracy under fp8 KV (unverified, with the diff harness to run), AIME25 not re-run on 0.5.17, and DP-attention + DeepEP still untried.
DROP THIS COMMIT IF YOU DISAGREE -- it is the only one here that changes a
number you published, and it is deliberately last so it lifts cleanly. If it
goes, the new gfx950 rows should drop total_tok_s too, or the compare table's
'peak total' column reads a 19x MI300X -> MI355X gap that is a units change.
app.js captions the column 'total tok/s adds input tokens', and
gen_cookbook_rows.py emits it from bench_serving's total_throughput. The
glm-5.2-fp8 rows carry aggregate OUTPUT throughput there instead. Three
independent confirmations, no measurement involved:
1. The gfx942 cell holds the same c=1 point twice, once per convention:
total_tok_s 47.8 in one row and 430.7 in the other. 47.8 x 9 = 430.2.
2. bench_serving defines total_throughput = (in+out)/duration and
output_throughput = out/duration, so the ratio is exactly 1 + isl/osl
independent of duration, and exact for --random-range-ratio 1.0.
3. Re-running the gfx950 recipe on a newer image reproduces its published
c=16 figure as OUTPUT throughput to 0.7% (531.7 vs 535.66), while that
run's actual total_throughput is 4785.6.
So nothing measured is rewritten: the published figure moves to output_tok_s,
where it is correct, and total_tok_s / tok_s_per_gpu are recomputed. The row
that was already input-inclusive is detected and left alone -- which is also
the check that the factor is right, since the recomputed twin lands on 430.2
against its 430.7.
Reviewer's GuideRe-measures and extends the GLM-5.2-FP8 MI355X (gfx950) cookbook support on SGLang 0.5.17/ROCm 7.2.4 by retiring the old bpreshuffle patches, enabling MTP/NEXTN and fp8 KV on ROCm, filling the previously empty balanced/high-throughput cells via an automated row generator, updating playbooks and model metadata, and refining CLI flag docs and performance/accuracy notes. Sequence diagram for NEXTN speculative decoding on ROCmsequenceDiagram
participant Client
participant Server as SGLangServer
participant Draft as NEXTN_Draft
participant Target as GLM52_Target
Client->>Server: Request with --speculative-algorithm NEXTN
Server->>Draft: Draft model forward pass
Draft-->>Server: Propose 4 tokens
Server->>Target: Verify proposed tokens
Target-->>Server: Accepted prefix and correction
Server-->>Client: Stream verified output
State diagram for gfx950 bpreshuffle compatibilitystateDiagram-v2
[*] --> LegacyImage
LegacyImage: Older SGLang/aiter or ROCm below 7.2
LegacyImage --> Patched: Disable bpreshuffle at both sites
Patched --> Correct: GSM8K remains valid
ModernImage
ModernImage: SGLang 0.5.17 / ROCm 7.2.4 / aiter d9e5ef7ce
ModernImage --> FastPath: Keep bpreshuffle enabled
FastPath --> Correct: CK fix is present
ModernImage --> SlowPath: Apply legacy patches
SlowPath --> Correct: Accuracy unchanged, throughput reduced
Flow diagram for MI355X recipe selection and KV capacityflowchart TD
Start["Choose MI355X recipe"] --> Latency["Need low single-stream latency?"]
Latency -->|yes| Low["low-latency\nbfloat16 KV + NEXTN\nKV pool 1,645,440 tokens"]
Latency -->|no| Spec["Need speculative decode?"]
Spec -->|yes| Balanced["balanced\nfp8_e4m3 KV + NEXTN\nKV pool 3,194,368 tokens"]
Spec -->|no| High["high-throughput\nfp8_e4m3 KV\nKV pool 3,717,888 tokens"]
Low --> Capacity["Check peak context"]
Balanced --> Capacity
High --> Capacity
Capacity --> Rule["capacity = floor(KV pool tokens / peak context tokens)"]
Rule --> Safe["Keep runnable concurrency at or below capacity"]
File-Level Changes
Tips and commandsInteracting with Sourcery
Customizing Your ExperienceAccess your dashboard to:
Getting Help
|
![sourcery-ai[bot]](https://avatars.githubusercontent.com/in/48477?s=60&v=4){kind=small}
There was a problem hiding this comment.
Hey - I've found 1 issue
Prompt for AI Agents
Please address the comments from this code review:
## Individual Comments
### Comment 1
<location path="glm52_fp8_mi355x_playbook.md" line_range="147-148" />
<code_context>
+
+```bash
+# gfx950 / MI355X, balanced: MTP speculative decode ON TOP OF an fp8_e4m3 KV
+# cache. fp8 KV is ROCm-only on the tilelang DSA path. mem-fraction is 0.88
+# and not higher on purpose -- see the lazy-kernel gotcha.
+export SGLANG_USE_AITER=1
+export PYTORCH_HIP_ALLOC_CONF=expandable_segments:True
</code_context>
<issue_to_address>
**issue:** The balanced recipe comment says `--mem-fraction-static 0.88` is intentional, but the executable command uses `0.85`; the same playbook documents that the speculative fp8 configuration aborts at 0.88 during a long prefill. A reader copying the documented setting can therefore reproduce the process-aborting configuration instead of the shipped safe recipe.
**Triggers:** When a user follows the balanced recipe's inline comments or copies the stated memory fraction rather than the command argument.
**Suggested fix:** Change the comments to state `0.85`, and explicitly identify `0.88` as the failing tested value.
```suggestion
# cache. fp8 KV is ROCm-only on the tilelang DSA path. mem-fraction is 0.85.
# The tested value 0.88 fails during long prefill -- see the lazy-kernel gotcha.
```
</issue_to_address>Sourcery assessment
Approval pending. 1 finding to address first.
Blocking findings: glm52_fp8_mi355x_playbook.md:148
Help me be more useful! Please click 👍 or 👎 on each comment and I'll use the feedback to improve your reviews.
Comment on lines
+147
to
+148
| # cache. fp8 KV is ROCm-only on the tilelang DSA path. mem-fraction is 0.88 | ||
| # and not higher on purpose -- see the lazy-kernel gotcha. |
There was a problem hiding this comment.
issue: The balanced recipe comment says --mem-fraction-static 0.88 is intentional, but the executable command uses 0.85; the same playbook documents that the speculative fp8 configuration aborts at 0.88 during a long prefill. A reader copying the documented setting can therefore reproduce the process-aborting configuration instead of the shipped safe recipe.
Triggers: When a user follows the balanced recipe's inline comments or copies the stated memory fraction rather than the command argument.
Suggested fix: Change the comments to state 0.85, and explicitly identify 0.88 as the failing tested value.
Suggested change
| # cache. fp8 KV is ROCm-only on the tilelang DSA path. mem-fraction is 0.88 | |
| # and not higher on purpose -- see the lazy-kernel gotcha. | |
| # cache. fp8 KV is ROCm-only on the tilelang DSA path. mem-fraction is 0.85. | |
| # The tested value 0.88 fails during long prefill -- see the lazy-kernel gotcha. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
The gfx950 cell declares two gaps. Both are closed on
rocm/sgl-dev:v0.5.17-rocm724-mi35x-20260820: no fork, no patch, same 8× MI355X shape you published on.Calibration first. Your published recipe, re-run unchanged on the new image, reproduces your own ISL 8192 / c=16 row to 0.7% (535.66 → 531.74 aggregate tok/s; TPOT 25.22 → 25.01 ms). Everything below is a delta measured on that one harness.
1. The two mandatory
bpreshufflepatches are retiredYour gotcha named its own exit condition — "not in aiter
7d604afe5". This image shipsd9e5ef7ce, with the CK fix (ROCm/rocm-libraries#8639) present. Keeping the gate forced off now costs 7–9% across six shapes, with GSM8K identical (0.980 both ways). The gotchais rewritten rather than deleted: below ROCm 7.2 it is still required.
2.
gaps[1]"MTP / speculative decode — not enabled on AMD" — it isNEXTN works. BS=1, ISL 8192 / OSL 1024: your published 69.3 tok/s → 204.0 with fp8 KV + NEXTN (2.94×).
Isolated with a one-flag control (
--speculative-algorithm NEXTNon/off, everything else identical), the prefill cost is 1.03–1.06×, not the 2× a cross-image comparison suggests.Accept length is workload-dependent, and
--dataset-name randomoverstates it — exactly as your Kimi-K3 §5.5b says. Your owndegeneracy_probe.py, run unmodified against GLM-5.2, returns 9 unique tokens out of 512, top 8-gram ×27.random, ISL 8192The published rows use the synthetic figure for comparability with your existing cells; 1.72× is the number to plan against.
3. "KV-cache dtype must be bfloat16 with the DSA tilelang backend" — that is a CUDA rule
_check_tilelang_dsa_fp8_kvraises onlywhen not hip. On ROCm,fp8_e4m3is legal and roughly doubles the pool: 1,645,440 → 3,194,368 tokens at the same--mem-fraction-static.4.
gaps[0]— both cells filledlow-latencybalancedhigh-throughputThree honest new gaps replace the two closed ones — chiefly that fp8 KV is not cleared for long-context accuracy; the ISL ladder measures speed, not correctness.
The gotcha worth the most time
MTP + fp8 KV at
--mem-fraction-static 0.88aborts the server process on a 131k prompt — not the request, the process. Speculative graphs are captured out of the remainder after the pool is sized, which is also where a long chunked prefill draws its working set,so the correction runs downward.
0.85clears 262k for 5.5% of pool and no measurable speed cost, and that is what the shipped recipe uses.One change that is not mine, and is easy to drop
The last commit adjusts 6 published
total_tok_svalues, which currently hold output-only throughput against your own caption andgen_cookbook_rows.py. The proof is inside the cell: the same c=1 point appears as 47.8 and 430.7, and 47.8 × 9 = 430.2. It is last in theseries and invents no data — drop it and the rest applies cleanly.
Checks
verify-cookbook.jspasses at every commit,test_render.jsat the tip; all three cells render with all flags documented. gfx942 rows untouched. New rows are generated bygen_glm52_mi355x_rows.pyfrombench_servingJSON, not transcribed.glm52_fp8_playbook.mdkeeps its gfx950 sections and gains supersede pointers, in the style of Kimi-K3 §5.4.Environment: 8× MI355X (gfx950), TP=8 · ROCm 7.2.4 · SGLang
0.5.17.dev20260820+g47fc97d754· aiterd9e5ef7ce· tilelang 0.1.7.post3Summary by Sourcery
Re-measure GLM-5.2-FP8 on MI355X with current ROCm and SGLang support, replacing obsolete limitations with three verified deployment recipes and updated benchmark metadata.
New Features:
Bug Fixes:
Enhancements:
Documentation:
Tests: