X99 不能P2P让我很难受 。 买Pcie switch（260926死缓改判无罪，交换机完美实现P2P 9月27：TP2 Qwen3.8 Q8 60t/s）

[张光璞]

@laobenxiong 说:

@张光璞 说:

T7910 PCIE SWITCH 双卡推理续篇：自编 ReBAR 内核吃下 32G 大 BAR、7900XTX 解锁 2000MHz，附一整天的实测数据

期待对 switch 的继续测试. 我感觉 "内部 gen2" 很可疑啊... lspci -tv 看看 topology 到底长什么样? 链路的 bottlenect 在哪里?

把每一跳的 LnkCtl2（Target Link Speed）和 LnkSta 都 dump 了一遍，结论是**"内部 Gen2"这个说法是我读错了**：

PEX880xx 会把自己一堆内部虚拟端口（Port #225/#233/#247）也枚举出来，那些行 Target=5GT/s、EqualizationComplete-（没做过均衡）—— 它们只是下行真实端口的镜像值，不能当证据。

真凶是交换机下行口 显卡上游口那一跳：两端 LnkCap 16GT/s x16、Target 16GT/s，但 LnkSta = 5GT/s x16 (downgraded) ⇒ 要求 Gen4、训练后自己回落，属物理层，不是配置。

宽度全程 x16 没降，所以原文"x8"也写错了，已更正。
交叉验证：该跳单卡 H2D/D2H = 6.04/6.82 GB/s（Gen2 x16 理论上限 8.0，效率 7585%）；直插根口那对是 12.47/13.27（Gen3 x16 上限 15.75，效率 7984%）——效率一致、比值 2.1× ≈ 带宽比 2.0×。

还不止：两张卡共用一个根口（直插那对是各自独占），所以这段上行本身是 Gen3 x16 共享。按实测单卡 6.0 GB/s，两卡满跑 12 < 15.75，当前瓶颈仍是那个 Gen2 跳。
下一步按成本排：换线/转接板重插 → 换交换机下行口 → pcie_aspm=off（现在全链路 ASPM L1 开着）→ 最后才考虑 setpci 强制重训（边际链路可能直接掉卡，且绝不能在跑推理时做）。
卡对卡 P2P 我没下结论：AllReduce 7.04 vs 直连那对 7.19 几乎一样，更像被 RCCL 开销卡住而非 Gen2，要 rocm-bandwidth-test 横幅才定。

══ 0000:05:00.0 ★ R9700（直插槽）
0000:00:02.0 Intel 根口 (Xeon E7v4)
0000:03:00.0 Navi31 上行口 (卡内) ← 上行链路 Gen3 x16 (能力 Gen5 x16) ✓ 正常：上游平台只到 Gen3
0000:04:00.0 Navi31 下行口 (卡内) ← 上行链路 Gen5 x16 (能力 Gen5 x16)
0000:05:00.0 ★ R9700 ← 上行链路 Gen5 x16 (能力 Gen5 x16)

══ 0000:08:00.0 ★ R9700（直插槽）
0000:00:03.0 Intel 根口 (Xeon E7v4)
0000:06:00.0 Navi31 上行口 (卡内) ← 上行链路 Gen3 x16 (能力 Gen5 x16) ✓ 正常
0000:07:00.0 Navi31 下行口 (卡内) ← 上行链路 Gen5 x16 (能力 Gen5 x16)
0000:08:00.0 ★ R9700 ← 上行链路 Gen5 x16 (能力 Gen5 x16)

══ 0000:aa:00.0 ★ 7900 XTX（走交换机）
0000:a0:02.0 Intel 根口 (Xeon E7v4)
0000:a1:00.0 交换机上游口 ← 上行链路 Gen3 x16 (能力 Gen4 x16) ✓ 正常
0000:a2:04.0 交换机端口 [内] ← 上行链路 Gen2 x16 (能力 Gen4 x16) ［镜像值，非物理链路］
0000:a6:00.0 交换机端口 [内] ← 上行链路 Gen2 x16 (能力 Gen4 x16) ［镜像值，非物理链路］
0000:a7:00.0 交换机下端口 ← 上行链路 Gen2 x16 (能力 Gen4 x16) ✗✗ 真实物理链路降速
0000:a8:00.0 卡上行口 ← 上行链路 Gen2 x16 (能力 Gen4 x16) ✗✗ 同一跳的另一端
0000:a9:00.0 卡下行口 ← 上行链路 Gen4 x16 (能力 Gen4 x16)
0000:aa:00.0 ★ 7900 XTX ← 上行链路 Gen4 x16 (能力 Gen4 x16)

══ 0000:b0:00.0 ★ 7900 XTX（走交换机）
0000:a0:02.0 Intel 根口 (Xeon E7v4)
0000:a1:00.0 交换机上游口 ← 上行链路 Gen3 x16 (能力 Gen4 x16) ✓ 正常
0000:a2:08.0 交换机端口 [内] ← 上行链路 Gen2 x16 (能力 Gen4 x16) ［镜像值，非物理链路］
0000:ac:00.0 交换机端口 [内] ← 上行链路 Gen2 x16 (能力 Gen4 x16) ［镜像值，非物理链路］
0000:ad:00.0 交换机下端口 ← 上行链路 Gen2 x16 (能力 Gen4 x16) ✗✗ 真实物理链路降速
0000:ae:00.0 卡上行口 ← 上行链路 Gen2 x16 (能力 Gen4 x16) ✗✗ 同一跳的另一端
0000:af:00.0 卡下行口 ← 上行链路 Gen4 x16 (能力 Gen4 x16)
0000:b0:00.0 ★ 7900 XTX ← 上行链路 Gen4 x16 (能力 Gen4 x16)

[Grayson Ren]

继续加油 持续关注中

[laobenxiong]

@张光璞 如果只能协商到gen2x16, 可以联系一下厂家， 看看能不能调整一下 switch 的 firmware...我觉得这个明显是不科学的。喊他们搞一张 7900xtx 自己复现一下，修一下 firmware...

[张光璞]

我现在还在测试Q6 模型。并且缺少一个机架 。这一套实验台得慢慢收拾 。

[张光璞]

T7910 双卡第三篇：Q8 全精度 + 真张量并行跑通（比层切分快 19.7%），以及一套「三档可切换」的本地推理服务

上篇：《T7910 双卡推理续篇：自编 ReBAR 内核吃下 32G 大 BAR、7900XTX 解锁 2000MHz，附一整天的实测数据》；更早：《双 7900XTX + PCIe 交换机 + SGLang TP=2 实测与掉卡复盘》《T7910 双卡 AI 推理全记录：从 7900XTX 掉总线到 R9700 打通 P2P》。
本篇只写增量：并行方式与精度档的重新选型 —— 从「层切分 + 中间量化」换成「真张量并行 + Q8 全精度」，以及由此衍生出的一套三档服务、它的固化方式和一批带口径的数据。平台、装机、采购价、大 BAR 补丁上篇都讲过，不再重复。

TL;DR

事项	结果	一句话
双 7900XTX 上跑真张量并行（-sm tensor）	成了	上游 llama.cpp 在 ROCm 下必崩；换实验 fork（自研内部 all-reduce + P2P）后跑通
**Q8 全精度（29 GB 权重）**装进双 24 GB 卡	成了	关键不是「显存刚好够」，而是张量并行把显存摊平：22.2 / 22.2 GiB
同二进制 A/B：tensor vs layer	+19.7%	58.16 vs 48.60 tok/s（Q8_0 / 128K / 单流），代码类 +30%
上下文与并发	️ 要读对口径	ctx=196608 是两槽共享的池，不是「每槽 192K」；单请求能吃到满
主机内存提示缓存 48 GiB	️ 超配	淘汰 0 次、命中 83% ⇒ 再加大没用；两个满槽其实 24~33 GiB 就够
客户端出现空回答	踩了	直连推理端口会因「思考通道」把正文吐空 ⇒ 客户端一律走网关别名
「开机即用」的固化	️ 只 enable 服务不够	GRUB 默认内核指向另一版 ⇒ 重启后 P2P 能力全丢，必须连内核一起钉住
拿 R9700 换下单张 7900XTX 跑	不划算	单卡 32.05 tok/s / 只能 64K，对比双卡 36.85 tok/s / 192K（详见 §2.6）

给赶时间的人：这一轮真正改命的是换并行方式（层切分 → 张量并行）——它一次解决了「慢」和「装不下」两件事；而看上去最该调的那几个参数（提示缓存大小、并发槽、草稿深度），实测下来不是已经超配就是本来就不是瓶颈。先量再改，第二次。

---

一、想解决的三件事（思路，不是结论）

1. 精度底线 vs 24 GB 显存

我们这边的偏好是「尽量不降智」：模型从 27B 级稠密模型起步，量化档只考虑 Q6/Q8 这一档，不接受把主力模型换成 4bit。问题很直接：

27B 稠密 · Q8_0  →  权重 ~29 GB
27B 稠密 · Q6_K  →  权重 ~24 GB
双 7900XTX 可用显存 = 2 × 24 GB = 48 GB（还要扣 KV / 计算缓冲 / 混合 SSM 的 rs cache）

