X99 不能P2P让我很难受 。 买Pcie switch（260926死缓改判无罪，交换机完美实现P2P 9月27：TP2 Qwen3.8 Q8 60t/s）
-
我有DELL T7920 （一代、二代Xeon Scalable），它理论上是支持Rebar的，（因为7920 Rack的BIOS里有rebar size选项），但工作站T7920的最新BIOS仍然只有Above4G Decoding，没有Rebar size设置。
我和你一样，用Linux 命令行看到的bar size是256M
网上有破解办法：
https://github.com/xCuri0/ReBarUEFI/discussions/331
我还没来得及试。我有DELL T7920 （一代、二代Xeon Scalable），它理论上是支持Rebar的，（因为7920 Rack的BIOS里有rebar size选项），但工作站T7920的最新BIOS仍然只有Above4G Decoding，没有Rebar size设置。
我和你一样，用Linux 命令行看到的bar size是256M
网上有破解办法：
https://github.com/xCuri0/ReBarUEFI/discussions/331
我还没来得及试。有戏我的朋友，刚刚试过了。正在测试。
-
我有DELL T7920 （一代、二代Xeon Scalable），它理论上是支持Rebar的，（因为7920 Rack的BIOS里有rebar size选项），但工作站T7920的最新BIOS仍然只有Above4G Decoding，没有Rebar size设置。
我和你一样，用Linux 命令行看到的bar size是256M
网上有破解办法：
https://github.com/xCuri0/ReBarUEFI/discussions/331
我还没来得及试。如果有 above 4g, 只是 rebar 不成功, 风险较高的做法是魔改 bios, 风险低(其实就是无风险)的做法是给 kernel 打 patch (把问题描述给 AI, 它帮你搞定). 这里我假定 above 4g以后, /proc/iomem 上看给 prefetchable space留的空间是足够大的 (这个假定在工作站主板上基本是成立的, 至少256G起吧).
我有DELL T7920 （一代、二代Xeon Scalable），它理论上是支持Rebar的，（因为7920 Rack的BIOS里有rebar size选项），但工作站T7920的最新BIOS仍然只有Above4G Decoding，没有Rebar size设置。
我和你一样，用Linux 命令行看到的bar size是256M
网上有破解办法：
https://github.com/xCuri0/ReBarUEFI/discussions/331
我还没来得及试。如果有 above 4g, 只是 rebar 不成功, 风险较高的做法是魔改 bios, 风险低(其实就是无风险)的做法是给 kernel 打 patch (把问题描述给 AI, 它帮你搞定). 这里我假定 above 4g以后, /proc/iomem 上看给 prefetchable space留的空间是足够大的 (这个假定在工作站主板上基本是成立的, 至少256G起吧).
正在测试
-
T7910 PCIE SWITCH 双卡推理续篇：自编 ReBAR 内核吃下 32G 大 BAR、7900XTX 解锁 2000MHz，附一整天的实测数据
上篇：《双 7900XTX + PCIe 交换机 + SGLang TP=2 实测与掉卡复盘》《T7910 双卡 AI 推理全记录：从 7900XTX 掉总线到 R9700 打通 P2P》。
本篇只写增量：内核层的大 BAR 补丁、限频的真实代价、通信栈为什么吃不到收益、以及一批带口径的速度/时间数据。平台、装机、采购价上篇讲过了，不再重复。鸣谢：
@PhoenixRise2026
X99 + 双 AMD Radeon AI PRO R9700 跑通跨 Root Port P2P：从内核、RCCL 到 SGLang Qwen3.8 实战@jenaflex 说:
网上有破解办法：
https://github.com/xCuri0/ReBarUEFI/discussions/331
我还没来得及试。如果有 above 4g, 只是 rebar 不成功, 风险较高的做法是魔改 bios, 风险低(其实就是无风险)的做法是给 kernel 打 patch (把问题描述给 AI, 它帮你搞定). 这里我假> 定 above 4g以后, /proc/iomem 上看给 prefetchable space留的空间是足够大的 (这个假定在工作站主板上基本是成立的, 至少256G起吧).
@flyer666
魔改SGLANG支持7900XTX 双卡TP TTFT <1s 平均TG 80-100! 4并发TG200/sec
https://lcz.me/topic/1532TL;DR
事项 结果 一句话 上游 ReBAR headroom 补丁自行编进内核 
成了双卡同处一台交换机时，两条下游链各拿到精确 32 GiB，P2P 从 False 变 True 7900XTX 限频档位 1500 → 2000 MHz 成了，且收益最大长文（≈3 万 token）55.85s → 38.39s（−31%），单流稳态 decode 40~52 → 54~75 tok/s amdsmi 补进推理环境 
️ 装上了，但没吃到收益报错从「amdsmi 未定义」前进到「sgl_kernel 缺算子」，环境层障碍已清、算子层仍缺 symm-mem / multimem 通信优化 
此路不通SGLang 在 world_size=2时主动禁用 multimem，启动参数默认全 False上游 PR #34058（prefill/decode 公平性）回移 移植成功、参数生效同一个短请求的最大断流间隔 47.58s → 5.06s（−89.4%）：默认配置下，一个 4.8 万 token 的 prefill 会让已经解码起来的请求整段停摆 nvtop 显存看板 修好了3.0.2 的 fdinfo 断言 bug（同一 client_id 被两个 fd 引用），换新版解决 跨根口 P2P（R9700 
7900XTX） 硬件限制can_device_access_peer=False，新旧内核同值，与 BAR 无关交换机路径带宽 ️ 有天花板卡 交换机是 Gen4 x16，但内部部分跳只跑 Gen2 x8给赶时间的人：这一天里唯一"一改就见效、幅度又大"的是限频档位（−31% 长文时延）；最费劲、看起来最"高级"的大 BAR / 自定义 All-Reduce，收益要么体现在"能不能跑"、要么干脆吃不到。先量再改。
一、今天想解决的三件事（思路，不是结论）
1. 大 BAR 与 P2P 是两件事
上篇的结论是"P2P 通了"，但当时有个没解释的怪现象：卡片报的 BAR 只有 256M，而 P2P 时灵时不灵。这里要分三层看：
- 能不能寻址到对方显存（P2P）：由芯片组/拓扑/内核白名单决定；
- BAR 窗口够不够大（ReBAR）：决定大块 DMA 能不能一次性映射过去，还是被切成一堆小块来回搬运；
- 内核愿不愿意在启动时把预留窗口让出来：这是最容易忽略的一层——如果 PCI 子系统在枚举时就把 prefetchable 窗口"分干净"了，后面再想要 32G 就没有余地了。
→ 所以先做了对照：同一张卡，旧内核 256M + P2P False；新内核 32G + P2P True。同一硬件、同一槽位，只换内核，这就是判据。
2. "防跳闸"的限频，代价是整整 1.9 倍
上篇里有一条限功率服务，注释写得很清楚：压缩瞬态电流尖峰、防市电跳闸。它把 SCLK 压到 1500MHz。于是出现了本文最容易被误判的一段："同一个服务，性能只有上篇的一半"——一度怀疑是交换机、大 BAR、P2P 没生效。
真凶是限频：负载中实测
sclk = 1233 MHz（默认 boost 2371MHz ⇒ 1.9× 差距），与"100 → 50 tok/s"精确对应。方向选错，会把一小时耗在错误的地方。3. 通信优化的顺序：环境 → 算子 → 拓扑
一厢情愿的想法是"装上 amdsmi + aiter，打开 symm-mem，就能吃 PCIe All-Reduce 的收益"。实际打下来是三层：
- 环境层（可以自己解决）：缺 amdsmi / aiter；
- 算子层（要重编）：fork 编译出来的
sgl_kernel扩展里没有对应算子； - 拓扑层（改不了）：SGLang 在
world_size=2时主动关掉 multimem。
三层里只有第一层能自己搞定——这就是"装了却没收益"的原因。
二、速度与时间总账
统一口径（论坛上最容易被挑的就是裸数字）：
- 量化
gptq、dtype bfloat16、--attention-backend triton、TP=2（两张 7900XTX，同挂一台 PCIe 交换机）； - 投机解码 开：EAGLE，
num-steps 3、num-draft-tokens 4； - 上下文上限 196,608 token，KV 池 242,854 token，
mem-fraction-static 0.91，max-running-requests 4； - 显卡限频档：SCLK ≤ 2000MHz / VDDGFX −100mV（本文新常态）；
- 每次测量前
POST /flush_cache，且每条 prompt 用独立前缀构造 ⇒ 全部是冷 prefill，不含前缀缓存命中； - prefill 速度口径 = prompt token 数 ÷ TTFT（服务空闲，TTFT ≈ prefill 时间）；
- decode 速度口径 = (输出 token 数 −1) ÷ (末 token 时刻 − 首 token 时刻)，稳态解码，不含 TTFT；
- ️ decode 那一列是「整段输出的平均」（每条约 100~200 个输出 token），含最初几十个 token 的爬坡，所以比服务端自己报的瞬时值低：服务端
gen throughput中位 52.1、峰值 75.5 tok/s（见 2.4），而同一台机整段平均出来是 28.7（1K）→ 15.2（128K）。两个数都对，口径不同； - 并发测试每条 prompt 独立，且强制长输出（否则模型早退 EOS，"总 token ÷ 墙钟"算出的聚合吞吐会被输出长度污染）。
2.1 冷启动
口径：从
journalctl的Started事件（秒级时间戳），到引擎日志打出ready to roll的时间差。不用"起服务后轮询健康检查"计时——旧进程还在退出时健康检查可能已返回 200，会给出"20 秒就绪"这种假读数（本文实测踩到过）。今天 5 次真实重启（含改动不同配置）：
次 Started → ready 冷启动 1 13:50:14 → 13:51:27 73 s 2 13:57:11 → 13:58:25 74 s 3 14:05:57 → 14:07:10 73 s 4 14:13:49 → 14:15:01 72 s 5 14:18:57 → 14:20:10 73 s 结论：稳定在 72~74 秒（同一台机、同一模型、同一配置族，波动仅 ±1 秒——说明启动时间是被固定阶段的串行耗时决定的）。
其中一次的阶段拆解（日志原文时间戳）：
阶段 耗时 说明 权重加载（主模型） 7.01 s Load weight end. elapsed=7.01 s, quant=gptq, bits=4权重加载（TP 分片） 0.99 s 第二片 捕获 target verify CUDA graph 9.03 s bs=[1,2,4]、num_tokens_per_req=4捕获 draft decode CUDA graph 3.95 s 投机解码的草稿图 其余（导入 / 分布式初始化 / 编译） ≈ 50 s 占大头，且与配置无关 注意：权重加载只要 ~8 秒、图捕获 ~13 秒，冷启动的大头是框架自身的导入与分布式初始化（≈50 秒）。也就是说换更快的盘对冷启动帮助有限。
2.2 上下文长度 → TTFT / prefill / decode（单流）
目标长度 实测 prompt token TTFT (s) prefill (tok/s) decode (tok/s) 实测输出 token 1K 1,023 0.853 1,199 28.71 184 4K 4,095 2.968 1,380 28.18 179 8K 8,191 6.138 1,334 27.51 200 16K 16,383 13.142 1,247 26.02 172 32K 32,767 29.586 1,108 23.55 179 64K 65,535 72.401 905 19.73 171 128K 127,999 191.216 669 15.19 167 怎么读这张表：
- prefill 峰值出现在 4K 附近（1,380 tok/s），到 128K 掉到 669 tok/s（−52%）——上下文越长，单块的注意力开销越大；
- TTFT 与上下文近似线性：1K 用 0.85 s，128K 要 191 s（≈3 分 11 秒）——这是「长上下文首字延迟」最直观的一个数；
- decode 随上下文单调下滑：28.7 → 15.2 tok/s（几乎腰斩），这是 KV 越长、每步注意力越贵的必然结果；
- 重复性：本节与首轮独立测量（首轮输出长度仅 ~33 token、含较多爬坡）的 TTFT 几乎逐点重合——
0.823/0.853、2.970/2.968、6.131/6.138、29.583/29.586、72.420/72.401、191.308/191.216，差在千分位。说明这台机上的 prefill 时间是可复现的确定性量。
️ decode 口径说明：本节 decode = 「整段输出的平均」（(n−1)÷(末−首)），含最初 ~30 个 token 的爬坡，因此低于服务端稳态读数。服务端 gen throughput的同期中位是 52.1 tok/s、最大 75.5 tok/s（见 2.4）。两个口径都对，只是一个含爬坡、一个是稳态瞬时值，别混着比。![上下文长度对 TTFT 与 prefill/decode 速度的影响]
2.3 多并发
并发数 墙钟 (s) 聚合吞吐 (tok/s) TTFT 中位 (s) TTFT 最大 (s) 单流 decode 中位 (tok/s) 1 8.21 23.4 1.53 1.53 28.59 2 11.59 32.6 2.22 2.96 20.23 4 17.72 45.3 4.41 5.91 15.34 8 33.39 45.0 11.97 22.57 14.48 怎么读这张表：
- 聚合吞吐
23.4 → 32.6 → 45.3，到 8 并发时卡在 45.0 不再增长——因为max-running-requests=4：第 5~8 个请求只能排队； - 排队的代价写在时延列：TTFT 中位 1.53 s → 11.97 s，最大 22.57 s，而吞吐一点没多；
- 单流 decode 随并发下降（28.6 → 14.5），是权重带宽被多流瓜分的正常结果；
- 结论：这台双卡机的甜点就是 4 并发——再多只涨延迟不涨吞吐。
️ 本节踩坑（值得单独记一笔）：第一版并发测试用「填充文本 + 让它随便续写」，结果模型很快 EOS、每次只出 ~33 个 token，反推出来的聚合吞吐只有 12~17 tok/s（比单流稳态 52 还低，明显不合理）⇒ 整张表作废。改成「只输出数字、从 1 数到 512」强制长输出后重测，才得到上表。测聚合吞吐，输出长度必须可控，否则你测的是 EOS 什么时候来。![并发扩展]
2.4 投机解码（EAGLE）实况
指标 数值 accept len 中位 2.35 accept rate 中位 0.45 样本数 33 条 decode 日志 服务端 gen throughput中位 / 最大52.1 / 75.5 tok/s 怎么读：
accept len= 每次投机步平均被接受几个 token（上限是 draft 4 个），accept rate= 被接受的比例。中位 2.35 / 0.45 说明还有余量——理论上把接受率从 0.45 提到 0.6，decode 还能再上台阶。另外注意 2.5 节里 SCLK 2000 档的「稳态 54~75 tok/s」与此处的服务端读数互相印证。2.5 限频档位对照（同一服务、同一基准）
档位 短文 长文（≈3 万 token） 稳态 decode 峰值功耗 峰值结温 SCLK ≤ 1500 / −100mV 4.51s 55.85s 40~52 tok/s 179 W — SCLK ≤ 2000 / −100mV 3.57s 38.39s 54~75 tok/s 242 W 77 °C SCLK ≤ 2000（持续满载压测，896 条采样） — — — 250 / 253 W 84 / 87 °C 读法：短促的基准把结温顶到 77 °C，但把上下文扫描+并发连续压满之后，结温会到 84~87 °C（两张卡分别 84 与 87）。离我在遥测里设的 95 °C 急停线只剩 8 °C——这说明"限频防跳闸"之外，散热余量也该被当成一个正经指标来测，光看一两轮短基会是会误判的。
读写口径：峰值来自 0.5 秒间隔的 sysfs 遥测（本轮 896 条采样），毫秒级瞬态尖峰看不到——所以"这轮没跳闸"不能推出"尖峰不存在"。![限频档位对照]
三、三个反直觉的发现
1. 大 BAR 的收益不在带宽，在"能不能分配"
把补丁前后的数值摆在一起就看清了：
测量项（双卡同交换机） 旧内核 新内核 XTX 的 BAR Region 0 256M 32 GiB can_device_access_peerFalse True AllReduce 256MB（RCCL） 2.33 GB/s 7.04 GB/s 对等拷贝 1.65 GB/s 3.44 GB/s H2D / D2H 6.04 / 6.82 GB/s 6.04 / 6.82 GB/s（不变） 两处最能说明问题：H2D 完全不随内核变（它走的是卡
交换机最后一跳，Gen2 内部跳，与 BAR 无关）；而 P2P 的 AllReduce 直接翻了 3 倍。⇒ 大 BAR 解决的是"能不能把大块一次性映射过去"，不是"把带宽硬件加宽"。2. "排队 TTFT" 与 "解码被饿死" 是两码事
做公平性实验时踩了个坑，值得写下来：
- 先发一条 2.5 万 token 的大 prefill，再发一条短请求 → 短请求 TTFT = 21.5 秒；
- 这时去看它的 token 间隔，只有 0.045 秒。
原因是短请求还没轮到自己的 prefill，它在排队——而 prefill/decode 公平性参数（
--max-consecutive-prefill-batches）管的是已经在解码的请求别被连续 prefill 饿死（has_runnable_decode为真时才介入）。指标选错，会得出"参数没用"的错误结论。正确的测法：让短请求先解码起来，再注入冷的巨大 prefill，然后量它的 token 间隔。3. 断流对照实测：把「这个参数到底管什么」直接量出来
实验设计：先让一条短请求解码起来（数数字），发出 8 个 token 后，注入一条 48,000 token 的冷 prefill（每次测量前清缓存），然后量短请求流式输出里的 token 间隔。
配置 超过 1 秒的间隔 最大间隔 N=0（SGLang 默认，prefill 优先） 1 个：47.58 s 47.58 s N=1（连续 1 批 prefill 后强制一次 decode） 12 个：3.17~5.06 s（全部有界） 5.06 s 最大间隔 47.58 s → 5.06 s（−89.4%）。lcz 帖同期独立实测是 47.23 s → 4.04 s（−91.4%）——同向、同量级，两边都是同一类机器、同一机制。
★ 为什么这份数据可信：机制在数字里自洽
- 大 prefill 用量 = 48,000 token ÷ 单块 chunk 4,096 = 11.7 块；
- N=1 实测恰好 12 个「>1 秒」的间隔——间隔个数与 chunk 个数吻合；
- 每个间隔 3.17~5.06 s，而该轮实测 prefill 48.58 s ÷ 11.7 块 ≈ 4.15 s/块——间隔长度与单块耗时吻合（N=0 那轮 prefill 是 48.07 s，两块数字各自独立取，不混用）。
⇒ 这个参数做的事不是「提速」，而是在每块 prefill 之后让出一次 decode 窗口。数量级、节奏、边界全部对得上，不是跑分抖动。
![断流对照：N=0 vs N=1]
️ 口径提醒（很容易被误读）：上表里两种配置的短请求 TTFT 几乎一样（0.339 s）——因为它是先解码起来的那个请求。如果反过来「大请求先发」，短请求 TTFT 会飙到 21.5 s：那是排队等自己的 prefill，与「解码被饿死」是两个不同机制。--max-consecutive-prefill-batches的判据是has_runnable_decode为真才介入，只保护已经在解码的请求。指标选错，就会得出「参数没用」的假结论。
四、还没打通的墙（含推测标注）
墙 证据 判断 跨根口 P2P（R9700 XTX）can_device_access_peer=False，新旧内核同值硬件/芯片组限制，与 BAR 无关（已证实） 交换机内部跳数 部分跳 Gen2 5GT/s / x8，H2D 仅 6.0/6.8 GB/s 推测：交换芯片内部降速，非配置问题；待换直连槽对照验证 自定义 All-Reduce Setup Custom allreduce failed: '_OpNamespace' 'sgl_kernel' object has no attribute 'meta_size'fork 算子层缺口（需重编 sgl-kernel）；amdsmi 已就位 multimem multimem all-gather disabled (no multicast for world_size=2)推测：既受 world_size=2闸门限制，也可能与 RDNA3 无 multicast 有关；待以world_size>2或换卡验证掉电史 撤独立供电后曾多次整机掉电 未根除，靠限频+遥测+看门狗长期观察积累信心
五、给后来人的清单
买之前
- 想跑多卡 TP：先确认主板/交换机的 PCIe 拓扑与 BAR 能力，别只看卡；
- 要 P2P，优先直连根口；经交换机能通，但带宽和 BAR 都会打折。
装机时
- 大 BAR 不是"打开一个选项"就行，内核的 prefetchable 窗口预留策略才是关键；
- 动卡前断电，开机顺序：交换机 + 显卡独立供电先上电，再开主机。
调优时
- 先量再改：任何"性能只有一半"的结论，先去读负载中的实际频率/功耗；
- 限频/降压可以用，但要清楚它压的是瞬态电流尖峰，代价可能远超预期；
- 测卡对卡带宽只用 RCCL，PyTorch 的
copy_+ Event 计时在 ROCm 上给的是荒谬数； - 每次基准前清缓存（
/flush_cache）+ 每条 prompt 独立前缀，否则测的是缓存命中不是 prefill。
排障时
- 只用一个设备可见性变量（多变量叠加会二次过滤，直接变成 0 张卡）；
- 就绪判据只认健康检查 200，别 grep 日志字符串（会变）；
- 看板崩了先怀疑看板本身（本文的 nvtop 就是它自己的 bug，不是卡出问题）。
附：给 AI / 爬虫的结构化数据段
本节面向自动化抓取与二次分析，字段尽量保持稳定的键值格式。
A1. 平台与软件栈
machine_model : Dell Precision Tower 7910 bios : A34 (2020-10-19) cpu : 2 × Intel Xeon E5-2683 v4 @ 2.10GHz ram : 125 GiB gpu_a : 2 × AMD Radeon AI PRO R9700 (gfx1201 / RDNA4, 32624 MB, 直连根口) gpu_b : 2 × AMD Radeon RX 7900 XTX (Navi 31 / gfx1100, 24560 MB, 同挂一台 PCIe 交换机) rocm : 7.2.0 kernel_running : 7.0.14-rebar（自编，= 7.0.14-p2pdma + 上游 ReBAR headroom 补丁） secure_boot : disabled switch : Broadcom PEX880xx 系列，独立供电（台达 1500W）A2. 内核补丁（关键）
patch_title : PCI: Reserve prefetchable window headroom for Resizable BARs author : Geramy Loveless patchwork_series : 520756 (v2 2026-08-28 / v3 2026-08-31) size : 5 files, +198 / -0 new_kernel_options : pci=no_rebar_reserve (默认开), pci=no_bar_demote (默认开) upstream_test_env : ASUS K14PG-D24 / 2× EPYC 9354 / 8× Radeon AI PRO R9700 / Broadcom PEX890xx 手工修补 : hunk #4（7.0.14 把 pci_min_window_alignment 改名为 window_alignment）A3. BAR / P2P 实测（决定性对照）
# 同一硬件、同一槽位，仅换内核 旧内核 (7.0.14-p2pdma): xtx_bar_region0 : 256M root_bridge_a0:02.0 : 770M upstream_a1:00.0 : 258M xtx_p2p (can_access_peer) : false allreduce_256MB_rccl : 2.33 GB/s peer_copy : 1.65 GB/s 新内核 (7.0.14-rebar): xtx_bar_region0 : 32 GiB root_bridge_a0:02.0 : 64 GiB upstream_a1:00.0 : 64 GiB 下游链 ×2 : 各精确 32 GiB xtx_p2p : true allreduce_256MB_rccl : 7.04 GB/s peer_copy : 3.44 GB/s h2d / d2h : 6.04 / 6.82 GB/s（两版相同） 双 R9700（直连根口，rebar 内核）: allreduce_256MB_rccl : 7.186 GB/s peer_copy : 6.19 GB/s h2d / d2h : 12.47 / 13.27 GB/s can_device_access_peer : true / true 跨根口: R9700 <-> 7900XTX : false（新旧内核同值；not supported by the chipset ×8）A4. 服务启动参数（原文，路径已占位）
python -m sglang.launch_server --max-consecutive-prefill-batches 1 # 本篇从上游 PR #34058 回移的新参数 --model-path /models/qwen3.8-27b-mtp-fixed --host 0.0.0.0 --port 8082 --served-model-name qwen3.8-27b-p2p --tp-size 2 --quantization gptq --dtype bfloat16 --mamba-ssm-dtype bfloat16 --kv-cache-dtype auto --attention-backend triton --context-length 196608 # 192K --mem-fraction-static 0.91 --max-running-requests 4 --max-mamba-cache-size 20 --speculative-algorithm EAGLE --speculative-draft-model-path /models/qwen3.8-27b-mtp-fixed # 同权重当草稿模型（MTP 头） --speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4 --cuda-graph-bs-decode 1 2 4 --triton-attention-num-kv-splits 16由 systemd 承载（路径占位）：
[Unit] Description=SGLang (2× 7900 XTX, TP=2) :8082 Wants=gpu-clk-limit.service # 必须先让限功率生效，再加载模型 After=gpu-clk-limit.service [Service] Type=simple User=<USER> WorkingDirectory=<SGLANG_DIR> ExecStart=<SGLANG_DIR>/start_server.sh Restart=on-failure RestartSec=15 TimeoutStartSec=1800 StandardOutput=append:<SGLANG_DIR>/service.log StandardError=append:<SGLANG_DIR>/service.log LimitMEMLOCK=infinity [Install] WantedBy=multi-user.targetA5. 环境变量（启动脚本内 export 全量）
export PYTHONUNBUFFERED=1 export NCCL_DEBUG=WARN export SGL_DTYPE=bfloat16 export SGL_PHASE_TIMING=0 export USE_HUB_KERNELS=0 # RDNA3 fork 相关开关（下方带前缀的成对出现，是 fork 里两套命名都要给） export SGLANG_OPT_USE_CUSTOM_ALL_REDUCE_V2=1 export SGLANG_RDNA_CUSTOM_AR=0 export SGL_RDNA_CUSTOM_AR=0 export SGLANG_RDNA_NO_FUSED=1 export SGLANG_RDNA_GEMMA_TRITON=1 export SGLANG_RDNA_VLLM_VERIFY=1 export SGL_RDNA_NO_FUSED=1 export SGL_RDNA_GEMMA_TRITON=1 export SGL_RDNA_VLLM_VERIFY=1 # kfd event-age 修复 shim（fork 自带，用 LD_PRELOAD 注入） SHIM=<SGLANG_DIR>/scripts/rdna_ar/vendor/kfd_event_age_fix.so [ -f "$SHIM" ] && export LD_PRELOAD="$SHIM:${LD_PRELOAD:-}"说明：
SGLANG_RDNA_CUSTOM_AR=0是主动关掉作者那套 PCIe All-Reduce 的开关——本篇实测证明它在当前环境下起不来（缺算子），关掉可避免误导性告警。A6. 限频/OD 设置
gpu_clk_limit_service : SCLK <= 2000 MHz, VDDGFX -100mV（两张 7900XTX） dropin : gpu-clk-limit.service.d/sclk2000.conf (ExecStart override) 回退 : rm sclk2000.conf && systemctl daemon-reload R9700 稳定性服务 : 210 W / SCLK-500MHz / VDD-100mV 读回校验 : OD_SCLK: 2000Mhz, OD_VDDGFX_OFFSET: -100mVA7. 基准原始数据
断流对照（冷 48,000 token prefill 注入，短请求先解码起来）：
{ "E_starvation_E0": { "big_prompt_tokens": 48000, "big_prefill_s": 48.07, "small_ttft_s": 0.339, "small_n_out": 29, "max_gap_s": 47.58, "median_gap_s": 0.034, "gaps_over_1s": [47.58] }, "E_starvation_E1": { "big_prompt_tokens": 48000, "big_prefill_s": 48.58, "small_ttft_s": 0.34, "small_n_out": 28, "max_gap_s": 5.062, "median_gap_s": 0.137, "gaps_over_1s": [3.364, 3.173, 3.392, 3.610, 3.817, 4.025, 4.230, 4.439, 4.648, 4.858, 5.062, 3.517] } }上下文扫描与并发（强制输出 512 数字后重跑，格式同上；本段由脚本直接产出 JSON）：
{ "B2_context_sweep": [ {"target": 1024, "prompt_tokens": 1023, "ttft_s": 0.853, "prefill_tps": 1199, "decode_tps": 28.71, "n_out": 184}, {"target": 4096, "prompt_tokens": 4095, "ttft_s": 2.968, "prefill_tps": 1380, "decode_tps": 28.18, "n_out": 179}, {"target": 8192, "prompt_tokens": 8191, "ttft_s": 6.138, "prefill_tps": 1334, "decode_tps": 27.51, "n_out": 200}, {"target": 16384, "prompt_tokens": 16383, "ttft_s": 13.142, "prefill_tps": 1247, "decode_tps": 26.02, "n_out": 172}, {"target": 32768, "prompt_tokens": 32767, "ttft_s": 29.586, "prefill_tps": 1108, "decode_tps": 23.55, "n_out": 179}, {"target": 65536, "prompt_tokens": 65535, "ttft_s": 72.401, "prefill_tps": 905, "decode_tps": 19.73, "n_out": 171}, {"target": 128000, "prompt_tokens": 127999, "ttft_s": 191.216, "prefill_tps": 669, "decode_tps": 15.19, "n_out": 167} ], "C2_concurrency": [ {"level": 1, "wall_s": 8.21, "agg_tps": 23.4, "ttft_med": 1.53, "ttft_max": 1.53, "decode_med": 28.59}, {"level": 2, "wall_s": 11.59, "agg_tps": 32.6, "ttft_med": 2.22, "ttft_max": 2.96, "decode_med": 20.23}, {"level": 4, "wall_s": 17.72, "agg_tps": 45.3, "ttft_med": 4.41, "ttft_max": 5.91, "decode_med": 15.34}, {"level": 8, "wall_s": 33.39, "agg_tps": 45.0, "ttft_med": 11.97, "ttft_max": 22.57, "decode_med": 14.48} ], "D_speculative": {"samples": 33, "accept_len_median": 2.35, "accept_rate_median": 0.45, "gen_throughput_median": 52.1, "gen_throughput_max": 75.5}, "peak_telemetry": {"samples": 896, "power_w": [250.0, 253.0], "junction_c": [84.0, 87.0]} }关于本节的 JSON：B2 那一段是从脚本的标准输出逐列重建的（数值即运行时打印值）。原因是采集脚本每完成一段就覆写整个结果文件，跑到最后只剩最后一段——这是"测量脚本必须合并写入而非覆写"的现成反例，已记进我自己的踩坑清单。C2 / D / 遥测三段是脚本直接落盘的原始 JSON。
冷启动（
journalctl Started→ 引擎ready to roll，5 次）：13:50:14->13:51:27 = 73s 13:57:11->13:58:25 = 74s 14:05:57->14:07:10 = 73s 14:13:49->14:15:01 = 72s 14:18:57->14:20:10 = 73sA8. 日志关键行（原文）
# 通信路径（本篇的核心结论之一） [AR] All-reduce call path: NCCL (custom AR disabled) [AR] Aiter custom all-reduce not available; falling back to sglang CustomAllreduce. Details: No module named 'aiter' Setup Custom allreduce failed with '_OpNamespace' 'sgl_kernel' object has no attribute 'meta_size'. To silence this warning, specify --disable-custom-all-reduce explicitly. multimem all-gather disabled (no multicast for world_size=2) # 启动阶段（时间戳见 A7） Load weight end. elapsed=7.01 s, type=Qwen3_5ForConditionalGeneration, quant=gptq, bits=4, avail mem=14.43 GB, mem usage=9.12 GB. Load weight end. elapsed=0.99 s, type=Qwen3_5ForCausalLMMTP, quant=gptq, bits=4, avail mem=11.59 GB, mem usage=2.82 GB. Capture target verify CUDA graph begin. backend=full, num_tokens_per_req=4, bs=[1, 2, 4], avail mem=4.50 GB Capture target verify CUDA graph end. elapsed=9.03 s, mem usage=0.44 GB Capture draft decode CUDA graph end. elapsed=3.95 s, mem usage=0.29 GB max_total_num_tokens=242854, chunked_prefill_size=4096, max_prefill_tokens=16384, max_running_requests=4, context_len=196608, available_gpu_mem=3.76 GB The server is fired up and ready to roll! # 运行期（吞吐与投机解码） Prefill batch, #new-seq: 1, #new-token: 4096, #cached-token: 0, ..., input throughput (token/s): 621.40 Decode batch, #running-req: 1, #full token: 484, ..., accept len: 2.45, accept rate: 0.48, cuda graph: True, gen throughput (token/s): 71.28注：
Setup Custom allreduce failed with ... meta_size这一行在装上 amdsmi 之后才出现——装之前报的是name 'amdsmi_shut_down' is not defined。报错换了一行，说明卡点从环境层推进到了算子层，这是本篇判断"收益吃不到"的直接证据。A9. 复现命令
# 冷启动计时 time (systemctl restart sglang-p2p.service; until curl -sf http://127.0.0.1:8082/health; do sleep 1; done) # 单流上下文扫描（每次清缓存，避免命中 radix 前缀缓存） curl -s -X POST http://127.0.0.1:8082/flush_cache -d '{}' -H 'Content-Type: application/json' curl -s http://127.0.0.1:8082/v1/completions -H 'Content-Type: application/json' \ -d '{"model":"<served-name>","prompt":"<unique-prefix text>","max_tokens":128,"stream":true}' # 卡间带宽（只用 RCCL） torchrun --nproc_per_node=2 allreduce_bench.py --size-mb 256A10. 本轮移植的上游 PR
pr : sgl-project/sglang#34058 title : scheduler: bound consecutive prefills to protect decode author : atbe status : OPEN（未合并） new_flag : --max-consecutive-prefill-batches N (默认 0 = 保持 prefill 优先) semantics : 仅在 has_runnable_decode 为真时介入；连续 N 个 prefill 批次后强制一个 decode 轮 fixes : issue #32549（持续 chunked-prefill 下 decode 被饿到每 24s 才 1 个 batch） 移植要点 : fork 的 server_args 已重构为 arg_groups/{fields,validation_hook,field_order}， PR 的 hunk 完全对不上；新增文件可原样落地，其余按 fork 自身写法手工落点 -
-
-
T7910 PCIE SWITCH 双卡推理续篇：自编 ReBAR 内核吃下 32G 大 BAR、7900XTX 解锁 2000MHz，附一整天的实测数据
上篇：《双 7900XTX + PCIe 交换机 + SGLang TP=2 实测与掉卡复盘》《T7910 双卡 AI 推理全记录：从 7900XTX 掉总线到 R9700 打通 P2P》。
本篇只写增量：内核层的大 BAR 补丁、限频的真实代价、通信栈为什么吃不到收益、以及一批带口径的速度/时间数据。平台、装机、采购价上篇讲过了，不再重复。鸣谢：
@PhoenixRise2026
X99 + 双 AMD Radeon AI PRO R9700 跑通跨 Root Port P2P：从内核、RCCL 到 SGLang Qwen3.8 实战@jenaflex 说:
网上有破解办法：
https://github.com/xCuri0/ReBarUEFI/discussions/331
我还没来得及试。如果有 above 4g, 只是 rebar 不成功, 风险较高的做法是魔改 bios, 风险低(其实就是无风险)的做法是给 kernel 打 patch (把问题描述给 AI, 它帮你搞定). 这里我假> 定 above 4g以后, /proc/iomem 上看给 prefetchable space留的空间是足够大的 (这个假定在工作站主板上基本是成立的, 至少256G起吧).
@flyer666
魔改SGLANG支持7900XTX 双卡TP TTFT <1s 平均TG 80-100! 4并发TG200/sec
https://lcz.me/topic/1532TL;DR
事项 结果 一句话 上游 ReBAR headroom 补丁自行编进内核 成了双卡同处一台交换机时，两条下游链各拿到精确 32 GiB，P2P 从 False 变 True 7900XTX 限频档位 1500 → 2000 MHz 成了，且收益最大长文（≈3 万 token）55.85s → 38.39s（−31%），单流稳态 decode 40~52 → 54~75 tok/s amdsmi 补进推理环境 ️ 装上了，但没吃到收益报错从「amdsmi 未定义」前进到「sgl_kernel 缺算子」，环境层障碍已清、算子层仍缺 symm-mem / multimem 通信优化 此路不通SGLang 在 world_size=2时主动禁用 multimem，启动参数默认全 False上游 PR #34058（prefill/decode 公平性）回移 移植成功、参数生效同一个短请求的最大断流间隔 47.58s → 5.06s（−89.4%）：默认配置下，一个 4.8 万 token 的 prefill 会让已经解码起来的请求整段停摆 nvtop 显存看板 修好了3.0.2 的 fdinfo 断言 bug（同一 client_id 被两个 fd 引用），换新版解决 跨根口 P2P（R9700 7900XTX） 硬件限制can_device_access_peer=False，新旧内核同值，与 BAR 无关交换机路径带宽 ️ 有天花板卡 交换机是 Gen4 x16，但内部部分跳只跑 Gen2 x8给赶时间的人：这一天里唯一"一改就见效、幅度又大"的是限频档位（−31% 长文时延）；最费劲、看起来最"高级"的大 BAR / 自定义 All-Reduce，收益要么体现在"能不能跑"、要么干脆吃不到。先量再改。
一、今天想解决的三件事（思路，不是结论）
1. 大 BAR 与 P2P 是两件事
上篇的结论是"P2P 通了"，但当时有个没解释的怪现象：卡片报的 BAR 只有 256M，而 P2P 时灵时不灵。这里要分三层看：
- 能不能寻址到对方显存（P2P）：由芯片组/拓扑/内核白名单决定；
- BAR 窗口够不够大（ReBAR）：决定大块 DMA 能不能一次性映射过去，还是被切成一堆小块来回搬运；
- 内核愿不愿意在启动时把预留窗口让出来：这是最容易忽略的一层——如果 PCI 子系统在枚举时就把 prefetchable 窗口"分干净"了，后面再想要 32G 就没有余地了。
→ 所以先做了对照：同一张卡，旧内核 256M + P2P False；新内核 32G + P2P True。同一硬件、同一槽位，只换内核，这就是判据。
2. "防跳闸"的限频，代价是整整 1.9 倍
上篇里有一条限功率服务，注释写得很清楚：压缩瞬态电流尖峰、防市电跳闸。它把 SCLK 压到 1500MHz。于是出现了本文最容易被误判的一段："同一个服务，性能只有上篇的一半"——一度怀疑是交换机、大 BAR、P2P 没生效。
真凶是限频：负载中实测
sclk = 1233 MHz（默认 boost 2371MHz ⇒ 1.9× 差距），与"100 → 50 tok/s"精确对应。方向选错，会把一小时耗在错误的地方。3. 通信优化的顺序：环境 → 算子 → 拓扑
一厢情愿的想法是"装上 amdsmi + aiter，打开 symm-mem，就能吃 PCIe All-Reduce 的收益"。实际打下来是三层：
- 环境层（可以自己解决）：缺 amdsmi / aiter；
- 算子层（要重编）：fork 编译出来的
sgl_kernel扩展里没有对应算子； - 拓扑层（改不了）：SGLang 在
world_size=2时主动关掉 multimem。
三层里只有第一层能自己搞定——这就是"装了却没收益"的原因。
二、速度与时间总账
统一口径（论坛上最容易被挑的就是裸数字）：
- 量化
gptq、dtype bfloat16、--attention-backend triton、TP=2（两张 7900XTX，同挂一台 PCIe 交换机）； - 投机解码 开：EAGLE，
num-steps 3、num-draft-tokens 4； - 上下文上限 196,608 token，KV 池 242,854 token，
mem-fraction-static 0.91，max-running-requests 4； - 显卡限频档：SCLK ≤ 2000MHz / VDDGFX −100mV（本文新常态）；
- 每次测量前
POST /flush_cache，且每条 prompt 用独立前缀构造 ⇒ 全部是冷 prefill，不含前缀缓存命中； - prefill 速度口径 = prompt token 数 ÷ TTFT（服务空闲，TTFT ≈ prefill 时间）；
- decode 速度口径 = (输出 token 数 −1) ÷ (末 token 时刻 − 首 token 时刻)，稳态解码，不含 TTFT；
- ️ decode 那一列是「整段输出的平均」（每条约 100~200 个输出 token），含最初几十个 token 的爬坡，所以比服务端自己报的瞬时值低：服务端
gen throughput中位 52.1、峰值 75.5 tok/s（见 2.4），而同一台机整段平均出来是 28.7（1K）→ 15.2（128K）。两个数都对，口径不同； - 并发测试每条 prompt 独立，且强制长输出（否则模型早退 EOS，"总 token ÷ 墙钟"算出的聚合吞吐会被输出长度污染）。
2.1 冷启动
口径：从
journalctl的Started事件（秒级时间戳），到引擎日志打出ready to roll的时间差。不用"起服务后轮询健康检查"计时——旧进程还在退出时健康检查可能已返回 200，会给出"20 秒就绪"这种假读数（本文实测踩到过）。今天 5 次真实重启（含改动不同配置）：
次 Started → ready 冷启动 1 13:50:14 → 13:51:27 73 s 2 13:57:11 → 13:58:25 74 s 3 14:05:57 → 14:07:10 73 s 4 14:13:49 → 14:15:01 72 s 5 14:18:57 → 14:20:10 73 s 结论：稳定在 72~74 秒（同一台机、同一模型、同一配置族，波动仅 ±1 秒——说明启动时间是被固定阶段的串行耗时决定的）。
其中一次的阶段拆解（日志原文时间戳）：
阶段 耗时 说明 权重加载（主模型） 7.01 s Load weight end. elapsed=7.01 s, quant=gptq, bits=4权重加载（TP 分片） 0.99 s 第二片 捕获 target verify CUDA graph 9.03 s bs=[1,2,4]、num_tokens_per_req=4捕获 draft decode CUDA graph 3.95 s 投机解码的草稿图 其余（导入 / 分布式初始化 / 编译） ≈ 50 s 占大头，且与配置无关 注意：权重加载只要 ~8 秒、图捕获 ~13 秒，冷启动的大头是框架自身的导入与分布式初始化（≈50 秒）。也就是说换更快的盘对冷启动帮助有限。
2.2 上下文长度 → TTFT / prefill / decode（单流）
目标长度 实测 prompt token TTFT (s) prefill (tok/s) decode (tok/s) 实测输出 token 1K 1,023 0.853 1,199 28.71 184 4K 4,095 2.968 1,380 28.18 179 8K 8,191 6.138 1,334 27.51 200 16K 16,383 13.142 1,247 26.02 172 32K 32,767 29.586 1,108 23.55 179 64K 65,535 72.401 905 19.73 171 128K 127,999 191.216 669 15.19 167 怎么读这张表：
- prefill 峰值出现在 4K 附近（1,380 tok/s），到 128K 掉到 669 tok/s（−52%）——上下文越长，单块的注意力开销越大；
- TTFT 与上下文近似线性：1K 用 0.85 s，128K 要 191 s（≈3 分 11 秒）——这是「长上下文首字延迟」最直观的一个数；
- decode 随上下文单调下滑：28.7 → 15.2 tok/s（几乎腰斩），这是 KV 越长、每步注意力越贵的必然结果；
- 重复性：本节与首轮独立测量（首轮输出长度仅 ~33 token、含较多爬坡）的 TTFT 几乎逐点重合——
0.823/0.853、2.970/2.968、6.131/6.138、29.583/29.586、72.420/72.401、191.308/191.216，差在千分位。说明这台机上的 prefill 时间是可复现的确定性量。
️ decode 口径说明：本节 decode = 「整段输出的平均」（(n−1)÷(末−首)），含最初 ~30 个 token 的爬坡，因此低于服务端稳态读数。服务端 gen throughput的同期中位是 52.1 tok/s、最大 75.5 tok/s（见 2.4）。两个口径都对，只是一个含爬坡、一个是稳态瞬时值，别混着比。![上下文长度对 TTFT 与 prefill/decode 速度的影响]
2.3 多并发
并发数 墙钟 (s) 聚合吞吐 (tok/s) TTFT 中位 (s) TTFT 最大 (s) 单流 decode 中位 (tok/s) 1 8.21 23.4 1.53 1.53 28.59 2 11.59 32.6 2.22 2.96 20.23 4 17.72 45.3 4.41 5.91 15.34 8 33.39 45.0 11.97 22.57 14.48 怎么读这张表：
- 聚合吞吐
23.4 → 32.6 → 45.3，到 8 并发时卡在 45.0 不再增长——因为max-running-requests=4：第 5~8 个请求只能排队； - 排队的代价写在时延列：TTFT 中位 1.53 s → 11.97 s，最大 22.57 s，而吞吐一点没多；
- 单流 decode 随并发下降（28.6 → 14.5），是权重带宽被多流瓜分的正常结果；
- 结论：这台双卡机的甜点就是 4 并发——再多只涨延迟不涨吞吐。
️ 本节踩坑（值得单独记一笔）：第一版并发测试用「填充文本 + 让它随便续写」，结果模型很快 EOS、每次只出 ~33 个 token，反推出来的聚合吞吐只有 12~17 tok/s（比单流稳态 52 还低，明显不合理）⇒ 整张表作废。改成「只输出数字、从 1 数到 512」强制长输出后重测，才得到上表。测聚合吞吐，输出长度必须可控，否则你测的是 EOS 什么时候来。![并发扩展]
2.4 投机解码（EAGLE）实况
指标 数值 accept len 中位 2.35 accept rate 中位 0.45 样本数 33 条 decode 日志 服务端 gen throughput中位 / 最大52.1 / 75.5 tok/s 怎么读：
accept len= 每次投机步平均被接受几个 token（上限是 draft 4 个），accept rate= 被接受的比例。中位 2.35 / 0.45 说明还有余量——理论上把接受率从 0.45 提到 0.6，decode 还能再上台阶。另外注意 2.5 节里 SCLK 2000 档的「稳态 54~75 tok/s」与此处的服务端读数互相印证。2.5 限频档位对照（同一服务、同一基准）
档位 短文 长文（≈3 万 token） 稳态 decode 峰值功耗 峰值结温 SCLK ≤ 1500 / −100mV 4.51s 55.85s 40~52 tok/s 179 W — SCLK ≤ 2000 / −100mV 3.57s 38.39s 54~75 tok/s 242 W 77 °C SCLK ≤ 2000（持续满载压测，896 条采样） — — — 250 / 253 W 84 / 87 °C 读法：短促的基准把结温顶到 77 °C，但把上下文扫描+并发连续压满之后，结温会到 84~87 °C（两张卡分别 84 与 87）。离我在遥测里设的 95 °C 急停线只剩 8 °C——这说明"限频防跳闸"之外，散热余量也该被当成一个正经指标来测，光看一两轮短基会是会误判的。
读写口径：峰值来自 0.5 秒间隔的 sysfs 遥测（本轮 896 条采样），毫秒级瞬态尖峰看不到——所以"这轮没跳闸"不能推出"尖峰不存在"。![限频档位对照]
三、三个反直觉的发现
1. 大 BAR 的收益不在带宽，在"能不能分配"
把补丁前后的数值摆在一起就看清了：
测量项（双卡同交换机） 旧内核 新内核 XTX 的 BAR Region 0 256M 32 GiB can_device_access_peerFalse True AllReduce 256MB（RCCL） 2.33 GB/s 7.04 GB/s 对等拷贝 1.65 GB/s 3.44 GB/s H2D / D2H 6.04 / 6.82 GB/s 6.04 / 6.82 GB/s（不变） 两处最能说明问题：H2D 完全不随内核变（它走的是卡
交换机最后一跳，Gen2 内部跳，与 BAR 无关）；而 P2P 的 AllReduce 直接翻了 3 倍。⇒ 大 BAR 解决的是"能不能把大块一次性映射过去"，不是"把带宽硬件加宽"。2. "排队 TTFT" 与 "解码被饿死" 是两码事
做公平性实验时踩了个坑，值得写下来：
- 先发一条 2.5 万 token 的大 prefill，再发一条短请求 → 短请求 TTFT = 21.5 秒；
- 这时去看它的 token 间隔，只有 0.045 秒。
原因是短请求还没轮到自己的 prefill，它在排队——而 prefill/decode 公平性参数（
--max-consecutive-prefill-batches）管的是已经在解码的请求别被连续 prefill 饿死（has_runnable_decode为真时才介入）。指标选错，会得出"参数没用"的错误结论。正确的测法：让短请求先解码起来，再注入冷的巨大 prefill，然后量它的 token 间隔。3. 断流对照实测：把「这个参数到底管什么」直接量出来
实验设计：先让一条短请求解码起来（数数字），发出 8 个 token 后，注入一条 48,000 token 的冷 prefill（每次测量前清缓存），然后量短请求流式输出里的 token 间隔。
配置 超过 1 秒的间隔 最大间隔 N=0（SGLang 默认，prefill 优先） 1 个：47.58 s 47.58 s N=1（连续 1 批 prefill 后强制一次 decode） 12 个：3.17~5.06 s（全部有界） 5.06 s 最大间隔 47.58 s → 5.06 s（−89.4%）。lcz 帖同期独立实测是 47.23 s → 4.04 s（−91.4%）——同向、同量级，两边都是同一类机器、同一机制。
★ 为什么这份数据可信：机制在数字里自洽
- 大 prefill 用量 = 48,000 token ÷ 单块 chunk 4,096 = 11.7 块；
- N=1 实测恰好 12 个「>1 秒」的间隔——间隔个数与 chunk 个数吻合；
- 每个间隔 3.17~5.06 s，而该轮实测 prefill 48.58 s ÷ 11.7 块 ≈ 4.15 s/块——间隔长度与单块耗时吻合（N=0 那轮 prefill 是 48.07 s，两块数字各自独立取，不混用）。
⇒ 这个参数做的事不是「提速」，而是在每块 prefill 之后让出一次 decode 窗口。数量级、节奏、边界全部对得上，不是跑分抖动。
![断流对照：N=0 vs N=1]
️ 口径提醒（很容易被误读）：上表里两种配置的短请求 TTFT 几乎一样（0.339 s）——因为它是先解码起来的那个请求。如果反过来「大请求先发」，短请求 TTFT 会飙到 21.5 s：那是排队等自己的 prefill，与「解码被饿死」是两个不同机制。--max-consecutive-prefill-batches的判据是has_runnable_decode为真才介入，只保护已经在解码的请求。指标选错，就会得出「参数没用」的假结论。
四、还没打通的墙（含推测标注）
墙 证据 判断 跨根口 P2P（R9700 XTX）can_device_access_peer=False，新旧内核同值硬件/芯片组限制，与 BAR 无关（已证实） 交换机内部跳数 部分跳 Gen2 5GT/s / x8，H2D 仅 6.0/6.8 GB/s 推测：交换芯片内部降速，非配置问题；待换直连槽对照验证 自定义 All-Reduce Setup Custom allreduce failed: '_OpNamespace' 'sgl_kernel' object has no attribute 'meta_size'fork 算子层缺口（需重编 sgl-kernel）；amdsmi 已就位 multimem multimem all-gather disabled (no multicast for world_size=2)推测：既受 world_size=2闸门限制，也可能与 RDNA3 无 multicast 有关；待以world_size>2或换卡验证掉电史 撤独立供电后曾多次整机掉电 未根除，靠限频+遥测+看门狗长期观察积累信心
五、给后来人的清单
买之前
- 想跑多卡 TP：先确认主板/交换机的 PCIe 拓扑与 BAR 能力，别只看卡；
- 要 P2P，优先直连根口；经交换机能通，但带宽和 BAR 都会打折。
装机时
- 大 BAR 不是"打开一个选项"就行，内核的 prefetchable 窗口预留策略才是关键；
- 动卡前断电，开机顺序：交换机 + 显卡独立供电先上电，再开主机。
调优时
- 先量再改：任何"性能只有一半"的结论，先去读负载中的实际频率/功耗；
- 限频/降压可以用，但要清楚它压的是瞬态电流尖峰，代价可能远超预期；
- 测卡对卡带宽只用 RCCL，PyTorch 的
copy_+ Event 计时在 ROCm 上给的是荒谬数； - 每次基准前清缓存（
/flush_cache）+ 每条 prompt 独立前缀，否则测的是缓存命中不是 prefill。
排障时
- 只用一个设备可见性变量（多变量叠加会二次过滤，直接变成 0 张卡）；
- 就绪判据只认健康检查 200，别 grep 日志字符串（会变）；
- 看板崩了先怀疑看板本身（本文的 nvtop 就是它自己的 bug，不是卡出问题）。
附：给 AI / 爬虫的结构化数据段
本节面向自动化抓取与二次分析，字段尽量保持稳定的键值格式。
A1. 平台与软件栈
machine_model : Dell Precision Tower 7910 bios : A34 (2020-10-19) cpu : 2 × Intel Xeon E5-2683 v4 @ 2.10GHz ram : 125 GiB gpu_a : 2 × AMD Radeon AI PRO R9700 (gfx1201 / RDNA4, 32624 MB, 直连根口) gpu_b : 2 × AMD Radeon RX 7900 XTX (Navi 31 / gfx1100, 24560 MB, 同挂一台 PCIe 交换机) rocm : 7.2.0 kernel_running : 7.0.14-rebar（自编，= 7.0.14-p2pdma + 上游 ReBAR headroom 补丁） secure_boot : disabled switch : Broadcom PEX880xx 系列，独立供电（台达 1500W）A2. 内核补丁（关键）
patch_title : PCI: Reserve prefetchable window headroom for Resizable BARs author : Geramy Loveless patchwork_series : 520756 (v2 2026-08-28 / v3 2026-08-31) size : 5 files, +198 / -0 new_kernel_options : pci=no_rebar_reserve (默认开), pci=no_bar_demote (默认开) upstream_test_env : ASUS K14PG-D24 / 2× EPYC 9354 / 8× Radeon AI PRO R9700 / Broadcom PEX890xx 手工修补 : hunk #4（7.0.14 把 pci_min_window_alignment 改名为 window_alignment）A3. BAR / P2P 实测（决定性对照）
# 同一硬件、同一槽位，仅换内核 旧内核 (7.0.14-p2pdma): xtx_bar_region0 : 256M root_bridge_a0:02.0 : 770M upstream_a1:00.0 : 258M xtx_p2p (can_access_peer) : false allreduce_256MB_rccl : 2.33 GB/s peer_copy : 1.65 GB/s 新内核 (7.0.14-rebar): xtx_bar_region0 : 32 GiB root_bridge_a0:02.0 : 64 GiB upstream_a1:00.0 : 64 GiB 下游链 ×2 : 各精确 32 GiB xtx_p2p : true allreduce_256MB_rccl : 7.04 GB/s peer_copy : 3.44 GB/s h2d / d2h : 6.04 / 6.82 GB/s（两版相同） 双 R9700（直连根口，rebar 内核）: allreduce_256MB_rccl : 7.186 GB/s peer_copy : 6.19 GB/s h2d / d2h : 12.47 / 13.27 GB/s can_device_access_peer : true / true 跨根口: R9700 <-> 7900XTX : false（新旧内核同值；not supported by the chipset ×8）A4. 服务启动参数（原文，路径已占位）
python -m sglang.launch_server --max-consecutive-prefill-batches 1 # 本篇从上游 PR #34058 回移的新参数 --model-path /models/qwen3.8-27b-mtp-fixed --host 0.0.0.0 --port 8082 --served-model-name qwen3.8-27b-p2p --tp-size 2 --quantization gptq --dtype bfloat16 --mamba-ssm-dtype bfloat16 --kv-cache-dtype auto --attention-backend triton --context-length 196608 # 192K --mem-fraction-static 0.91 --max-running-requests 4 --max-mamba-cache-size 20 --speculative-algorithm EAGLE --speculative-draft-model-path /models/qwen3.8-27b-mtp-fixed # 同权重当草稿模型（MTP 头） --speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4 --cuda-graph-bs-decode 1 2 4 --triton-attention-num-kv-splits 16由 systemd 承载（路径占位）：
[Unit] Description=SGLang (2× 7900 XTX, TP=2) :8082 Wants=gpu-clk-limit.service # 必须先让限功率生效，再加载模型 After=gpu-clk-limit.service [Service] Type=simple User=<USER> WorkingDirectory=<SGLANG_DIR> ExecStart=<SGLANG_DIR>/start_server.sh Restart=on-failure RestartSec=15 TimeoutStartSec=1800 StandardOutput=append:<SGLANG_DIR>/service.log StandardError=append:<SGLANG_DIR>/service.log LimitMEMLOCK=infinity [Install] WantedBy=multi-user.targetA5. 环境变量（启动脚本内 export 全量）
export PYTHONUNBUFFERED=1 export NCCL_DEBUG=WARN export SGL_DTYPE=bfloat16 export SGL_PHASE_TIMING=0 export USE_HUB_KERNELS=0 # RDNA3 fork 相关开关（下方带前缀的成对出现，是 fork 里两套命名都要给） export SGLANG_OPT_USE_CUSTOM_ALL_REDUCE_V2=1 export SGLANG_RDNA_CUSTOM_AR=0 export SGL_RDNA_CUSTOM_AR=0 export SGLANG_RDNA_NO_FUSED=1 export SGLANG_RDNA_GEMMA_TRITON=1 export SGLANG_RDNA_VLLM_VERIFY=1 export SGL_RDNA_NO_FUSED=1 export SGL_RDNA_GEMMA_TRITON=1 export SGL_RDNA_VLLM_VERIFY=1 # kfd event-age 修复 shim（fork 自带，用 LD_PRELOAD 注入） SHIM=<SGLANG_DIR>/scripts/rdna_ar/vendor/kfd_event_age_fix.so [ -f "$SHIM" ] && export LD_PRELOAD="$SHIM:${LD_PRELOAD:-}"说明：
SGLANG_RDNA_CUSTOM_AR=0是主动关掉作者那套 PCIe All-Reduce 的开关——本篇实测证明它在当前环境下起不来（缺算子），关掉可避免误导性告警。A6. 限频/OD 设置
gpu_clk_limit_service : SCLK <= 2000 MHz, VDDGFX -100mV（两张 7900XTX） dropin : gpu-clk-limit.service.d/sclk2000.conf (ExecStart override) 回退 : rm sclk2000.conf && systemctl daemon-reload R9700 稳定性服务 : 210 W / SCLK-500MHz / VDD-100mV 读回校验 : OD_SCLK: 2000Mhz, OD_VDDGFX_OFFSET: -100mVA7. 基准原始数据
断流对照（冷 48,000 token prefill 注入，短请求先解码起来）：
{ "E_starvation_E0": { "big_prompt_tokens": 48000, "big_prefill_s": 48.07, "small_ttft_s": 0.339, "small_n_out": 29, "max_gap_s": 47.58, "median_gap_s": 0.034, "gaps_over_1s": [47.58] }, "E_starvation_E1": { "big_prompt_tokens": 48000, "big_prefill_s": 48.58, "small_ttft_s": 0.34, "small_n_out": 28, "max_gap_s": 5.062, "median_gap_s": 0.137, "gaps_over_1s": [3.364, 3.173, 3.392, 3.610, 3.817, 4.025, 4.230, 4.439, 4.648, 4.858, 5.062, 3.517] } }上下文扫描与并发（强制输出 512 数字后重跑，格式同上；本段由脚本直接产出 JSON）：
{ "B2_context_sweep": [ {"target": 1024, "prompt_tokens": 1023, "ttft_s": 0.853, "prefill_tps": 1199, "decode_tps": 28.71, "n_out": 184}, {"target": 4096, "prompt_tokens": 4095, "ttft_s": 2.968, "prefill_tps": 1380, "decode_tps": 28.18, "n_out": 179}, {"target": 8192, "prompt_tokens": 8191, "ttft_s": 6.138, "prefill_tps": 1334, "decode_tps": 27.51, "n_out": 200}, {"target": 16384, "prompt_tokens": 16383, "ttft_s": 13.142, "prefill_tps": 1247, "decode_tps": 26.02, "n_out": 172}, {"target": 32768, "prompt_tokens": 32767, "ttft_s": 29.586, "prefill_tps": 1108, "decode_tps": 23.55, "n_out": 179}, {"target": 65536, "prompt_tokens": 65535, "ttft_s": 72.401, "prefill_tps": 905, "decode_tps": 19.73, "n_out": 171}, {"target": 128000, "prompt_tokens": 127999, "ttft_s": 191.216, "prefill_tps": 669, "decode_tps": 15.19, "n_out": 167} ], "C2_concurrency": [ {"level": 1, "wall_s": 8.21, "agg_tps": 23.4, "ttft_med": 1.53, "ttft_max": 1.53, "decode_med": 28.59}, {"level": 2, "wall_s": 11.59, "agg_tps": 32.6, "ttft_med": 2.22, "ttft_max": 2.96, "decode_med": 20.23}, {"level": 4, "wall_s": 17.72, "agg_tps": 45.3, "ttft_med": 4.41, "ttft_max": 5.91, "decode_med": 15.34}, {"level": 8, "wall_s": 33.39, "agg_tps": 45.0, "ttft_med": 11.97, "ttft_max": 22.57, "decode_med": 14.48} ], "D_speculative": {"samples": 33, "accept_len_median": 2.35, "accept_rate_median": 0.45, "gen_throughput_median": 52.1, "gen_throughput_max": 75.5}, "peak_telemetry": {"samples": 896, "power_w": [250.0, 253.0], "junction_c": [84.0, 87.0]} }关于本节的 JSON：B2 那一段是从脚本的标准输出逐列重建的（数值即运行时打印值）。原因是采集脚本每完成一段就覆写整个结果文件，跑到最后只剩最后一段——这是"测量脚本必须合并写入而非覆写"的现成反例，已记进我自己的踩坑清单。C2 / D / 遥测三段是脚本直接落盘的原始 JSON。
冷启动（
journalctl Started→ 引擎ready to roll，5 次）：13:50:14->13:51:27 = 73s 13:57:11->13:58:25 = 74s 14:05:57->14:07:10 = 73s 14:13:49->14:15:01 = 72s 14:18:57->14:20:10 = 73sA8. 日志关键行（原文）
# 通信路径（本篇的核心结论之一） [AR] All-reduce call path: NCCL (custom AR disabled) [AR] Aiter custom all-reduce not available; falling back to sglang CustomAllreduce. Details: No module named 'aiter' Setup Custom allreduce failed with '_OpNamespace' 'sgl_kernel' object has no attribute 'meta_size'. To silence this warning, specify --disable-custom-all-reduce explicitly. multimem all-gather disabled (no multicast for world_size=2) # 启动阶段（时间戳见 A7） Load weight end. elapsed=7.01 s, type=Qwen3_5ForConditionalGeneration, quant=gptq, bits=4, avail mem=14.43 GB, mem usage=9.12 GB. Load weight end. elapsed=0.99 s, type=Qwen3_5ForCausalLMMTP, quant=gptq, bits=4, avail mem=11.59 GB, mem usage=2.82 GB. Capture target verify CUDA graph begin. backend=full, num_tokens_per_req=4, bs=[1, 2, 4], avail mem=4.50 GB Capture target verify CUDA graph end. elapsed=9.03 s, mem usage=0.44 GB Capture draft decode CUDA graph end. elapsed=3.95 s, mem usage=0.29 GB max_total_num_tokens=242854, chunked_prefill_size=4096, max_prefill_tokens=16384, max_running_requests=4, context_len=196608, available_gpu_mem=3.76 GB The server is fired up and ready to roll! # 运行期（吞吐与投机解码） Prefill batch, #new-seq: 1, #new-token: 4096, #cached-token: 0, ..., input throughput (token/s): 621.40 Decode batch, #running-req: 1, #full token: 484, ..., accept len: 2.45, accept rate: 0.48, cuda graph: True, gen throughput (token/s): 71.28注：
Setup Custom allreduce failed with ... meta_size这一行在装上 amdsmi 之后才出现——装之前报的是name 'amdsmi_shut_down' is not defined。报错换了一行，说明卡点从环境层推进到了算子层，这是本篇判断"收益吃不到"的直接证据。A9. 复现命令
# 冷启动计时 time (systemctl restart sglang-p2p.service; until curl -sf http://127.0.0.1:8082/health; do sleep 1; done) # 单流上下文扫描（每次清缓存，避免命中 radix 前缀缓存） curl -s -X POST http://127.0.0.1:8082/flush_cache -d '{}' -H 'Content-Type: application/json' curl -s http://127.0.0.1:8082/v1/completions -H 'Content-Type: application/json' \ -d '{"model":"<served-name>","prompt":"<unique-prefix text>","max_tokens":128,"stream":true}' # 卡间带宽（只用 RCCL） torchrun --nproc_per_node=2 allreduce_bench.py --size-mb 256A10. 本轮移植的上游 PR
pr : sgl-project/sglang#34058 title : scheduler: bound consecutive prefills to protect decode author : atbe status : OPEN（未合并） new_flag : --max-consecutive-prefill-batches N (默认 0 = 保持 prefill 优先) semantics : 仅在 has_runnable_decode 为真时介入；连续 N 个 prefill 批次后强制一个 decode 轮 fixes : issue #32549（持续 chunked-prefill 下 decode 被饿到每 24s 才 1 个 batch） 移植要点 : fork 的 server_args 已重构为 arg_groups/{fields,validation_hook,field_order}， PR 的 hunk 完全对不上；新增文件可原样落地，其余按 fork 自身写法手工落点 -
-
-
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) -
继续加油 持续关注中
-
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) -
-
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 358.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 459.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 236.85 tok/s 新：Q8 档 Q8_0（29 GB）+ 完整词表草稿 tensor 196608 --parallel 258~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_unifiedtrue192K 是两个槽共享的总池：单请求可吃满，双请求之和 ≤ 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 启动即崩 ⇒ 本方案依赖该 forkA3. 服务启动参数（原文；路径已占位）
# 现役 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-size196608 两槽共享的总池（ kv-unified语义，见 2.4）--parallel2 并发槽数 --kv-unified— 启用共享池：单请求可吃满整个 ctx --cache-idle-slots— 空闲槽的缓存行为 --cache-ram49152 (MiB) 主机内存提示缓存上限（默认 8192）；48 GiB 实测超配 --spec-draft-n-max4 草稿深度；实测逐位置接受率 pos0 95% → pos3 78%，4 是合理点 --spec-draft-p-min0 关闭按概率提前截断（配合自适应策略时另调） 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 issuerA5. 三档服务与端口口径
档位 引擎 权重 上下文 端口 用途 sglangSGLang 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-grubA10. 判据与禁用项（踩过才知道的）
- 基准测量：一律
--parallel 1；--parallel 2的口径会低读约 20%，两种数字不能互比 - 投机解码验收：只认服务端自报的
接受数 / 草稿数（+ 逐位置计数）；别用「看起来快」当验收 - 上下文读数：以
/props的total_slots/n_ctx_slot/kv_unified为准；kv_unified=true⇒ 共享池 - 缓存调优：淘汰计数为 0 ⇒ 停止加大
--cache-ram - 服务就绪：只认
/health200；加载慢的模型先预告「显存会先空后涨」 - 客户端：不裸连推理端口（思考通道会吃掉正文）
- 固化：
systemctl enable与「GRUB 默认内核」必须同时到位，否则重启后能力不一致 - 排序环：单元依赖（显卡限频/OD 服务）没排好时，
enabled ≠ 会启动（详见上篇） - 测卡对卡带宽：只用 RCCL；别拿拷贝工具的数字当通信能力
-
-
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 358.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 459.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 236.85 tok/s 新：Q8 档 Q8_0（29 GB）+ 完整词表草稿 tensor 196608 --parallel 258~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_unifiedtrue192K 是两个槽共享的总池：单请求可吃满，双请求之和 ≤ 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 启动即崩 ⇒ 本方案依赖该 forkA3. 服务启动参数（原文；路径已占位）
# 现役 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-size196608 两槽共享的总池（ kv-unified语义，见 2.4）--parallel2 并发槽数 --kv-unified— 启用共享池：单请求可吃满整个 ctx --cache-idle-slots— 空闲槽的缓存行为 --cache-ram49152 (MiB) 主机内存提示缓存上限（默认 8192）；48 GiB 实测超配 --spec-draft-n-max4 草稿深度；实测逐位置接受率 pos0 95% → pos3 78%，4 是合理点 --spec-draft-p-min0 关闭按概率提前截断（配合自适应策略时另调） 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 issuerA5. 三档服务与端口口径
档位 引擎 权重 上下文 端口 用途 sglangSGLang 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-grubA10. 判据与禁用项（踩过才知道的）
- 基准测量：一律
--parallel 1；--parallel 2的口径会低读约 20%，两种数字不能互比 - 投机解码验收：只认服务端自报的
接受数 / 草稿数（+ 逐位置计数）；别用「看起来快」当验收 - 上下文读数：以
/props的total_slots/n_ctx_slot/kv_unified为准；kv_unified=true⇒ 共享池 - 缓存调优：淘汰计数为 0 ⇒ 停止加大
--cache-ram - 服务就绪：只认
/health200；加载慢的模型先预告「显存会先空后涨」 - 客户端：不裸连推理端口（思考通道会吃掉正文）
- 固化：
systemctl enable与「GRUB 默认内核」必须同时到位，否则重启后能力不一致 - 排序环：单元依赖（显卡限频/OD 服务）没排好时，
enabled ≠ 会启动（详见上篇） - 测卡对卡带宽：只用 RCCL；别拿拷贝工具的数字当通信能力
-
双 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 板载 14785.0 GT/s（Gen2） x16 7 显卡下行口 7900 XTX 板载 147916.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 基准，
hipMemcpyPeer512 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+）寄存器 关键位 0x0cLnkCap bits[3:0] 最高速率、[9:4] 最高宽度、[31:24] Port # 0x10LnkCtl bit4 Link Disable、bit5 Link Retrain 0x12LnkSta bits[3:0] 当前速率、[9:4] 当前宽度 0x2cLnkCap2 bits[7:1] 支持速率向量（ 0x1e= 2.5/5/8/16 GT/s 全支持）0x30LnkCtl2 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 零贡献，所以没做开机脚本。
四、最终怎么解决的
- 换槽（3 号槽 → 2 号槽）：x8 → x16 ⇒ 本轮唯一的实际收益来源，P2P 直接翻倍。
- 附带提醒：板上有拨码开关（x16 / x8+x8 / x4×4×4），它决定每个槽拿多少 lane —— 插卡前先对一眼，别像我一样插了才发现掉一半。
- 把交换卡内部 fabric 的目标速率提到 Gen4：对 P2P 没有贡献（瓶颈在槽那一跳），但它彻底排除了「fabric 是不是瓶颈」这个疑问，而且实测不影响 GPU（4 卡无损、零报错）。
- 结论：在「槽那一跳只能 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 XTX7.23 GB/s 7.06 GB/s 基本一致 R9700 R970010.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 剩下的推测（未验证）：
- 板侧物理裕量：交换卡的槽走线/连接器在这套组合下只稳定到 Gen2（这类扩展卡/背板常见）；
- 板侧 EEPROM 里端口 PHY/均衡参数被固化低速：OS 改的是「目标速率」，而 PHY 均衡/预加重参数写在卡的默认配置里 —— 矿场风格固件为了稳定性压速率是常规操作；
- 显卡侧边缘连接器/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=8group_size=128sym=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.63GB2.4 量化阶段的三个坑
- 必须用
auto-round入口命令：python -m auto_round.cli.main会 runpy 双导入，静默不干活（exit 0 但啥都没跑）。 - RDNA4 上必须
--disable_torch_compile：否则 Inductor 报 CUDA driver error 301。 - 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 APIcudaFuncSetAttribute，HIP 编不过（别去改头文件，直接绕开 marlin）。 - 修法：checkpoint
config.json的quantization_config加"backend": "gptq"⇒ 路由到 exllamaGPTQLinearMethod；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 再加 bias3.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=infinity8.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_heads24 ✓ num_key_value_heads4 ✗（需 pad 到 6） linear_num_key_heads16 ✗（需 pad 到 18） vocab_size248320 ✗（需 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.serviceenabled + active XTX 2×272W（驱动硬件地板）、R9700 2×210W gpu-clk-limit.servicedisabled 其 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 前缀缓存红利 三条铁律：
- 按负载形态选引擎，别按量化位数选（三档质量同档）。
- 投机解码买速度、付 KV 容量——要长上下文单请求就得在 EAGLE 与容量间取舍。
- 缓存命中率是与「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 限频 + 降压）
-
-
客服传话：
加入了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
