魔改SGLANG支持7900XTX 双卡TP V2 DFLASH2 单流220Token/s
-
这个是跑的是之前的昨天的分支，尝试修复FP8 KV 的问题，但是效果不太好，就不提交了，发数据发一下，看看大家是否有参考。（补充说明： 这个就是D4 ， EPYC 7452 Z2, H12D-8D , 内存 4通道 ，128G ）
SGLang DFlash2 双卡 7900 XTX (RDNA3 / gfx1100) 深度调优与全梯度测速报告
1. 实验背景与核心结论
针对双 AMD Radeon RX 7900 XTX (各 24GB VRAM，总 48GB，TP=2) 在运行
Qwen3.8-27B-W4A16-AutoRound及块扩散投机解码Qwen3.8-27B-DFlash2-W4A16时的推理吞吐、算子瓶颈及长上下文退化现象进行了系统排查、源码修改与基准实测。核心结论速览
- 短文本极致爆发：BF16 模式下，结合 Triton 算子补丁，短文本 (~1K) 稳态 Decode 达到 104.94 tok/s（成功破百）。
- 最佳速度平衡点：120K 上下文 + 并发 2 是当前双卡 7900 XTX 下的绝对物理黄金平衡位（单流 120K 稳态 Decode 72.6 ~ 74.8 tok/s，作者基准均值 83.8~88.6 tok/s）。
- 极限深度物理墙 (160K ~ 180K)：在 160K 与 180K 极限物理深度下，受限于双卡 ~1440 GB/s 显存带宽需在每步验证中扫描 11.5GB 历史 KV 数据，Decode 速度物理收敛于 35.4 ~ 36.5 tok/s（与作者官方记录 160K 下的 39.0 tok/s 完全吻合）。
- FP8 KV 的物理局限：因 RDNA3 (gfx1100) 缺少硬件级 FP8 WMMA 矩阵核心，FP8 模式下 Triton 需在通用寄存器内进行逐点软件反量化（Dequant），导致无论如何优化，深上下文 Decode 均被压死在 20~40 tok/s。待社区后续提供硬件对齐的高效 FP8/BF8 算子前，生产极速模式锁死 BF16 KV。
2. 源码修改与工程修复清单
测试镜像已固化为
sglang-gfx1100:pr34058，启动脚本位于~/sglang-gfx1100/launch-dflash2.sh。2.1 RDNA3 Custom All-Reduce JIT 修复
- 原问题：运行时 JIT 编译
rdna_custom_all_reduce.cu时硬编码了 CDNA 专用的__builtin_amdgcn_global_store_b128，导致 gfx1100 报错并降级至慢速 NCCL PCIe。 - 修复：替换为 RDNA3 专用的 128 位向量化 volatile store，打通双卡低延迟内存通信扩展
sgl_rdna_ar_v18_83f10ec8.so。
2.2 Triton 3D Verify 算子性能调优
rdna_unified_verify.py：将 3D Verify 下硬编码的TILE_SIZE = 16改为32，匹配 7900 XTX Wave32 并发粒度，内层循环次数直接减半。triton_backend.py与rdna_verify_adapter.py：将原本针对 16K 上下文硬编码的segments = 32 / 16扩充为segments = 64，将双卡 192 个计算单元 (WGP) 在 Verify 阶段全部吃满。
2.3 参数虚高修正
- 将启动脚本中虚高的
--context-length 225280正式修正为196608(192K)，与实际物理显存池（197,144 Tokens）精准对齐，消除单流超界导致的静默 OOM 风险。
3. 全梯度实测数据对比表 (BF16 KV 稳态)
测试环境：
Qwen3.8-27B-W4A16+DFlash2-W4A16，TP=2，BF16 KV Cache，Radix Cache 前缀命中。上下文深度 Cold Prefill 耗时 Prefill 吞吐速率 显存缓存驻留率 稳态 Decode 耗时 (128 tok) 稳态有效 Decode 速度 短文本 (~1K) 0.05s - 0.1% 1.220s 104.94 tok/s
8K 深度 4.82s 1,701 tok/s 4.1% 1.502s 85.24 tok/s16K 深度 9.71s 1,688 tok/s 8.3% 1.554s 82.37 tok/s30K 深度 13.68s 1,727 tok/s 15.2% 1.965s 65.14 tok/s60K 深度 17.85s 2,648 tok/s 30.4% 2.274s 56.30 tok/s120K 深度 48.66s 1,942 tok/s 60.9% 1.719s 74.45 tok/s160K 极限 174.77s 913 tok/s 81.0% 3.510s 36.47 tok/s180K 极限 37.21s (增量) 4,825 tok/s 91.1% 3.615s 35.41 tok/s -
我弟的测试非常精髓，这是本站神卡7900XTX的封神之作，最后一块短板倍补齐，还是可以折腾下FP8 KV，社区找找方案，到时候让大家抄作业，还有你不是要测试hicache吗？失败了吗？这玩意和FP8KV必须搞定一个
-
@terry 哈哈哈 这周在搞一个更好玩的project: 把qwen3.8 flash next 从我的这台机器上拉起来，等搞完再看看。
fp8 kv cache估计是不太行，主要是sglang架构问题和7900xtx的硬件支持导致的，我试了几次，就算拉起来16k的prefill就会掉到200几乎不可用，vllm的page attention没这个问题，可以跑fp8 kv （参考我最早的帖子）
-
根据楼主的文章，我也测试了SGLANG和VLLM，发现SGLANG只是纯 Prefill 吞吐、尤其是多并发短 Prompt，vLLM 更高
文章的 PP=2048、并发4 测试里：
vLLM baseline：1891 tok/s
而：
SGLang DFlash2：1275 tok/s
最终还是选择SGLang是嘛 -
根据楼主的文章，我也测试了SGLANG和VLLM，发现SGLANG只是纯 Prefill 吞吐、尤其是多并发短 Prompt，vLLM 更高
文章的 PP=2048、并发4 测试里：
vLLM baseline：1891 tok/s
而：
SGLang DFlash2：1275 tok/s
最终还是选择SGLang是嘛这两个数不是同一口径，别直接下结论：
- vLLM 1891 是 baseline（无投机），SGLang 1275 是 DFlash2 路径。投机解码在 prefill 阶段不省反增——要多跑一个 drafter、还要验证，收益全在 decode。所以多并发短 prompt 的纯 prefill SGLang 输是正常的，不代表引擎本身弱。
- 选型按负载：
- 多并发、短 prompt、TTFT/prefill 敏感，用 vLLM，或者 SGLang 关掉投机；
- 单流或少并发、长上下文、decode(TG) 敏感，用 SGLang + DFlash2（也就是 flyer666 那 220 tok/s 的场景）；
- 自己一个人用，多并发通常不是刚需，长 ctx 下的 TG 才决定体感，SGLang 更划算。
- 想公平比，把 SGLang 的投机关掉再跑一遍同样的 PP=2048/c=4，才能把「引擎差距」和「投机开销」分开。
-
@flyer666 我发的数据就是 D4 的数据，已经把补充的信息贴上去了
-
，系统 取消固定了此主题
-
，F franklee006 引用了 此主题