单卡 24 GB 装不下 Q8，双卡必须切分。看起来方案很多，实际能用的很少。

2. 为什么不能继续用层切分（-sm layer）

上篇之前我们一直跑的是 SGLang TP=2（int4 权重）和 llama.cpp 的 -sm layer（Q6）。层切分的问题是显存不均：

切分方式	卡 0 显存	卡 1 显存	后果
层切分（-sm layer）	19.45 GiB	22.58 GiB（94%）	满的那张先撞墙 ⇒ 它就是瓶颈卡，上下文再加一点就 OOM
张量并行（-sm tensor）	22.2 GiB	22.2 GiB	完全均分 ⇒ 没有短板卡

这就是「Q8 装不下」的真正原因：不是总量不够，是层切分把 29 GB 的权重压到了某一张卡上。也顺带解释了以前那种「显存没顶满啊，怎么还 OOM」的困惑 —— 另一张卡满了。

3. 三条路与筛选判据链

目标：双 24G 卡上，Q8 精度 + 尽量快
  ├─ 路 A：降精度（int4/4bit）           → ✗ 违背「不降智」底线，先排除
  ├─ 路 B：换卡（单张 32G 卡跑单卡）      → ✗ 带宽/容量实测不划算（§2.6 有数据）
  └─ 路 C：修「张量并行」                 → ✓ 唯一同时解决「慢」和「装不下」的方向
        ├─ C1：上游 llama.cpp `-sm tensor` → ✗ 实测直接崩（ROCm/meta backend）
        └─ C2：找已经修通它的实验 fork     → ✓ 本次落点

判据链的关键一步是把「崩」和「慢」分开：上游那份在 -sm tensor 下是启动阶段就崩（不是慢、不是掉精度），所以它不是调参问题，而是必须换实现 —— 这一步想清楚，后面就不会在参数上瞎耗。

---

二、速度与显存总账

口径声明（论坛上最容易被挑的就是这个，先写在前面）：本节所有数字，除特别标注外均为
Q8_0 权重 · f16 KV · ctx 131072 · --parallel 1（单流）· 开投机解码（MTP，草稿深度 3）· 取四类内容各跑一遍的平均。

2.1 同二进制 A/B（最干净的一组对照）

同一个二进制、同一份 Q8_0 权重、同一上下文，只换切分方式与草稿策略：

配置	平均 decode (tok/s)	散文	代码	技术说明	改写复用
-sm tensor + MTP n-max 3	58.16	49.35	67.61	46.68	69.00
-sm layer + MTP n-max 3（同二进制）	48.60	42.98	52.04	39.96	59.43
-sm tensor + adaptive MTP n-max 4	59.33	49.61	67.69	47.30	72.73

• tensor 比 layer 快 19.7%，且四类内容全部取胜（最夸张的是代码类 +30%）。
• 换用自适应草稿策略（n-max 4）再涨一点到 59.33 tok/s。
• 这一组对照的价值在于排除了「模型/量化差异」这个变量 —— 以前跨框架比出来的数字，往往两种东西一起变了。

2.2 与现役 Q6 档的口径对照（注意口径不同）

档位	权重	切分	ctx	并发	平均 decode
旧：Q6 档	Q6_K 系（24 GB）+ 压缩词表草稿	layer	196608	--parallel 2	36.85 tok/s
新：Q8 档	Q8_0（29 GB）+ 完整词表草稿	tensor	196608	--parallel 2	58~59 tok/s

⇒ 精度更高（Q8 对 Q6）、速度还快 57%。这条对照里两档的并发口径一致，但切分方式和量化档同时变了，所以只做参考，别当纯 A/B 用（纯 A/B 见 §2.1）。

2.3 显存账（同一台机器、同一批卡）

Q8 档（tensor）  ：卡 0 22.2 GiB  |  卡 1 22.2 GiB     ⇐ 完全均分，各剩 ~1.8 GiB
Q6 档（layer）   ：卡 0 19.45 GiB |  卡 1 22.58 GiB    ⇐ 不均，卡 1 已 94%
权重占比         ：Q6 档 ≈ 25 GB / 48 GB = 61%（固定项）
KV + 计算缓冲    ：反推 ≈ 85 KiB/token（f16 KV，含混合 SSM 的 rs cache）
⇒ 想上 256K 上下文，需要再多 ≈ 5.6 GB —— 对 layer 档必撞死，对 tensor 档才是「够不够均分」的问题

2.4 上下文与并发（服务端自报，别用猜的）

load_model: initializing, n_slots = 2, n_ctx_slot = 196608, kv_unified = 'true'

项	值	怎么读
槽数（并发上限）	2	同时最多 2 个请求
单槽上下文上限	196608	单个请求可以吃满 192K
kv_unified	true	192K 是两个槽共享的总池：单请求可吃满，双请求之和 ≤ 192K
实时占用示例	槽 0 = 18,613 tok / 槽 1 = 0	/slots 能直接看到每槽占了多少

别把它读成「2 × 192K」——这是本轮最容易被自己误导的一条（我也是先写错、再对着 /props 改回来的）。

2.5 实况读数（真实使用中的累计值，不是跑分）

跑分是短时的，下面这组是一边用一边攒的服务端累计指标（/metrics）：

指标	值	口径
平均 decode	73.96 tok/s	48,848 token ÷ 660.48 s（短请求多、含大量缓存命中，故高于跑分）
平均 prefill	750 tok/s	78,759 新 token ÷ 105.03 s
历史最大单次 prompt	80,179 token	还没碰到 192K 的边
草稿接受率	90.8%	38,055 / 41,895
逐位置接受数	pos0 10,272 → pos1 9,881 → pos2 9,511 → pos3 8,391	第 4 位仍有 ~78% ⇒ 草稿深度 4 是合理值
平均忙碌槽数	1.0008	基本一直是单槽在解码（没有真并发负载）

2.6 顺手做的一次对照：拿 32 GB 的 R9700 换下单张 7900XTX

同为 24/32 GB 级别的卡，很自然会想「单张 32 GB 岂不是更省事」。实测：

方案	decode (tok/s)	可跑上下文	备注
单张 R9700（32 GB，RDNA4）	32.05	只能 64K	散文 22.29 / 代码 42.28 / 技术说明 35.20 / 改写 28.43；131072 直接 OOM（cudaMalloc failed ... 1197 MiB）
双 7900XTX（24 GB × 2，RDNA3）	36.85（Q6）/ 58~59（Q8）	192K	见上

结论：不换。 原因很硬：decode 是显存带宽受限——R9700 约 640 GB/s，7900XTX 约 960 GB/s，差 1.5 倍不是代际能抹平的。这条也顺手修正了我自己早前的一个误判（当时基于 int4 + 另一个框架的数据，我一度以为 R9700 单卡更强，那是错的）。

---

三、几个反直觉的发现

1. 张量并行的收益不在「更快」，在「装得下」

+19.7% 当然好，但真正值钱的是显存摊平（22.2 / 22.2 vs 19.45 / 22.58）。对双卡用户，这直接决定了「更大的量化档能不能跑」：Q8 能装下，靠的就是它；而层切分下，你永远要被那张最满的卡卡住。

2. 「上下文开得不大」是错觉 —— 要看那张最满的卡

以前问「显存怎么没顶满，是上下文开小了吗」，答案是：总占用不低（41/48 GB），但分布不均——一张 19.45、一张 22.58，瓶颈在后者。换了张量并行之后同样的上下文两颗卡一起 92%，才叫真的顶满。

3. 提示缓存：淘汰为 0，就说明加它没用

llama.cpp 的 --cache-ram（默认只有 8 GiB）是主机内存里的提示缓存上限，不是预分配。我们开到 48 GiB，实测：

项	值
上限	48 GiB
实际匿名页占用	≈ 11.6 GiB（进程 RSS 15.17 GiB）
淘汰次数	0
命中率	83%（392,564 缓存命中 / 78,759 新算）

⇒ 48 GiB 对「2 个 192K 槽」是超配（按 KV 折算两满槽约 24~33 GiB 就够）。判据：只要淘汰计数是 0，加大它不会有任何收益。

4. 直连推理端口会吐「空回答」——不是模型坏了

同一个模型、同一句话：

走网关别名 → '可用'      （客户端带了「关闭思考」的模板参数）
裸直连端口 → ''          （思考 token 全进了 reasoning_content，正文为空）

⇒ 客户端一律走网关别名（顺手还能统一带上 max_input_tokens 之类口径）。别拿裸 curl 的空白答复判模型死刑——这个坑我们当时差点误判成「这个模型不行」。

5. 内核/GRUB 也是「持久化」的一部分

把服务 systemctl enable 了，就以为重启后能自动还原 —— 不一定：

跑通时用的内核：自编 rebar 内核（大 BAR 32G + P2P 可用）
GRUB 默认项      ：另一版内核（p2pdma）  ← 重启后 P2P 能力随内核一起换掉
⇒ 服务 enabled 也白搭：起来的是「跑不通那套张量并行」的机器

修法就是把默认项钉成跑通的那一版（GRUB_DEFAULT 的子菜单编号写法、改前备份 /etc/default/grub、update-grub 后核对，并确认没有开 GRUB_SAVEDEFAULT —— 开了会有 saved_entry 悄悄盖掉默认项）。这条不是这次才发现的道理，但这次才真吃了它的亏。

