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

flyer666

周中被老特翻了牌子哈哈哈，所以周六早早爬起来水一篇续写前帖：

给梁圣冲了30块，狠狠地搞了一下dflash2的优化，现在基本差不多满意了，pp变化不大，但是tg比mtp提速了25%左右，还是很爽的。

运行方式和之前一样, pull一下之前的代码：https://github.com/StevenChenSE/sglang/blob/gfx1100-support/README.zh-CN.md
找你最爱的LLM agent harness 重新build一下SGLANG 引擎和kernel，下载对应的dflash2 drafter:

推荐这个：https://huggingface.co/syvai/Qwen3.8-27B-DFlash2-W4A16 节省~2GB显存，实测速度和精度没有损失
原版也可以：https://huggingface.co/incoai/Qwen3.8-27B-DFlash2

Disclailmer: 标题所说速度是仅仅在极短上下文，特定workload下有效，中等长度prompt (16k-32k)的decode速度大概在90-150tok/s之间。
先秀一下：

做数学题：（200-220）

写代码：（180-250tok/s)

普通问答(英文）：(90-120 token/s)

普通问答(中文）：(70-90 token/s)

最后贴一下benchmark测试：

实测性能基准对比

全部基准均在双卡 RX 7900 XTX (TP=2) 实机测试，运行 Qwen3.8-27B-W4A16 模型。
SGLang DFlash2 列于 2026-09-11 在当前构建上刷新（65b16b3df7，Prefill 第 1–5 轮
迭代 + 热 L2 Autotune 修复；llama-benchy 0.4.0，--no-cache 全新提示词，
每个深度取 3 次运行均值）；数学与 120k 回放为单次运行。MTP-3 列与 c=4 取自
2026-09-09 套件（较早构建）。
vLLM 基线列为早期实测数据，如需精确对比请使用相同工具版本重新压测。
上方 DFlash2 列运行 W4A16 草稿模型（自 2026-09-08 起经 zz-w4a16-draft.conf
systemd drop-in 成为生产默认；KV 容量增至 214,284 tokens）。同日同构建的 bf16
草稿模型对照实测为：数学 181.0 tok/s / 120k 回放均值 119.3 tok/s / 深度均值
114.3 tok/s（池容量 189,172）——全面等于或低于 W4A16 列，故维持 W4A16。
第 5 节保留 2026-09-10 的 W4A16 原始验证记录及其取代的 bf16 列。

1. 标准化上下文深度衰减测试 (`llama-benchy`)

标准 Prompt Prefill ($PP=2048$) 与 Token Generation ($TG=128$), 并发数 = 1

上下文深度	SGLang MTP-3 (本分支)	SGLang DFlash2 (本分支)	vLLM MTP-3 基线	vLLM DFlash2 基线
Depth 0	93.6 tok/s	128.3 tok/s	88.6 tok/s	71.1 tok/s
Depth 4,096	90.1 tok/s	136.7 tok/s	83.3 tok/s	68.4 tok/s
Depth 8,192	85.9 tok/s	112.8 tok/s	93.3 tok/s	71.8 tok/s
Depth 16,384	78.2 tok/s	108.8 tok/s	75.9 tok/s	62.2 tok/s
速度留存率 (16k / 0k)	83.6%	84.8%	85.7%	87.5%

MTP-3 对比 vLLM MTP-3：+5.6 / +8.2 / −7.9 / +3.0%（2026-09-09 构建）。DFlash2 于 2026-09-11 刷新：深度 0 绝对速度较 09-09 读数 +19%（128.3 对 107.9），16k 深度 +3.5%（108.8 对 105.1）；留存率比值读数偏低仅因深度 0 基线涨幅大于深上下文点位（本机单次采样波动约 ±15 tok/s，4k 点位恰逢波峰）。DFlash2 对比 vLLM DFlash2：+80.5 / +99.9 / +57.1 / +74.9%。

2. 真实 120k Agent 多轮会话回放（16 轮离散交互）

采样自真实 120k 长文本 Agent 对话（332 $\to$ 120,443 tokens），启用 Radix 前缀缓存

评估指标	SGLang MTP-3 (本分支)	SGLang DFlash2 (本分支)	vLLM MTP-3	vLLM DFlash2
平均生成速度 (Mean TG)	89.46 tok/s	118.7 tok/s	63.50 tok/s	67.96 tok/s
中位数速度 (Median TG)	88.38 tok/s	112.8 tok/s	61.52 tok/s	67.16 tok/s
抖动率 (CV Jitter %)	11.74%	23.6%	45.34%	36.56%
最差轮次保底速度	72.88 tok/s	56.9 tok/s	16.94 tok/s	32.94 tok/s

MTP-3：较 vLLM MTP-3 平均提速 +40.9%、平稳度 3.9 倍、保底速度 4.3 倍（2026-09-09 构建）。DFlash2 于 2026-09-11 刷新：较 vLLM DFlash2 平均提速 +74.7%（118.7 对 67.96），平稳度 1.5 倍，保底速度较 09-09 读数 40.6 提升 40% 至 56.9。注意事项：最后一轮 120k 的前缀缓存命中降至 32%（09-09 启动约 93%），导致该轮 TTFT 膨胀至约 91 秒（TG 仍保持 56.9 tok/s）——当前显存池布局下 120k 前缀树驻留比例下降。

3. 数学思维链推理测试 (GSM8K & MATH-500)

Greedy 贪婪采样，temperature = 0.0，max_tokens = 1024

数据集用例	生成 Token 数量	首字延迟 (TTFT)	Prefill 速度	生成速度 (TG)	准确率
GSM8K #1	48	0.128s	701.6 tok/s	97.3 tok/s	100% 正确
GSM8K #2	196	0.139s	814.8 tok/s	115.5 tok/s	100% 正确
MATH-500 #1	201	0.137s	605.5 tok/s	120.8 tok/s	100% 正确
MATH-500 #2	212	0.128s	568.2 tok/s	110.9 tok/s	100% 正确
综合平均	—	0.133s	672.5 tok/s	113.9 tok/s	100% 正确

SGLang DFlash2（同一测试套件，2026-09-11 刷新）：

数据集用例	生成 Token 数量	首字延迟 (TTFT)	Prefill 速度	生成速度 (TG)	准确率
GSM8K #1	63	0.134s	1510.5 tok/s	144.2 tok/s	✗（答 96，正确为 72）
GSM8K #2	200	0.144s	1567.5 tok/s	208.9 tok/s	✓
MATH-500 #1	195	0.133s	1475.7 tok/s	224.0 tok/s	✓
MATH-500 #2	182	0.131s	1414.7 tok/s	195.6 tok/s	✓
综合平均	—	0.136s	1492.1 tok/s	193.2 tok/s	75%（3/4）

DFlash2 数学生成速度较 09-09 读数快约 70%（193.2 对 166.0 tok/s；两次套件运行复现
193.15 / 193.62），TTFT 在更长提示词下保持平稳，但 GSM8K #1 在贪婪采样下仍会失误
（答 96，正确 72）——这是其窗口式草稿路径在量化稠密 GEMM 上的已知权衡；MTP-3 四题全对。

4. 多并发吞吐实测 ($c=4$)

使用 llama-benchy 进行多并发请求压测 ($PP=2048, TG=128$, Depth 0)

推理服务引擎	并发数	Prefill 总吞吐 (PP)	生成总吞吐 (TG)	峰值生成吞吐	运行稳定性与说明
SGLang MTP-3 (本分支)	c = 4	1,357.4 tok/s	78.5 tok/s	118.0 tok/s	100% 稳定运行，Decode CUDA 图与 MTP-3 正常工作
SGLang DFlash2 (本分支)	c = 4	1,275.5 tok/s	92.1 tok/s	146.0 tok/s	100% 稳定运行，单次采样
vLLM Baseline (无投机)	c = 4	1,891.1 tok/s	83.1 tok/s	180.0 tok/s	原生稳定，但解码速度较低
vLLM DFlash2	c = 4	1,693.2 tok/s	75.9 tok/s	188.0 tok/s	投机开销导致多并发总 TG 吞吐反而低于 Baseline
llama.cpp (MTP)	c = 4	636.3 tok/s	50.1 tok/s	—	受限于插槽并发队列瓶颈 (`-np 2`)

多并发核心结论：SGLang 两套推测解码在 RDNA3 上 4 并发均保持 100% 稳定——无非法显存访问、无图捕获失败（vLLM 原生 MTP-3 在 batch > 1 时仍会崩溃）。DFlash2 以 92.1 tok/s 总生成吞吐领先（较 vLLM 无投机基线 +10.8%，较 vLLM DFlash2 +21.3%），峰值 146 tok/s；MTP-3 为 78.5 tok/s。vLLM 基线列为旧版 llama-benchy 实测，引用精确跨引擎对比前请用 0.4.0 重新压测基线。

5. W4A16 草稿模型（`syvai/Qwen3.8-27B-DFlash2-W4A16`，融合 KV）

2026-09-10，同一硬件与 llama-benchy 0.4.0。深度曲线与 120k 回放在无外部流量的
隔离端口实测；数学、c=4、深上下文与多模态取自 2026-09-09 套件（顺序 KV 路径——
准确率与该路径无关，融合 KV 仅影响草稿上下文 KV 构建）。

深度曲线（PP=2048/TG=128，c=1；两次运行取平均，16k 为三次采样均值）：

指标	bf16 草稿模型（2026-09-09）	W4A16 草稿模型（融合 KV）
Depth 0	107.9 tok/s	118.9 tok/s
Depth 4,096	96.8 tok/s	98.9 tok/s
Depth 8,192	97.7 tok/s	95.9 tok/s
Depth 16,384	105.1 tok/s	98.0 tok/s
速度留存率 (16k / 0k)	97.4%	82.4%

本机单次采样 TG 波动约 ±15 tok/s（深度 0 的背靠背采样分别读得 101.7 与 136.1），解读单点数据时请保留该误差量级。在同一隔离端口关闭融合 KV 后，16k 均值降至 88.1 tok/s——融合 KV 在深度场景带来约 +11% 收益，并将与 bf16 草稿模型的 16k 留存率差距基本抹平。

120k Agentic 会话回放（两次运行，融合 KV）：平均 88.6 tok/s、中位数
82.8 tok/s、CV 25.3%、最差轮次 52.5 tok/s——2026-09-09 的 bf16 草稿
模型为 83.8 / 79.8 / 23.0% / 40.6；2026-09-11 单次复测中 bf16 / W4A16 的均值
分别为 119.3 / 125.5 tok/s。平均速度持平或略优。

数学思维链（greedy，单次运行）：3/4 正确——与 bf16 草稿模型出现 相同的
GSM8K #1 失误（答 96，正确 72），量化草稿模型未带来额外精度损失。平均 TG
148.9 tok/s（63 token 的 GSM8K #1 短答案拉低均值至 54.0 tok/s；其余三题
为 155–200 tok/s）。

深上下文扩展（单次运行，10k→160k TG）：78.4 / 117.3 / 90.4 / 86.1 / 60.1 /
39.0 tok/s——约 80k 之后 TG 开始衰减，与该窗口式草稿家族的长上下文行为一致。

c=4 多并发： 总 TG 89.3 tok/s、峰值 121 tok/s、PP 1,266 tok/s、
100% 稳定（bf16 草稿模型：92.1 / 146.0 / 1,275.5）。

多模态矩阵： 单图 / 多图与多轮交错工作负载均无循环 / 格式缺陷。

坤坤

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

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

抡锤者