魔改SGLANG支持7900XTX 双卡TP V2 DFLASH2 单流220Token/s

坤坤

好家伙，兄弟你是啥主板来着

terry

我弟的测试非常精髓，这是本站神卡7900XTX的封神之作，最后一块短板倍补齐，还是可以折腾下FP8 KV，社区找找方案，到时候让大家抄作业，还有你不是要测试hicache吗？失败了吗？这玩意和FP8KV必须搞定一个

farmer node

这个是跑的是之前的昨天的分支，尝试修复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/s`
16K 深度	9.71s	1,688 tok/s	8.3%	1.554s	`82.37 tok/s`
30K 深度	13.68s	1,727 tok/s	15.2%	1.965s	`65.14 tok/s`
60K 深度	17.85s	2,648 tok/s	30.4%	2.274s	`56.30 tok/s`
120K 深度	48.66s	1,942 tok/s	60.9%	1.719s	`74.45 tok/s`
160K 极限	174.77s	913 tok/s	81.0%	3.510s	`36.47 tok/s`
180K 极限	37.21s (增量)	4,825 tok/s	91.1%	3.615s	`35.41 tok/s`

kos or

中等长度prompt (16k-32k)的decode速度大概在90-150tok/s之间。

這已經具有足夠的生產力了感謝分享！

flyer666

@坤坤华擎x670e tachi 还可以两根32g ddr5超到6000没啥问题

坤坤

我现在用着z590 用ddr4跑着速度勉强可以

flyer666

@terry 哈哈哈这周在搞一个更好玩的project: 把qwen3.8 flash next 从我的这台机器上拉起来，等搞完再看看。

fp8 kv cache估计是不太行，主要是sglang架构问题和7900xtx的硬件支持导致的，我试了几次，就算拉起来16k的prefill就会掉到200几乎不可用，vllm的page attention没这个问题，可以跑fp8 kv （参考我最早的帖子）

terry

@坤坤 X570就可以了PCIE4*8的口两个，他的是X670，D5平台太贵了，好处是体验好一点。

terry

@flyer666 你还是要多研究下，把FP8KV再往前推下，它比Hicache会更实惠一点，HiCache也很重要。还有你2张卡，如果是4比特模型，怎么占用着么大显存呢？按理说权重18G左右，框架开销和预留加起来5个G，还剩下20多G，也不至于KV缓存如此紧张啊。还得让AI帮你研究下，让DSV4.1来搞。

X一棵树

必须d5？你这个数据内存有加分？d4不行？

懒人烘培

根据楼主的文章，我也测试了SGLANG和VLLM，发现SGLANG只是纯 Prefill 吞吐、尤其是多并发短 Prompt，vLLM 更高
文章的 PP=2048、并发4 测试里：
vLLM baseline：1891 tok/s
而：
SGLang DFlash2：1275 tok/s
最终还是选择SGLang是嘛

flyer666

@X一棵树我只有D5，欢迎贡献d4的数据点。

D5和D4区别只在hicache上，如果不用内存kv cache 基本都一样，推理速度除了卡就是主板的pcie带宽了

flyer666

@懒人烘培我的sglang dflash2 单流 pp可以到1900 tok/s，好像并发pp不太行，看你的场景

Xiaote

这两个数不是同一口径，别直接下结论：

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，才能把「引擎差距」和「投机开销」分开。

farmer node

@flyer666 我发的数据就是 D4 的数据，已经把补充的信息贴上去了

Ben Lee

用本地实测来膜拜致敬 @flyer666 大神orz！

https://lcz.me/topic/1814

抡锤者