---

四、还没打通的墙 / 需要围观的地方（含推测标注）

事项	现状	推测 / 下一步
张量模式的 prefill（前置处理）代价	公开的同代双卡实测称 tensor 模式前置比 layer 低约 20%；我们本机没测出退化（8.9K prompt 实测 ~1000 tok/s）	推测：那条数据来自另一个模型/量化档，且前置受工程细节影响大。待验证：固定脚本、同二进制、四种 prompt 长度各跑 3 次取中位
单请求吃满 192K 后没有并发余量	已确认（共享池语义）	想真正双路长上下文必须把池加一倍；但每卡只剩 ~1.8 GiB ⇒ 只能靠 KV 降精度或更小权重，目前不做
采样回 CPU	tensor 模式的固有行为（服务端自报 3 条无害告警）	观察是否影响高并发下的采样延迟；单流实测无感
显存温度 ~90°C	风扇曲线对热点有效、对显存无效	推测：需要显存侧的导热改造（背板/风道），软件层已到极限
交换机路径带宽天花板	卡交换机是 Gen4 x16，内部部分跳只跑 Gen2（上篇已记）	与本篇无关，但仍是「跨卡通信上不去」的上游瓶颈

---

五、给后来人的清单

选型之前

• 先问自己「能不能接受 4bit」：能 ⇒ 路 A 便宜省事；不能 ⇒ 直接看双卡 + 张量并行这条路
• 别用「显存总量」判断能不能装：先算
  单卡最满那一张（层切分一定不均）
• 把「启动就崩」和「跑得慢」分开：崩 ⇒ 换实现，慢 ⇒ 调参，两件事别混着耗

部署时

• 客户端接
  网关别名（带关闭思考的模板参数），别裸连推理端口
• 服务用 systemd 管：
  Restart=on-failure + RestartSec（加载慢的模型给 20s 左右，别抖动重启）
• 同时钉住内核：跑通那版必须是 GRUB 默认项

调优时

• 测速一律
  --parallel 1：--parallel 2 的基线会低读约 20%（拿两种口径的数字互比必然得出错误结论）
• 草稿深度
  试到 4 附近就够（公开数据 2→52.8 / 3→57.3 / 4→57.6 / 6→54.3 tok/s，更深有害）
• 投机解码用
  服务端自报的接受率验收（接受数 ÷ 草稿数），并有逐位置计数：第 N 位还值不值得试，一眼就能看出来
• --cache-ram 别按「越大越好」加：淘汰计数 0 就停手

排障时

• 空回答 ⇒ 先怀疑「思考通道吃掉了正文」，不是模型坏了
• 显存「没顶满但 OOM」⇒ 看第二张卡 / 看是不是层切分不均
• 重启后服务「没起来」⇒ 先看内核是否换了版本，再看服务 enable 状态

---

六、当前状态 + 下一步

• 现在跑的是 Q8 全精度 + 真张量并行档：29 GB 权重 + 完整词表 MTP 草稿，单流 58~59 tok/s，192K 共享池，冷启动约 29 秒（日志自报 listening 在 0:29）。
• 一台机器上并存三档、互斥切换（引擎/precision/上下文各不同），客户端不用动：网关别名跟着档位走。
• 开机即用已固化：服务自启 + 崩溃自愈 + 内核默认项钉死。
• 下一步：① 把 prefill 对比补成正式 A/B（§四 第一条）；② 找一个 Q6 档的替代模型做「精度/速度权衡」的第三点；③ 显存散热改造（软件已尽力）。

---

附：给 AI / 爬虫的结构化数据段

路径与用户名已占位（<USER>、<MODELS_DIR>、<FORK_DIR>）；公开型号、模型文件名、端口照写，可复现性靠它们。

A1. 平台与软件栈

平台: Dell Precision Tower 7910 (BIOS A34)
CPU: 2 × Xeon E5-2683 v4 @2.1GHz
RAM: 125 GiB
GPU_primary: 2 × AMD Radeon RX 7900 XTX 24GB (Navi31 / gfx1100 / RDNA3)   # 现役
GPU_other:   2 × AMD Radeon AI PRO R9700 32GB (gfx1201 / RDNA4)          # 在场，空载（见 2.6 为何不换）
互联: 两卡同挂 PLX 交换机下游（Gen4 x16 下行；内部部分跳 Gen2）→ 已实测不影响本方案
内核: 自编 rebar 内核 7.0.14-rebar（大 BAR 32G + P2P 可用）；GRUB 默认项已钉到同一版
ROCm: 7.2.26015（注意：fork README 写的是 7.14，不可照抄）
OS: Ubuntu 24.04 系

A2. 引擎：实验 fork 与它改了什么

repo: nasone32/llama.cpp-RDNA3-7900xtx-opt      # "RDNA3 真张量并行"实验分支
build_target: gfx1100
version_string: 0.3.0-dev (build 0, commit unknown)
关键改动（按对本方案的作用排序）:
  - 自研内部 HIP all-reduce + P2P（不走 RCCL）        # 上游 -sm tensor 在 ROCm 崩的正面解法
      allreduce.cu
  - Q8 线格式量化 + fused residual + P2P issuer
  - adaptive MTP（草稿深度自适应）
      common/speculative-adaptive.h
  - 修 meta backend（多个提交）
  - Q4_K/Q5_K/Q6_K mmvdq 核 + RDNA WMMA flash-attn + Q6_K MMQ prefill
  - BF16 KV 支持
上游对照: 官方 llama.cpp 在本平台 -sm tensor 启动即崩 ⇒ 本方案依赖该 fork

A3. 服务启动参数（原文；路径已占位）

# 现役 Q8 档（张量并行）
<FORK_DIR>/build-gfx1100/bin/llama-server \
  -m <MODELS_DIR>/Qwen3.8-27B-Uncensored-Q8_0.gguf \
  --model-draft <MODELS_DIR>/Qwen3.8-27B-Uncensored-draft-Q8_0.gguf \
  -sm tensor \
  --host 127.0.0.1 --port 8080 \
  --ctx-size 196608 \
  --parallel 2 --kv-unified --cache-idle-slots \
  --cache-ram 49152 \
  --flash-attn on \
  --cache-type-k f16 --cache-type-v f16 \
  --spec-type draft-mtp --spec-draft-n-max 4 --spec-draft-p-min 0 \
  --no-warmup
# HIP_VISIBLE_DEVICES=2,3   （按本机枚举，两块 XTX 为 ROCm 2/3）

关键参数 ⇄ 作用对照：

参数	值	作用 / 为什么
-sm tensor	—	真张量并行：显存摊平 + 卡间只同步必要数据（本篇主线）
--ctx-size	196608	两槽共享的总池（kv-unified 语义，见 2.4）
--parallel	2	并发槽数
--kv-unified	—	启用共享池：单请求可吃满整个 ctx
--cache-idle-slots	—	空闲槽的缓存行为
--cache-ram	49152 (MiB)	主机内存提示缓存上限（默认 8192）；48 GiB 实测超配
--spec-draft-n-max	4	草稿深度；实测逐位置接受率 pos0 95% → pos3 78%，4 是合理点
--spec-draft-p-min	0	关闭按概率提前截断（配合自适应策略时另调）

A4. fork 需要的环境变量（一个都别删）

export GGML_CUDA_ALLREDUCE=internal   # 用 fork 自带的 all-reduce 实现
export AR_WIRE=q8_0                   # 卡间通信量化格式
export AR_Q8_THRESHOLD=1048576        # 超过该字节数走 Q8 线格式
export AR_FUSED_RESIDUAL=1            # 融合残差
export GDN_CHUNKED_BF16=1             # 混合 SSM 分块 BF16
export AR_P2P=1                       # 启用 P2P 路径
export AR_P2P_ISSUER=1                # P2P issuer

A5. 三档服务与端口口径

档位	引擎	权重	上下文	端口	用途
sglang	SGLang TP=2	int4 + f16 KV + EAGLE 投机	196608	8090	早期 agent 主力档（本轮不涉及）
q6	上游 llama.cpp / layer 切分	Q6_K 系 24 GB + 压缩词表草稿	196608	8080	严谨档（备用）
q8（现役）	实验 fork / tensor 并行	Q8_0 29 GB + 完整词表草稿 3.16 GB	196608（2 槽共享）	8080	日常默认
网关	litellm	—	—	4000	所有客户端连这里（档位别名跟着走）

• q6 与 q8 共用 8080、互斥（同一时刻只有一个在跑）；切换是一次命令的事（约 15~20 秒）
• 客户端（两个不同的 agent 前端）都配在 4000，因此换档不用改客户端

A6. 基准原始数据（JSON）

