Integration for GLM-4.6-FP8 and performance tune on H20
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 1.2k
- Forks
- 142
- Avg merge
- 23h 15m
- Merged PRs (30d)
- 8
Description
Description
This issue documents the integration of HPC-Ops kernels with SGLang for GLM-4.6-FP8 serving on NVIDIA H20. The work spans two branches:
- SGLang side: hpc-glm46-integration
- HPC-Ops side: glm46-hpg12-mtp3
The table below summarizes the major gaps identified in HPC-Ops during this integration and the fixes applied. Several of these optimizations are not specific to GLM-4.6 and may be applicable to other model serving configurations.
Configuration note: All results were measured under
--tp 8 --speculative-eagle-topk 1 --speculative-num-steps 3 --speculative-num-draft-tokens 4.
| # | Gap in HPC-Ops | Fix | Result |
|---|---|---|---|
| 1 | Decode kernels support only heads-per-group ∈ {4, 8} (GMMA atom M-constraint); GLM-4.6 requires hpg=12 | Native fused-row tiles (kTileM = seq_q × 12) with 2 consumer warpgroups |
seq_q=4: 217.8 → 153.1 µs vs. the padding approach |
| 2 | No register budgeting for multi-consumer configs: 288-thread block caps all threads at 168 regs, but the consumer path needs ~236, causing 272 register spills | 384-thread padded block + setmaxnreg role rebalance |
153.1 → 131.1 µs |
| 3 | Launcher is wave-quantization-blind: full-regfile CTA → 1 CTA/SM, but grid stayed at 156 CTAs = 2 sequential waves | Single-wave grid sizing (78 CTAs at seq_q=4; per-shape maps) | main: 123.9 → 117.7 µs; seq_q=1: 67.2 → 61.5 µs; seq_q=2: 86.1 → 77.7 µs |
| 4 | Per-tile chain waste: separate 24-FMUL P-rescale pass; dead split_flag zero-fill launch on the dynamic path |
Fold P×256 into softmax exp2; skip dead fill | Combine tail: ~9 → ~6 µs |
| 5 | Combine kernel is latency-bound: serial chunk scan; two-pass max/sum with an extra DRAM round-trip + sync | Warp-parallel chunk splits, register prefetch, single guarded butterfly pass | seq_q=1: 72.4 → 68.4 µs; tail ≈ 6 µs |
| 6 | MoE is per-tensor FP8 only — numerically broken on GLM-4.6 (outlier tokens saturate; tok_rel = 0.986); no blockwise path, no support for non-128-divisible intermediate dims (TP8 → 192/rank) | Blockwise FP8 MoE + zero-pad 192→256 + load-time weight requant | tok_rel 0.986 → 0.061; long-generation looping eliminated; routed MoE reduced from 5 → 1 kernel in-graph |
| 7 | Router (gate → TopK → dispatch) runs as a chain of small generic kernels; last-layer/draft all-reduce is exposed | Fused small-M router chain; all-reduce fused into the following RMSNorm | c=8 e2e: 23.95 → 23.43 ms/iter (cumulative with #3–#5) |
End-to-end (SGLang, accept-len-2.00-filtered decode iterations, 40K context): decode throughput is 1.33–1.38× vs. FA3 + Triton-MoE at concurrency 4–16.
Environment
- GPU: NVIDIA H20 (SM90), 78 SMs, ~4 TB/s HBM, 64K registers/SM
- Model: GLM-4.6-FP8, 92 layers; TP=8 → 12 Q-heads + 1 KV-head per GPU (hpg=12)
- Serving: SGLang with EAGLE (3 steps / 4 draft tokens), page-size 64, fp8_e4m3 KV cache, full CUDA-graph decode; production attention shape is seq_q=4 target verify
- Kernel micro-bench shape: B=16, KV=40,000 (163.8 MB KV/call; 43.1 µs bandwidth-ideal)
I hope some of this work proves useful for other models and configurations. Feedback and suggestions are welcome.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Review the hpc-glm46-integration and glm46-hpg12-mtp3 branches first, focusing on the SGLang/HPC-Ops integration points described in the issue. Validate the GLM-4.6-FP8 configuration on NVIDIA H20 with the listed TP, speculative-decoding, and context settings; done means the integration works and reproduces the reported throughput and latency improvements.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- ai-infra-agents, performance
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Quiet
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100