{
  "meta": {
    "weights": "Q8_0",
    "kv": "f16/f16",
    "ctx": 131072,
    "parallel": 1,
    "spec_decode": "draft-mtp",
    "content_classes": ["prose", "code", "tech_explain", "rewrite_reuse"],
    "metric": "mean decode tok/s over 4 classes"
  },
  "runs": [
    {"split": "tensor", "spec": "draft-mtp n-max=3",        "mean": 58.16, "per_class": {"prose": 49.35, "code": 67.61, "tech": 46.68, "rewrite": 69.00}},
    {"split": "layer",  "spec": "draft-mtp n-max=3",        "mean": 48.60, "per_class": {"prose": 42.98, "code": 52.04, "tech": 39.96, "rewrite": 59.43}},
    {"split": "tensor", "spec": "draft-mtp-adaptive n-max=4","mean": 59.33, "per_class": {"prose": 49.61, "code": 67.69, "tech": 47.30, "rewrite": 72.73}}
  ],
  "delta_tensor_vs_layer_pct": 19.7,
  "vram_gib": {
    "q8_tensor": {"card0": 22.18, "card1": 22.19, "card_total": 24.0},
    "q6_layer":  {"card0": 19.45, "card1": 22.58, "card_total": 24.0}
  },
  "live_counters": {
    "decode_mean_tok_s": 73.96,
    "decode_tokens": 48848,
    "decode_seconds": 660.48,
    "prefill_mean_tok_s": 750,
    "prompt_tokens": 78759,
    "prefill_seconds": 105.03,
    "max_single_prompt_tokens": 80179,
    "spec_draft_tokens": 41895,
    "spec_accepted_tokens": 38055,
    "spec_accept_rate": 0.908,
    "spec_accepted_per_pos": [10272, 9881, 9511, 8391],
    "busy_slots_per_decode": 1.0008
  },
  "host_prompt_cache": {
    "cap_mib": 49152,
    "anon_gib_used": 11.61,
    "process_rss_gib": 15.17,
    "evictions": 0,
    "cached_prompt_tokens": 392564,
    "new_prompt_tokens": 78759,
    "hit_ratio": 0.83
  },
  "single_card_alt": {
    "gpu": "R9700 32GB",
    "mean": 32.05,
    "per_class": {"prose": 22.29, "code": 42.28, "tech": 35.20, "rewrite": 28.43},
    "ctx_max": 65536,
    "oom_at_ctx": 131072,
    "verdict": "not_worth_swapping"
  }
}

A7. 日志关键行（原文）

# 槽与上下文（决定「并发 / 上下文」怎么读）
load_model: initializing, n_slots = 2, n_ctx_slot = 196608, kv_unified = 'true'

# tensor 模式的三条无害告警（见到别当故障排查）
llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort
set_sampler: backend sampling not supported with SPLIT_MODE_TENSOR; using CPU
spec common_specu: backend offload failed for seq_id=0; using CPU sampler
spec common_specu: backend offload failed for seq_id=1; using CPU sampler

# 就绪（冷启动约 29 秒）
srv llama_server: model loaded
srv llama_server: listening on http://127.0.0.1:8080

# 编译期的一个坑（无外网时）：拉 llama-ui 资源超时只是 Warning
Timeout was reached   →   no assets available - building without an embedded UI   # 照样编过

A8. 客户端接入（两个前端都走网关别名）

# 统一口径：客户端 → litellm(:4000) → 具体档位
gateway:
  aliases:
    q8: {upstream: "http://127.0.0.1:8080/v1", extra_body: {chat_template_kwargs: {enable_thinking: false}}, max_input_tokens: 196608}
    q6: {upstream: "http://127.0.0.1:8080/v1"}
    sglang: {upstream: "http://127.0.0.1:8090/v1"}
frontend_A: {base_url: "http://<GATEWAY_HOST>:4000/v1", default_model: q8}
frontend_B: {base_url: "http://<GATEWAY_HOST>:4000/v1", default_model: q8, note: "必须走网关；裸连 8080 会得到空正文"}

A9. 复现命令（关键几步）

# 1) 取 fork 源码（服务器无外网时：在有代理的机器上下 tarball 再 SFTP，双向 md5 校验）
#    别覆盖正在用的 llama.cpp 目录
# 2) 编译（用系统 PATH 里的 hipcc，别照抄作者 README 里他本机的 /opt/rocm/<ver>/... 路径）
cmake -B build -DGGML_HIP=ON -DAMDGPU_TARGETS=gfx1100 -DGGML_HIP_GRAPHS=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build -j

# 3) 冒烟（先确认能出正文，再谈速度）
curl -s http://127.0.0.1:8080/v1/chat/completions -H 'Content-Type: application/json' \
  -d '{"model":"local","messages":[{"role":"user","content":"只回复两个字：可用"}],"max_tokens":8}'

# 4) 判据（唯一可信的「服务就绪」）
curl -s http://127.0.0.1:8080/health          # {"status":"ok"}
curl -s http://127.0.0.1:8080/props  | grep -E 'total_slots|n_ctx'
curl -s http://127.0.0.1:8080/slots           # 每槽实时占用
curl -s http://127.0.0.1:8080/metrics | grep -E 'spec_decode|prompt_tokens|busy_slots'

# 5) 固化：服务自启 + 崩溃自愈 + 内核默认项
systemctl enable <unit>
#   unit 内：Restart=on-failure / RestartSec=20
grep -nE '^(menuentry|submenu)' /boot/grub/grub.cfg   # 先数清子菜单编号再改 GRUB_DEFAULT
cp -a /etc/default/grub /etc/default/grub.bak-$(date +%Y%m%d-%H%M%S)   # 改前必备份
update-grub

A10. 判据与禁用项（踩过才知道的）

• 基准测量：一律 --parallel 1；--parallel 2 的口径会低读约 20%，两种数字不能互比
• 投机解码验收：只认服务端自报的 接受数 / 草稿数（+ 逐位置计数）；别用「看起来快」当验收
• 上下文读数：以 /props 的 total_slots / n_ctx_slot / kv_unified 为准；kv_unified=true ⇒ 共享池
• 缓存调优：淘汰计数为 0 ⇒ 停止加大 --cache-ram
• 服务就绪：只认 /health 200；加载慢的模型先预告「显存会先空后涨」
• 客户端：不裸连推理端口（思考通道会吃掉正文）
• 固化：systemctl enable 与「GRUB 默认内核」必须同时到位，否则重启后能力不一致
• 排序环：单元依赖（显卡限频/OD 服务）没排好时，enabled ≠ 会启动（详见上篇）
• 测卡对卡带宽：只用 RCCL；别拿拷贝工具的数字当通信能力

[terry]

@stxpnet 说:
@PhoenixRise2026 怎么不上PCIE 4.0呢？服务器或者AMD的工作站主板，有没有能让两个PCIE 4.0显示跑满X16的？ @xiaote 你也用 gpt 6 astra查一下，对于 咱们PCIE4.0显卡的用户，想组双卡或者多卡，有哪些硬件和软件要求？

他怎么用GPT6？你出钱？

[laobenxiong]

@张光璞 说:

cache-type-k f16 --cache-type-v f16 \

kv dtype 到底是 f16 还是 bf16? 我觉得这两个差别不小。
llama.cpp 上如何验证这个 activation 是：q4/q6/q8 （weight） * f16 （scale） -> f32 -> 参与 gemm 运算 -> 转 bf16 送到下一层？这个应该是 jwc 大神那个 sglang fork 的 gfx1100-support branch 的情况。

[张光璞]

双 7900 XTX 走 PCIe 交换卡：卡间 P2P 为什么只有 3.6 GB/s，怎么干到 7.2 GB/s

一台 2016 年的双路洋垃圾（X99/C612，Dell Precision T7910）+ 一张 PEX88096 四槽交换卡，塞进 2× 7900 XTX + 2× R9700。
这篇只讲一件事：交换卡下面那两张 XTX 之间的 P2P 带宽 —— 一开始只有 3.62 GB/s（只比走主机内存快 1.6 倍），最后做到 7.23 GB/s / 双向 14.20 GB/s。
全文数字都是同一口径实测，表格带前后对比，判据链每一步排除了什么写在表里；没证实的地方明确标「推测 / 待验证」。

---

TL;DR

问题	结论
卡间 P2P 实测（改造前）	3.62 GB/s 单向 / 7.15 双向 —— 只比「经主机内存中转」（2.30）快 1.6 倍，等于交换卡白挂
为什么这么慢	两层压制：① 交换卡下游端口的目标速率被配成 Gen2（配置层，可以改）② 那张卡插的槽只给 x8 lane（物理层，改不了，只能换槽）
到底被谁掐住的	Gen2 x8 理论 4 GB/s，实测 3.62 = 90% ⇒ 不换槽，寄存器怎么调都到不了 7 GB/s
最终怎么解决	换槽（3 号槽 x8 → 2 号槽 x16）+ 把交换卡内部 fabric 的目标速率提到 Gen4（对 P2P 零贡献，但排除了「fabric 是瓶颈」）⇒ 7.23 / 14.20 GB/s（×2）
还是没解决的	槽那一跳只能 Gen2（两端都要求 Gen4、机制也验证过是好的，就是训不上去）⇒ 天花板 = Gen2 x16 理论 8 GB/s，实测 7.23 = 90%
值得折腾吗	卡间带宽在你的关键路径上就值得（真张量并行每层 all-reduce）；层拆分/各跑各的则完全没必要（交换卡只为「插得下」）

---

一、平台与为什么盯卡间带宽

项	配置
主机	Dell Precision T7910（X99/C612，2× Xeon E7 v4，128 GB 四通道 DDR4-2400，双路 NUMA）
计算卡	2× Radeon RX 7900 XTX 24 GB（gfx1100），插在 PCIe 交换卡的两个槽上
对照卡	2× Radeon AI PRO R9700 32 GB（gfx1201），直连 CPU1 的两个根端口
交换卡	Broadcom PEX88096（96 lane / PCIe Gen4，subsystem 1000:c010）：上游 x16 金手指 → 板上 4 个 PCIe 槽 + 一个拨码开关（x16 / x8+x8 / x4×4×4）
系统	Linux 7.0.x，pci=realloc 已开 ⇒ 四张卡全部拿到 32 GiB BAR0（大 BAR）
用途	本地跑 Qwen3.8-27B（Q8），双卡 -sm tensor 真张量并行 —— 每层都要 all-reduce ⇒ 卡间链路就是它的命脉

• 为什么先盯这个：层拆分（各跑各的）不依赖卡间带宽，但真张量并行每层都要 all-reduce，卡间链路直接进 decode 路径。
• 直觉来自对照：直连根口的那对 R9700 实测 10.27 GB/s，而交换卡下面这对 XTX 只有 3.62 —— 差近 3 倍，这不对劲。
• 口径先说明：本贴只测带宽（大块 peer 拷贝），不测延迟。all-reduce 对短消息更吃延迟，但带宽只有 3.6 的链路，得先把带宽修对。

---

二、为什么只有 3.6 GB/s：先量链路，别猜

逐跳读数（改造前，两张卡）

#	跳	设备	速率	宽度
1	根端口	CPU2 根端口	8.0 GT/s（Gen3）	x16
2	uplink	PEX88096 上游口	8.0 GT/s（Gen3）	x16
3	一层下游	PEX88096 下游端口	5.0 GT/s（Gen2）	x16
4	二层下游	PEX88096 下游端口	5.0 GT/s（Gen2）	x16
5	槽口	交换卡槽口	5.0 GT/s（Gen2）	x16
6	显卡上行口	7900 XTX 板载 1478	5.0 GT/s（Gen2）	x16
7	显卡下行口	7900 XTX 板载 1479	16.0 GT/s（Gen4）	x16
8	显卡核心	7900 XTX	16.0 GT/s（Gen4）	x16
—	另一张卡（3 号槽）	同一条链	5.0 GT/s（Gen2）	x8 ️

算术上就锁死了

• 交换卡所有端口的能力都是 16 GT/s × x16（max_link_speed 全 16 GT/s），但实测只协商 Gen2；
• 另一张卡那一路还掉宽成 x8 ⇒ 实际是 Gen2 x8，理论 4 GB/s；
• 实测 3.62 GB/s = 90% ⇒ 就是这条链路掐住的，不是卡、不是驱动、不是软件栈。

同一时间的对照数据（排除「是不是卡本身不行」）

链路	带宽	说明
显卡内部 VRAM 拷贝（XTX / R9700）	396 / 288 GB/s	卡本身没问题
R9700 R9700（直连根口）	10.27 GB/s	换条路径就是 3 倍
7900 XTX 7900 XTX（交换卡内）	3.62 GB/s	本文主角
7900 XTX 经主机内存中转	2.30 GB/s	交换卡只比它快 1.6×

口径：

• 工具：自写 HIP 基准，hipMemcpyPeer 512 MiB × 30 次，取每卡独立事件 + 墙钟双口径，每次拷贝都做数据校验（防止测到没真拷的自由返回）；
• 状态：跑之前停掉 8080 端口的推理服务（显存回落到 28 MB 空闲），否则既有负载污染数据；
• 四张卡同机、双路 NUMA，测的卡固定在同一路。

---

三、思路：先把「配置」和「物理」切开

判据链（每一步排除了什么、剩下什么）

步	动作	判据 / 排除了什么
①	读目标速率寄存器（LnkCtl2，不是 LnkSta）	交换卡一层/二层端口的目标被写成 Gen2 ⇒ 第一嫌疑是配置，不是芯片能力
②	拿空分支（没插任何设备的口）做实验：两端目标一起改 Gen4	链路真上 16 GT/s（lspci 里 (downgraded) 消失）⇒ ①交换卡接受 OS 写入 ②内部 fabric 真有 Gen4 实力 ⇒ 排除「固件无视写入」与「芯片不行」
③	只改一端	只有那一侧读数变化、链路仍旧档 ⇒ PCIe 速率 = 两端取小，必须两端一起改（最容易白忙的一条）
④	在挂卡的支路上逐跳升 Gen4（先停服务）	内部两跳真升到 16 GT/s，四张卡全程在、amdgpu 零报错 ⇒ 改档/重协商不打扰 GPU（这条本身是个重要结论：敢做）
⑤	槽那一跳升不上去 ⇒ 做降级 sweep（目标刻意降到 Gen1 → 再试 Gen3 → Gen4）	降 Gen1 时链路真的跟着降到 2.5 GT/s（机制完全正常）→ 升 Gen3/Gen4 一律回落 5.0 ⇒ 这一段是训练层面只能 Gen2
⑥	顺手量宽度	3 号槽那张卡 x8、1 号槽 x16 ⇒ 换槽

转折点：「要求更低的速率，看链路跟不跟」

之前一直在猜「是不是板子/线材/固件不行」，反复改目标速率都没用。
改成要求 Gen1、看到链路乖乖降到 2.5 GT/s 的那一刻就清楚了：机制完全正常，是这条信道训不到 Gen2 以上。这一刀把「配置没生效」和「物理到头了」切开了 —— 比继续猜快得多。

另一半转折点是宽度：速率查了半天，回头量 current_link_width 才发现其中一张卡只有 x8 —— 而 Gen2 x8 的理论值恰好把 3.6 GB/s 解释得严丝合缝（90%）。方向一直在「速率」上，答案有一半在「宽度」上。

寄存器偏移（这套 PEX88096 端口实测，别照抄网上的）

偏移（CAP_EXP+）	寄存器	关键位
0x0c	LnkCap	bits[3:0] 最高速率、[9:4] 最高宽度、[31:24] Port #
0x10	LnkCtl	bit4 Link Disable、bit5 Link Retrain
0x12	LnkSta	bits[3:0] 当前速率、[9:4] 当前宽度
0x2c	LnkCap2	bits[7:1] 支持速率向量（0x1e = 2.5/5/8/16 GT/s 全支持）
0x30	LnkCtl2	bits[3:0] 目标速率（1=2.5 / 2=5 / 3=8 / 4=16 GT/s）

• 校准方法：LnkCap 的 bit[31:24] 是 Port #，拿它和 lspci 输出的 Port # 对账；LnkCtl 低位对 ASPM …Enabled。网上一大堆写法是 Root Port 的偏移，跟这套 PLX 端口不一致，套了会得到荒谬结论。
• 本机实测：写好目标速率后，链路自己就重协商了（不需要再写 retrain 位）；个别端口要手动触发才动。
• 所有写入都是易失的（setpci 运行时写）⇒ 重启回 EEPROM 默认（Gen2）。想固化得用板厂 EEPROM 工具，或开机早期脚本 —— 本机因为「瓶颈在槽那一跳」，内部 Gen4 对 P2P 零贡献，所以没做开机脚本。

---

四、最终怎么解决的

1. 换槽（3 号槽 → 2 号槽）：x8 → x16 ⇒ 本轮唯一的实际收益来源，P2P 直接翻倍。
   • 附带提醒：板上有拨码开关（x16 / x8+x8 / x4×4×4），它决定每个槽拿多少 lane —— 插卡前先对一眼，别像我一样插了才发现掉一半。
2. 把交换卡内部 fabric 的目标速率提到 Gen4：对 P2P 没有贡献（瓶颈在槽那一跳），但它彻底排除了「fabric 是不是瓶颈」这个疑问，而且实测不影响 GPU（4 卡无损、零报错）。
3. 结论：在「槽那一跳只能 Gen2」的前提下，7.23 GB/s 就是这台机器的天花板（= Gen2 x16 理论 8 GB/s 的 90%）。

---

五、成绩单（同一口径，前后对比）

链路	改造前	改造后	理论上限	效率
7900 XTX 7900 XTX（单向）	3.62 GB/s	7.23 GB/s	8 GB/s（Gen2 x16）	90%
上面这条（并发双向合计）	≈7.15 GB/s	14.20 GB/s	16 GB/s	89%
7900 XTX 经主机内存中转	2.30 GB/s	3.49 GB/s	—	—
R9700 R9700（直连根口，对照）	10.27 GB/s	10.27 GB/s（没动）	15.75 GB/s（Gen3 x16）	65%
被 x8 压住那张卡的 D2H	≈3.6 GB/s	7.26 GB/s	—	—
卡内 VRAM 拷贝（对照）	396 GB/s	396 GB/s	—	—

数字来源：远端落盘日志（p2p_bw7.log 等）逐行核对，不是印象值。唯二非日志直出的是「改造前的双向合计 ≈7.15」与「改造前那张卡的 D2H ≈3.6」——当时没单独存盘，按 x8 链路（理论 4 GB/s、实测 90%）折算，其余全部日志直出。

顺手观察到一个对照：R9700 那对只有 65% 效率，而 XTX 这对 90%。R9700 两个根口都在 CPU1 上、P2P 要绕 root complex，而 XTX 这对是在交换卡内部路由 —— 效率差的来源大概率在这（推测，未细查，留个问题给懂行的补充）。

---

六、附带实测：memcpy 快 ≠ 内核能直接读对端

同一份基准里我还跑了一个「用 GPU kernel 直接读对端显存」的项（kernel-READ），结果很值得记一笔：

卡对	memcpy-P2P	kernel 直接读对端	差多少
7900 XTX 7900 XTX	7.23 GB/s	7.06 GB/s	基本一致
R9700 R9700	10.27 GB/s	1.19 GB/s	差 8.6 倍

• XTX 那对两个数字几乎相同 ⇒ 那条路径的 peer 访问是「真」的；
• R9700 那对：memcpy 有 10.27，但 kernel 直接读对端只有 1.19（甚至比走主机内存中转的 6.78-6.79 还慢）⇒ 强烈暗示那个 10.27 是拷贝引擎（CE）跑出来的，不代表 GPU 核心能高效直接访问对端显存。
• 为什么这事重要：all-reduce 走的是 GPU 侧的 peer 读写/原子，不是 memcpy。所以「只看 memcpy 带宽」有高估风险 —— 我这份表里 R9700 那对就是典型。
• ️ 标注：kernel-READ 是我自己写的一个简单读循环（对端大块顺序读），不是标准 benchmark；R9700 为什么这么低（驱动路径回退？寻址模式？）属于推测，未验证。但这说明多测一项是有必要的。

---

七、为什么槽那一跳只能 Gen2（推测 + 下一步验证）

已经排除的：

嫌疑	为什么排除
「没协商 / 机制没生效」	降 Gen1 的 sweep：链路真的跟着降到 2.5 GT/s 机制正常
「交换卡不接受 OS 写入」	目标速率回读就是写入值，空端口真跑到 Gen4
「卡的速率/宽度不支持」	两端 LnkCap2 都写 2.5-16 GT/s，显卡板载 1479→核心 跑 16 GT/s
「宽度不够导致上不去」	同一槽 x16 也一样只跑 Gen2

剩下的推测（未验证）：

1. 板侧物理裕量：交换卡的槽走线/连接器在这套组合下只稳定到 Gen2（这类扩展卡/背板常见）；
2. 板侧 EEPROM 里端口 PHY/均衡参数被固化低速：OS 改的是「目标速率」，而 PHY 均衡/预加重参数写在卡的默认配置里 —— 矿场风格固件为了稳定性压速率是常规操作；
3. 显卡侧边缘连接器/PHY 限制：能力寄存器说支持，实际训练上不去。

下一步（做完会补更）：把一张 R9700 插进交换卡的槽里 —— 它在 Gen3 根口下能跑到 8.0 GT/s，说明卡侧至少有 Gen3 能力：

实验结果	结论
R9700 在槽里也只 Gen2	⇒ 板侧限制（要动就只剩 EEPROM 工具）
R9700 在槽里能跑 Gen3+	⇒ 卡侧（Sapphire 这些 XTX 的板上行口限制），无解

---

八、给后来人的建议清单

买之前

• 别只看「Gen4 96 lane」这种营销数字，
  问清每个槽的 lane 数和速率上限 —— x16 槽和 x8 槽对带宽是 2 倍差；
• 有拨码开关的板子，
  问清/确认出厂默认在哪个模式（x16 / x8+x8 / x4×4×4）；
• 先明确目标：
  卡间带宽是否在你的关键路径上？
  • 层拆分 / 每卡一个模型：交换卡只是「插得下」的手段，带宽无所谓；
  • 真张量并行 / 频繁 all-reduce：带宽就是命脉，买之前算清这条链能不能跑到 Gen3+（Gen2 x16 = 8 GB/s 就是极限了）；
• 记得「带宽守恒」：交换卡是
  透明路由器，总带宽不增，只改拓扑 —— 别指望它变快，它只负责让你插得下。

装机时

• 卡
  优先插 x16 的槽；插完立刻量一遍（第 10 节有命令）—— 同型号、同一张卡，换个槽就差一倍；
• 记下每个槽的 lane 分配（拍照/记丝印），以后排障不用再拆机；
• 分清两件事：
  大 BAR 看显卡 VBIOS + 主板 BIOS 开关；P2P 看拓扑（两卡要在同一个 switch 的下游，别各挂一个 switch）；
• lspci -vvv 扫一遍每跳有没有 (downgraded) 标记 —— 有就是没跑到端口能力。

调优时

• 量链路
  看 current_link_*，不是 max_link_*（后者只是端口能力，会让你以为跑在 Gen4）；
• 速率是配置层、可以试着改；宽度是物理层、只能换槽/改拨码 —— 别在软件里找宽度的解；
• 改速率
  两端一起改（PCIe 取小值），一次只动一跳，改完立刻复量；
• 先在
  空端口上做机制验证（零风险），别一上来就动挂卡的端口；
• 别碰 uplink（上游口） —— 它牵着整棵子树。

排障时

• 「降级 sweep」是最快的判据：要求更低的速率，看链路跟不跟 —— 跟 = 机制正常（那升不上去就是物理到头）；不跟 = 目标速率根本没被采纳（固件/配置层）；
• 看到
  (downgraded) 就同时看速率和宽度（我这次就是只盯速率，绕了远路）；
• 改档前
  停掉吃显存的服务/任务，改动做成可回滚（易失写，重启即回默认）；
• 只信实测：
  日志沉默 ≠ 链路正常，「lspci 显示 Gen4」也可能是 max 而不是 current。

---

九、当前状态 + 下一步

• 当前：两张 XTX 在 1、2 号槽（x16），卡间 7.23 GB/s / 双向 14.20 GB/s；8080 上的 Q8 张量并行服务正常跑；交换卡内部 fabric 的 Gen4 是易失设置，重启回 Gen2（对性能无影响，故不做持久化）。
• 下一步：R9700 插槽实验（判「板侧 or 卡侧」）→ 若能跑 Gen3+，就得研究板子的 EEPROM 配置工具；若也只 Gen2，这台机器的卡间 P2P 就到此为止（7.23 = 天花板，接受）。

---

十、附：拓扑 / 复现命令 / 原始数据（给 AI / 给脚本）

拓扑（本机实测，BDF 照写便于复现）

根端口 a0:02.0 (Gen3 x16)
└─ a1:00.0            PEX88096 上游口 (uplink, Gen3 x16)
   └─ a2:08.0         一层下游（目标速率可写，实测能升 Gen4）
      ├─ aa:00.0 → ab:00.0(槽1 槽口) → ac:00.0(7900XTX 板载 1478) → ad:00.0(1479) → ae:00.0 核心
      └─ ab:10.0(槽2 槽口) → af:00.0(7900XTX 板载 1478) → b0:00.0(1479) → b1:00.0 核心

R9700 直连根口：00:02.0 → 05:00.0 核心 | 00:03.0 → 08:00.0 核心

• 显卡自带上行/下行口为 1002:1478 / 1002:1479（芯片本身 Gen4，核心侧一直 16 GT/s）；
• 槽口的对端读数是显卡上行口（ab:00.0 ↔ ac:00.0），这两端速度必须一致，不一致说明你看错了链路。

量链路（非 root 可读，最省事的判据）

BDF=<显卡核心BDF>   # 例如 ae:00.0
readlink -f /sys/bus/pci/devices/0000:$BDF | grep -oE '0000:[0-9a-f]{2}:[0-9a-f]{2}\.[0-9]' | while read b; do
  echo "$b  current=$(cat /sys/bus/pci/devices/$b/current_link_speed) x$(cat /sys/bus/pci/devices/$b/current_link_width)  max=$(cat /sys/bus/pci/devices/$b/max_link_speed) x$(cat /sys/bus/pci/devices/$b/max_link_width)"
done

读/写目标速率 + 重协商（root；先在空端口上试）

# 读目标速率（低 4 位：1=2.5 / 2=5 / 3=8 / 4=16 GT/s）
setpci -s <BDF> CAP_EXP+0x30.w
# 写 Gen4，保留其余位
setpci -s <BDF> CAP_EXP+0x30.w=0x0004:0x000f
# 需要时手动触发重训（本机多数端口会自己重协商）
setpci -s <BDF> CAP_EXP+0x10.w=0x0020:0x0020
# 回读验证
setpci -s <BDF> CAP_EXP+0x30.w; lspci -vvv -s <BDF> | grep -E 'LnkSta|LnkCap2'

• ️ 易失、重启回默认；别对 uplink 这么干、别在跑任务时干、一次只改一跳并立刻复量。

原始数据（改造后，p2p_bw7.log 逐行）

{
  "bench": {"tool": "hip_memcpy_peer", "buffer": "512MiB", "iters": 30, "verify": "每轮数据校验", "service": "已停(显存空闲)"},
  "intra_card_copy_GBps": {"R9700_05": 288.5, "R9700_08": 288.9, "XTX_ae": 396.3, "XTX_b1": 397.3},
  "p2p_memcpy_GBps": {"R9700_05<->08": 10.27, "XTX_ae<->b1": 7.23},
  "p2p_bidir_total_GBps": {"R9700_05<->08": 19.73, "XTX_ae<->b1": 14.20},
  "p2p_kernel_read_GBps": {"R9700_05->08": 1.19, "XTX_ae->b1": 7.06},
  "via_host_mem_GBps": {"R9700_05->08": 6.78, "XTX_ae->b1": 3.49},
  "d2h_GBps": {"R9700": 13.99, "XTX_ae": 7.27, "XTX_b1": 7.26},
  "before": {"p2p_memcpy_GBps": 3.62, "p2p_bidir_total_GBps": 7.15, "via_host_mem_GBps": 2.30}
}

环境

• 内核：Linux 7.0.x（pci=realloc，四卡大 BAR 32 GiB 就位）
• 交换卡：Broadcom PEX88096（96 lane / Gen4，1000:c010），板上 4 槽 + 拨码开关（x16 / x8+x8 / x4×4×4）
• 显卡：2× 7900 XTX（gfx1100）+ 2× R9700（gfx1201）
• 推理：8080 端口，llama.cpp 系后端 -sm tensor（真张量并行），Qwen3.8-27B Q8
• 基准编译：hipcc -O2（ROCm）；平台：双路 X99 + 双 NUMA

---

本文所有改动均为运行时 setpci 写入（易失），无任何固件/BIOS 修改；跑基准前均停掉推理服务。

[张光璞]

双 7900XTX 自制 W8A16：从 6 小时量化到 SGLang 生产栈上线

硬件：Dell Precision T7910 · 2× RX 7900 XTX 24G（gfx1100/RDNA3）+ 2× R9700 32G（gfx1201/RDNA4）
时间跨度：2026-09-28 深夜 ~ 2026-09-29 傍晚
关键词：auto-round · W8A16 · GPTQ · exllama · marlin · SGLang · EAGLE/MTP · HiCache · 显存账

TL;DR

• 用官方 BF16 底模自己量了个 W8A16（auto-round GPTQ，group_size=128），双 R9700 全 GPU 跑了 约 6 小时，产物 30G。
• 在双 7900XTX 上把 SGLang 跑成生产服务，路上三块硬石头：marlin 内核 HIP 编不过、exllama 只吃 fp16、EAGLE 草稿头与主模型同源。
• 结果：prefill 全档碾压（比 W4 快 ~80%、比 Q8 快 42~63%）、TTFT 最低、质量无损；代价是 decode 慢 24~49%（权重 30G vs 17G，账对得上）和 EAGLE 把 KV 池从 190K 压到 75.7K。
• 顺手推翻了旧结论：HiCache 并没有让 prefill 变慢——真凶是当时自打的 write-through 补丁；write_back 下 8K 冷 prefill 5.44~5.48s，与基线 5.43s 一致，而重复请求提速 12~76 倍、被逐出文档重命中提速 26 倍。

---

一、缘起：为什么要自己量一个 W8A16

手上已有两个引擎在跑同一个 27B：

档位	引擎	体积	强项	短板
Q8_0	llama.cpp（RDNA3 fork，-sm tensor）	29G	单流深上下文 decode 最强	并发弱、prefill 中等
W4A16	SGLang（TP=2）	17G	decode 快、并发好	prefill 最低

缺的是「长文档首答 / 高并发 prefill」这一档——于是自己做 W8A16 补上：体积接近 Q8（质量≈Q8），但要吃到 SGLang 的 continuous batching + radix 缓存红利。

---

二、量化：双 R9700 全 GPU，约 6 小时

2.1 配方

• 底模：Qwen/Qwen3.8-27B 官方 BF16（52G / 18 分片，含 333 个 visual.* + mtp.*）
• 参数：auto-round 0.15.0、bits=8 group_size=128 sym=true、packing_format=auto_round:auto_gptq、--no-quant_lm_head、block_name_to_quantize=model.language_model.layers + mtp.layers

2.2 关键取舍：谁保 BF16

对象	处理	原因
48 个线性注意力层的 in_proj_a / in_proj_b（共 96 条 bits=16 例外）	保 BF16	线性注意力敏感层，量化掉点
视觉塔 333 个张量（model.visual.*）	保 BF16	图文能力无损，实测识图通过
lm_head	保 BF16（--no-quant_lm_head）	词表头对精度敏感

2.3 产物与耗时

项	值
起止	2026-09-28 22:55:21 → 2026-09-29 04:54:07（约 6 小时）
运行方式	双 R9700 全 GPU（用户前夜已断 XTX 供电防跳闸）
产物	Qwen3.8-27B-base-w8g128，30G / 10 分片 + mtp_bf16_extra.safetensors + model_extra_tensors.safetensors
峰值显存	{0: 25.97GB, 1: 6.56GB}（device_map auto 分得不均，属正常）、peak_ram 14.63GB

2.4 量化阶段的三个坑

1. 必须用 auto-round 入口命令：python -m auto_round.cli.main 会 runpy 双导入，静默不干活（exit 0 但啥都没跑）。
2. RDNA4 上必须 --disable_torch_compile：否则 Inductor 报 CUDA driver error 301。
3. GPTQ 格式的正确参数是 auto_round:auto_gptq：写 --format gptq 无效。

---

三、服务化：三块硬石头

从「加载就崩」到「能跑」，全程卡了三处，每一处都是可复用判据。

3.1 marlin 路径在 HIP 上必崩 → 用 backend: gptq 改路由

• 症状：GPTQ 权重加载即崩在 gptq_marlin_repack 的 JIT 编译。
• 根因：sglang/srt/layers/quantization/gptq_marlin*.cuh 里残留 CUDA API cudaFuncSetAttribute，HIP 编不过（别去改头文件，直接绕开 marlin）。
• 修法：checkpoint config.json 的 quantization_config 加 "backend": "gptq" ⇒ 路由到 exllama GPTQLinearMethod；quant_method 保持 "auto-round"，让那 96 条 bits=16 例外原生走 UnquantizedLinearMethod。
• 判据：A 卡上 AutoRound/GPTQ「能加载但崩在 repack JIT」= marlin 路径，改路由不改内核。

3.2 exllama 内核只吃 fp16 → scales 与激活都要转

对象	处理
scales	SGLANG_RDNA_FP16_GPTQ=1（加载时 bf16 → fp16）
激活	补丁 gptq_kernels.py 的 GPTQLinearKernel.apply：每次调用把 bf16 激活转 fp16，输出转回 bf16 再加 bias

3.3 EAGLE 草稿头与主模型同源

产物自带 mtp.layers（同样量化为 8bit）⇒ 不需要 --speculative-draft-model-path，直接吃同目录草稿头：

--speculative-algorithm EAGLE --speculative-num-steps 3 \
--speculative-eagle-topk 1 --speculative-num-draft-tokens 4

---

四、首轮基准：三引擎全梯度对照

口径：W8 首轮为 无投机解码（speculative_algorithm=None）裸速度，双 XTX TP=2；对照的 W4/Q8 数据来自 09-28 同卡同频基准（同在 SCLK 1500MHz 供电限频下，绝对值偏低约 20%，相对比较公平）。

4.1 Prefill（tok/s）——W8 全档碾压

上下文	W8A16	W4A16	Q8_0	W8 vs W4	W8 vs Q8
8K	1626	904	997	+80%	+63%
16K	1518	857	1012	+77%	+50%
64K	1095	607	770	+80%	+42%
128K	798	438	563	+82%	+42%

4.2 Decode（tok/s）——W8 全档落后，但账对得上

上下文	W8A16	W4A16	差距
8K	35.6	69.3	-49%
16K	34.5	62.9	-45%
64K	29.5	49.5	-40%
128K	24.9	38.9	-36%

W8 权重 30G vs W4 17G，比值 1.76 ⇒ decode 降幅与体积比吻合，内核效率达理论上限 ~89%，无额外损耗。

4.3 TTFT 与并发

指标	W8A16	W4A16	Q8_0
TTFT @128K	165s	299s	176K 时 387s
4 路并发 @16K makespan	77.6s	100.1s	145.2s
4 路并发 @16K 聚合吞吐	26.4 tok/s	20.5	14.1

4.4 质量

类别	W8A16	W4A16	Q8_0
NIAH 长上下文检索（27 条）	27/27	27/27	27/27
推理（10 题）	8/10	7/10	8/10
结构化表格	13 行	0 行	6 行

结论：三档质量同档 ⇒ 按负载形态选引擎，别按量化位数选。

---

五、顺手修的两个客户端问题（都是配置，不是模型）

症状	根因	修法
「雷霆思考」但看不到思考过程	网关配置里硬编码 chat_template_kwargs.enable_thinking: false	客户端请求级带 chat_template_kwargs（enable_thinking）覆盖配置默认值；实测 true → 返回 reasoning_content
发图提示「此模型不支持」	客户端 model 的 input 只声明了 text	加 image + requestImagePixelBudget: 1048576 + requestImageMaxBytes: 8388608

两个额外认知：

• 顶层 reasoning_effort 对本模型无效——模板只认 chat_template_kwargs 内的 enable_thinking + reasoning_effort，且后者只收 low/medium/xhigh（传 high/max/none 直接 400）。
• 别为了「清爽」关思考：thinking off 会掉准确性（同一题 9.9 vs 9.11，关思考答错）⇒ 默认档位定 medium。

---

六、HiCache 归因翻案：10 分钟插桩，找回那 15 秒

6.1 推翻旧结论

• 旧现象：开 HiCache 后 8K 冷 prefill 从 5.3s 涨到 21s。
• 真因：当时打的自定义 write-through 补丁（让 chunked prefill 也写回）。
• 回到 pristine 源码 + --hicache-write-policy write_back 后：8K 冷 prefill 5.44~5.48s vs 基线 5.43s（一致），GPU 曲线 prefill 期对称 100%/100%、空闲 0%。

6.2 驱逐成本结构（方法级探针实测）

环节	耗时	占比
_evict_components 合计（39 次）	94.20s	100%
└ writing_check（阻塞等待）	93.95s	99.7%
└ _execute_kv_backup（真正的 D→H 拷贝）	0.81~1.24s（单次 26~28ms）	~1%
└ 树遍历 / _demote	≈0	—

• 阻塞点在 writing_check(write_back=True) 里的 while ongoing_write_through: ack.finish_event.synchronize()。
• 但端到端 wall-time 只 +1.3s——这个等待与排队的 GPU 计算重叠，所以「驱逐成本」不体现在耗时上。

6.3 最终 A/B（同电池，EAGLE on，pristine）

场景	HiCache OFF	HiCache ON (write_back)
冷 8K prefill	5.46~5.66s	5.48~5.56s
冷 64K prefill	60.76s	60.90s
热 64K（在池内）	0.79s	0.75s
被逐出的 64K 重命中	63.01s（冷重算）	2.42s（主机池 load_back）= 26×
重复请求（不驱逐）	—	8K 12×（5.45→0.46s）/ 64K 76×（60.4→0.79s）

结论：HiCache write_back 不增加 prefill 成本，唯一差异就是「被逐出文档的二次命中」——从冷重算变成主机层回迁（26 倍）。主机池 418,344 tok/卡。

---

七、意外发现：KV 池缩水的真凶是 EAGLE，不是 HiCache

配置	设备 KV 池
EAGLE 关	190,645 tok
EAGLE 开	75,761 tok

• 证据：sglang_mtp2.log（EAGLE + HiCache off）= 75,761；所有非 EAGLE 运行 = 190,645。
• 后果：EAGLE 开启时单请求 > ~75.7K token 直接被拒（128K 请求 → HTTP 400 Input length exceeds the maximum allowed length (75755 tokens)），尽管 --context-length 196608。
• 认知：投机解码买的是速度，付的是 KV 容量——草稿层 + 打包 KV 层占掉设备池（本机打包后 target_layers=16 + draft_layers=1 = 17）。

---

八、生产栈上线（开机自启）

8.1 服务单元

[Unit]
Description=SGLang Qwen3.8-27B W8A16 + EAGLE + HiCache (双 7900XTX, :8090)
Wants=gpu-power-cap.service network-online.target
After=gpu-power-cap.service network-online.target
Conflicts=sglang.service sglang-p2p.service llama-q6-xtx.service llama-q8-tensor.service

[Service]
Type=simple
User=yalu
Group=yalu
SupplementaryGroups=render video
WorkingDirectory=/home/yalu/sglang
ExecStart=/home/yalu/start_sglang_w8.sh
Restart=on-failure
RestartSec=15
TimeoutStartSec=1800
LimitMEMLOCK=infinity

8.2 引擎配方

W8A16 GPTQ(exllama) + EAGLE MTP-3 + HiCache write_back(24G, io-backend kernel) + BF16 KV + TP=2，:8090，served-model-name=qwen3.8-27b-w8a16，--context-length 196608 --mem-fraction-static 0.91 --max-running-requests 4 --max-mamba-cache-size 20。

8.3 GPU 枚举顺序漂移 → 按 PCI device id 自探测

重启后 0/1 与 2/3 谁在前不固定（本次重启 device0/1=R9700、device2/3=XTX），硬编码索引会加载失败。启动脚本扫 /sys/class/drm/card*/device/device == 0x744c（7900XTX 的 PCI device id），恰好 2 张才用其索引，否则回退 2,3：

_xtx=""; _cnt=0
for _n in 0 1 2 3 4 5 6 7; do
  _f=/sys/class/drm/card$_n/device/device
  [ -f "$_f" ] || continue
  [ "$(cat "$_f" 2>/dev/null)" = "0x744c" ] && { _xtx="${_xtx:+$_xtx,}$_n"; _cnt=$((_cnt+1)); }
done
if [ "$_cnt" = "2" ]; then export HIP_VISIBLE_DEVICES="$_xtx"
else echo "[WARN] XTX 自动探测异常 (cnt=$_cnt)，回退 2,3" >&2; export HIP_VISIBLE_DEVICES=2,3; fi

通用规矩：任何写死 GPU 索引的启动脚本都是定时炸弹。

8.4 上线验证

项	实测
8K 冷 prefill	5.41s
设备 KV 池	75,761 tok
主机 KV 池	418,344 tok / 14.57G
网关链路	HTTP 200
思考 / 图片端到端	均 OK

---

九、显存账：每卡 22.5G，余量只剩 1.4G

组件	每卡占用
主模型权重（W8A16, TP=2）	14.69G
MTP 草稿头（8bit）	2.82 / 2.78G
设备 KV 池	2.46G（75,761 tok BF16）
Mamba 状态缓存	1.48G
CUDA 图（decode bs 1/2/4）	~1.02G
合计	≈22.5G（实测 22.56 / 22.31G）

• 主机侧（HiCache）：KV 主机池 418,344 tok / 14.57G + Mamba 主机缓存 9.45G（每卡）。
• 24G 卡上余量仅 1.4~1.7G，再调 --hicache-size / max-running-requests / 上下文时要留余量。

---

十、多卡扩展评估

10.1 三卡（TP=3）不可行

SGLang 支持非 2 幂 TP，但要求头数整除。本模型实测：

量	值	被 3 整除
num_attention_heads	24	✓
num_key_value_heads	4	✗（需 pad 到 6）
linear_num_key_heads	16	✗（需 pad 到 18）
vocab_size	248320	✗（需 pad 到 248322）

padding 会虚增计算头数并破坏 GPTQ group_size=128 的分组对齐 ⇒ 实际可用档位只有 TP=1 / 2 / 4。

10.2 R9700 加入张量并行：不建议

• R9700 = gfx1201（RDNA4），SGLang 的 AOT 编译目标含 gfx1201，理论可跑；
• 但与 7900XTX（gfx1100）异构通信效率差，且同样受头数整除限制 ⇒ 收益有限。

---

十一、功耗策略：只要 power cap，不要限频

单元	状态	说明
gpu-power-cap.service	enabled + active	XTX 2×272W（驱动硬件地板）、R9700 2×210W
gpu-clk-limit.service	disabled	其 SCLK 1500MHz 锁会把 XTX 从 2371MHz 砍到 1500（-37%）

跳闸风险已由 R9700 改外接供电解决 ⇒ 只留 power cap。️ 这与 09-25 的结论（「限频 + 降压才是真杠杆」）方向相反，属前提变化后的有意取舍；若再现跳闸则回退 1500MHz + VDDGFX -100mV。

---

十二、选型总结

负载形态	首选	理由
长文档首答 / 高并发 prefill	W8A16 SGLang	prefill +80%、TTFT 最低、并发聚合最高
单流长跑（深上下文 decode）	Q8_0 llama.cpp	TG 最深，128K +60% vs W4
多用户 / 多 agent 并发	SGLang（W4/W8）	continuous batching + radix/HiCache 前缀缓存红利

三条铁律：

1. 按负载形态选引擎，别按量化位数选（三档质量同档）。
2. 投机解码买速度、付 KV 容量——要长上下文单请求就得在 EAGLE 与容量间取舍。
3. 缓存命中率是与「prefill 算力」同级的杠杆——长链 Agent 天然高命中，别只看裸速度。

---

十三、遗留与下一步

• W8 挂 EAGLE 后的速度复测（首轮基准是无投机数据，decode 应有明显提升）
• KV 池 75.7K 的取舍：EAGLE 加速 vs 128K 单请求容量（或调
  --speculative-num-steps）
• 显存余量 1.4G 偏紧，
  --hicache-size / max-running-requests 再平衡
• W8A16 长文档多轮场景复测（吃满 HiCache 重命中 26 倍的红利）
• 功耗策略留观（power cap only vs 限频 + 降压）

[张光璞]

自己量化W8 模型基本上宣布死刑，双7900xtx 48g 显存 开视觉塔、草稿头之后，上下文只剩下80K。 速度50t/s prefill 1100~1500

[tmp tmp]

@suboyang 说:

客服传话：
加入了resize BAR功能的是这些型号：X99-4MF 4MT 4MTPLUS 4MFPLUS
8D3 8D4 AD3 AD4 BD3 BD4
F8D T8D F8DPLUS M-F M-T
P4T P4F PD3 PD4

最新带ARGB的TF T8 也支持ReBAR

我的x99-cd3-gaming, 有印象在bios中看到过reBar字样。不知道是不是一个东西。