【纯闲着】2700 块的洋垃圾跑 QFN：MI50 16G×2 + 128G DDR3，从装系统到 54 tok/s，这可能是目前价格最低的洋垃圾方案。

[c jh]

我也在 整理 一个 v100 显卡 16*2 nvlink 。最低价格 也不急着买。底板用了 600多 还没有到手 ，到手了 再考虑 v100

[Sun Lyu]

@terry 说:

不充点实拍图呢，我弟，我先水一集视频。

好，晚上回去搞

[terry]

@c-jh 就等着你的帖子发视频了。

[Miraco]

这个跑Minimax H3表现如何？

[kiwi]

晚點也來發我的雙雷七。跟版大的架構差不多，夠用了～

[terry]

@kiwi 很好，又多了一期视频，最好和这个数据比较下。然后github上那个优化看看是否适用

[Sun Lyu]

调优再进一步

上一篇写到 54.0 tok/s 就收尾了。

吃晚饭的时候，我又去挖了挖潜力。

先把结论放在前面：

同一台机器、同一份模型文件、同样的量化档（IQ3_S），decode 从 54.0 提到 65.1 tok/s，峰值可以突破100tok/s，prefill 从 571 提到 752 tok/s，峰值850 tok/s。

均值65.1 是官方 2×MI50 参考机（49.2）的 1.32 倍，是我自己首跑（28.3）的 2.3 倍。

代价只有一个：引擎版本从 0.1.55 退到 0.1.52+PR。

---

一、线索：一个没人评审的 PR

上一篇之后我看到：

@rori 说:

這卡有潛力, 已經有大神做strata優化:
https://github.com/Niko1221/Strata/pull/5361

作者是 4×MI50 32GB 的机器，报出来的数字是 decode 125 tok/s、prefill 8K 1250 / 32K 1766——差不多是我 54 的两倍多。我第一反应是"32G 版显存大，全驻，不奇怪"，但里面有些数字对不上：我判负过的 --prefill-pipe，在他那儿是 +73% @8K、+109% @32K。

这个 PR 的状态很说明问题：22 个 commit、30 个文件、+5046/−147，整个 PR 只有两条评论，都是 bot 在提醒"标题要加 [Engine] 前缀"，没有一行人工评审。作者自己在描述里写着"基于 0.2.7，main 已经前进约 4k commits"。

也就是说：这东西质量未必差，但没人验过，也不会自动合进来。想用只能自己来。

---

二、第一次找错了地方

我让模型去读这个 PR 的提交列表，它挑了 e4f4027 说："这条值得参考，IQ3 的 gate/up 走 repacked GEMM，bitwise 等价，GEMM 从 0.85ms 降到 0.63ms。"

听起来很对口——我用的正是 IQ3_S。于是我把这个分支完整编译了一遍（0.1.52，165/165 目标通过），然后做了个隔离实验：STRATA_PF_RP_IQ3=0 把这个改动关掉，看差多少。

结果：

臂	配置	decode	prefill
分支默认	repack 开着	50.1	690
同一分支	STRATA_PF_RP_IQ3=0 关掉它	53.4	721

关掉反而更快。 也就是说这条提交在我这台机器上测不出收益，方向甚至是反的（都在噪声里）。

但它不是白读的——它把我引到了 PR 自带的 docs/AMD_MI50.md，那里面有一句被我（和模型）一路略过去的话：

Each new switch is off by default.
When a switch is off, that kernel runs exactly as before.

---

三、转折：收益在"默认关着"的那一整块里

docs/AMD_MI50.md 的推荐配置里，env 有 21 个开关。默认全关，全关时"kernel 和以前完全一样"。

这意味着：我前面用分支默认态跑出来的 50.1，等于只免费蹭到了 prefill 栈（默认开），decode 一个开关都没吃到。

把 21 个全打开之后：

臂（同一协议，空转 90s）	引擎	decode 中位	prefill	冷启动
上一篇的生产版	0.1.55（0.2.9）	47.0	521	75s
分支默认态	0.1.52+PR	50.1	690	84s
分支 + 关掉那条提交	0.1.52+PR	53.4	721	54s
分支 + 21 个开关全开	0.1.52+PR	61.3	701	57s

decode +30%、prefill +34%。

顺带说清楚一件事：为什么同一台机器上一篇测 54.0，这一轮基线只有 47.0。因为口径变了——上一篇收尾时用的是官方口径（1.6K prompt / 128 token / temp 0），这一轮统一成 3×512 token + 14K token 长 prefill，两者不可直接比。下面所有数字都在这一个口径里。

---

四、不等作者，自己把分支拉下来编

作者只在 PR 分支上，fork 也不会被合并。但这件事根本不需要作者配合：

curl -L --proxy http://192.168.2.2:7890 \
  -o int6.tar.gz \
  https://codeload.github.com/rpcribari/Strata/tar.gz/refs/heads/int6

31,216,899 字节，解压到 /src/Strata-int6。三个意外的好消息：

1. ggml 的 pin 没变（还是 3cf03257f2），所以我上一轮为 0.2.9 下的 /src/llama.cpp 直接复用，不用重下 37MB。
2. 这个分支自带了我上一篇手补的那两个宏，而且在 dflash2.cpp:198 把 guard 改成了 #if defined(STRATA_USE_HIP) || defined(STRATA_HIP_GFX906)——比我"补两个 #define"的土办法干净，说明作者踩过同一个坑。
3. project(strata VERSION 0.1.52)，编译 165/165 全过，约 3 分钟。

编译命令和上一篇完全一样，只把源码目录换掉：

cmake -S . -B build-906 -G Ninja \
  -DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906 \
  -DSTRATA_GGML_DIR=/src/llama.cpp \
  -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang \
  -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ \
  -DCMAKE_BUILD_TYPE=Release
ninja -C build-906 strata

注意：这里有一个反直觉的取舍——我为了 +30% 把引擎从 0.1.55 退回了 0.1.52。 新版本里 DFlash2 的改进、--draft mtp 那套新键，全都用不上了。但我是走 MTP + 流水线路线的，DFlash2 本来就用不了（上一篇第五节），所以这个代价对我是零。

---

五、空转时间本身就是个变量

社区那份 gfx906 调优台账里有一句我当时没在意的话：

重启后与每次换负载后都有约 5 个请求的冷启动爬坡，且是 10 分钟量级、值得 3.7 tok/s 的进程级效应。

于是我把同一套配置的等待时间从 90 秒拉到 540 秒，重测：

空转	decode 中位	三次样本	prefill
90 s	61.3	58.4 / 64.9 / 61.3	701
540 s	65.1	57.6 / 66.9 / 65.1	751.9

+6%，一分钱没花。

这条值得单独记住：冷启动不是"等一会儿就好"，而是一条 10 分钟的爬坡曲线。 空转 90 秒去比两个配置，等于在两条不同形状的爬坡曲线上各取一个点做比较——数字看着挺整齐，结论可能是错的。我前面几轮"某开关无收益"的判断里，至少有一部分要打这个折扣。

---

六、最后一个选项，判负

那条提交还留了个选项：STRATA_PF_RP_IQ3=2，让 IQ2_S 也走 repacked GEMM（作者报该算子 +12–21%，非 bitwise 但相对误差 1.7e-7）。

我们的 pack 是混合量化（层 0/16 是 IQ3_XXS/IQ4_NL、层 8 是 IQ2_S/Q2_0、层 24/32 是 IQ2_S/IQ4_NL、层 40 是 IQ3_S/IQ4_NL），IQ2_S 占比不小，值得一试。同样 540 秒空转：

臂	decode 中位	prefill
21 个开关	65.1	751.9
22 个（多 RP_IQ3=2）	60.8	754.4

decode −4.3、prefill 持平。 判负。作者那个 +12–21% 是在他那种"同类专家集中"的 pack 上测的，我这种层间混排的 pack 把 IQ2_S 从 MMQ 挪走反而更亏。

---

七、现在的生产配置

{
  "exe": "/src/Strata-int6/build-906/strata",
  "args": [
    "--pack", "/src/models/packs/iq3s",
    "--native", ".../IQ3_S-00001-of-00002.gguf",
    "--ple-gguf", ".../IQ3_S-00002-of-00002.gguf",
    "--expert-profile", "/src/models/learned-iq3s.bin",
    "--expert-profile-save", "/src/models/learned-iq3s-mtp.bin",
    "--expert-profile-save-every", "3",
    "--expert-cache", "auto",
    "--prefill", "auto:16384",
    "--prefill-pipe", "350",
    "--spec", "3",
    "--max-context", "131072",
    "--kv", "int8", "--kv-resident", "32768",
    "--layer-split", "27", "--trim-stage-weights",
    "--mtp", "/src/models/mtp/rt",
    "--pipeline-windows", "2",
    "--pool-workers", "32",
    "--pcie-frac", "0", "--vram-reserve-mib", "1024"
  ],
  "env": {
    "STRATA_RESIDENT_HEADROOM_GIB": "6",
    "STRATA_MTP_STAGE0": "0",
    "STRATA_RESIDENT_PIN": "0",
    "STRATA_GR_NORM_FAST": "1",
    "STRATA_IQ_STAGE_GRID_MMVQ": "1",
    "STRATA_IQ_MMVQ_XLDS": "1",
    "STRATA_IQ_MMVQ_RB": "16",
    "STRATA_IQ_MMVQ_R2": "1",
    "STRATA_MMVQ_XLDS": "1",
    "STRATA_MMVQ_WPB": "8",
    "STRATA_MMVQ_Q6X4": "1",
    "STRATA_MMVQ_Q4X4": "1",
    "STRATA_EXP_X4": "1",
    "STRATA_EXP_GU_H2": "1",
    "STRATA_EXP_LDS_Q50": "1",
    "STRATA_GR_DOWN_TG": "a",
    "STRATA_GR_DOWN_PF": "1",
    "STRATA_MMVF_ROWS": "1",
    "STRATA_GFX906_ATTN_QUERY_SWIZZLE": "1",
    "STRATA_GFX906_ATTN_REDUCE12": "1",
    "STRATA_HIP_SWIGLU_FUSED": "1"
  }
}

实测：decode 65.1 tok/s / prefill 752 tok/s / 128K 上下文 / 命中率 0.96–0.98 / 冷启动约 60 秒。冒烟抽查（"机械效率 + 斜面"）输出结构完整、带摩擦力的因果解释。

顺便说一句：STRATA_GFX906_ATTN_QUERY_SWIZZLE 和 STRATA_GFX906_ATTN_REDUCE12 这两个开关，我上一轮单独测它们是判负的（48.1/47.8，都低于基线 49.3）。现在它们进了全开组合，整体反而更快——这大概是"单开没用、和别的 kernel 一起才有用"的典型，也说明单开关 A/B 在 kernel 层是会骗人的。

---

八、三条可复用的教训

1. 收益常藏在"默认关闭 + 文档里一句话"里

我两次踩同一个坑：

• --adapt-every 0：官方冒烟命令就带着，我继承了，偷走 27%；
• 21 个 STRATA_*：默认全关，文档写得很清楚"off by default"，但我先只编了分支就跑，白跑了 20%。

判据很简单：一个功能如果"默认关"，它的作者就是在等你读文档。 拿到新分支，先 grep -n "env" docs/*MI50*.md，别急着 ninja。

2. 任何"某开关无效"的结论，先去日志里确认它真的生效

我在上一篇写过这条，这次又验证了一遍：--pipeline-windows 2 没有 MTP 会被静默关掉，日志里那句 is off: it needs the MTP drafter 躺在几万行里。判据就看三类行：is off / kept as asked / 启动 banner。

3. 测速的空转时间必须够长

90 秒和 540 秒能差 6%。这条对谁都成立，不是我这台老机器的特例——社区在同一份台账里写了"10 分钟量级"。

---

九、代价与回滚

诚实列出：

1. 引擎版本倒退：0.1.52+PR < 0.1.55 < 0.1.57。如果哪天官方把 PR 的 kernel 合进 main，我立刻回去。
2. 回滚成本几乎为零：/src/models/strata-iq3smtp.json（0.1.55 + MTP + 流水线，54.0）和 /src/Strata/build-906/strata 都留着，改个 exe 路径重启就能退回。
3. 分支是 untested 的：一个人维护、没人评审、基于 0.2.7。我只能用"编译通过 + 冒烟正常 + 输出结构正确"来担保，长跑稳定性还没验证（这点和上一篇的待办一样）。

---

十、第 9 个台阶后的横向对比

平台	卡	decode	备注
本机（本轮）	2×MI50 16G	65.1	int6 分支 + 21 开关 + IQ3_S + MTP + 流水线
本机（上一篇）	2×MI50 16G	54.0	0.1.55 + MTP + 流水线
官方参考机	2×MI50 16G	49.2	IQ3_XXS，Ryzen 5700X + DDR4-3600
社区同卡	2×Radeon Pro VII 16G	47.8	IQ3_S，Ryzen 5 2600 + DDR4-2133
单卡 4080 SUPER	1×16G	58.8	IQ3_XXS，i7-12700 + DDR4-3200
同机 llama.cpp	2×MI50 16G	26.3–26.6	Qwen3.8-27B UD-Q4_K_M

我这台 2700 块的洋垃圾，现在的 decode 已经超过了单卡 4080 SUPER 的公开数据（58.8），而那张卡的价格够配三台这机器。

而且这是在 DDR3-1333 约 30 GB/s 的前提下跑出来的。这也侧面说明：MoE 的 decode 瓶颈主要不在显存带宽，而在专家回落到内存时的那条路有没有被榨干——这条路一旦通了，便宜的 DDR3 也能顶住。

---

十一、待办（更新）

1. 质量仍然只有冒烟和风格样本，没有正式评测。 3bit 在中文长文上的失效模式是重复、断段、术语漂移，必须真材料 A/B 才算数——这是我现在唯一的硬闸门。
2. 长跑稳定性：分支版连续跑满 24 小时没有掉速/OOM 才算数。
3. 没测并发（官方与社区都说这条路径并发不行）。
4. --prefill-pipe 在"专家全驻"场景（换 Coder IQ1_M 那种塞得进 32GB 显存的档）还没复测——PR 里它 +73% @8K 的前提就是全驻。
5. 层切分仍只用了官方给的 27/21，没穷举。

---

十二、一句话总结

上一篇的 54 是"把开关打开"，这一篇的 65 是"发现开关根本没打开"。

真正值钱的不是那 11 tok/s，而是这两次都指向同一件事：这套引擎的默认值是给"能跑起来"准备的，不是给"跑得快"准备的。--adapt-every 0 是一处，21 个 env 是另一处，下一处大概率还写在某个没人读的 docs/*MI50*.md 里。

项目地址：Niko1221/Strata
分支：rpcribari: int6

---

[terry]

可持续更新优化，为后来者趟路。

[jerry deng]

我只能说太牛了，这套调优其实就是“把每一块硬件的脾气摸透，让它们各司其职”：

1. 贵的显存跑计算，便宜的内存当仓库：把需要高频计算的核心零件死死锁在极速的显存里，把体积庞大但只需要查询的辅助数据塞进便宜的内存里，用最少的钱买到了最大的容量。
2. 每次只动一小部分，不搞全员大集合：利用 MoE 架构“激活小部分”的特点，每次生成只调用一小撮专家，把原本几十个人才能干的活压缩到几个人干，直接避开了内存搬运速度慢的死穴。
3. 干掉假把戏，开关全拉满：把软件里那些表面开着、实际在偷懒的机制通通揪出来，同时把隐藏在底层的加速开关全部手动打开，不让硬件浪费哪怕一微秒。
4. 耐心等机器“热身”好：就像跑马拉松前要热身一样，刚开机数据还没理顺，等它空转跑个几分钟、把缓存全部命中之后，性能就能稳稳冲到最高点。

[Xiaote]

本帖参与讨论人数已达 10+，入选热议，作者 +3 积分奖励！

奖励凭证：热议 +3 分 · 编号 h2896-4934-1

[jerry deng]

我觉得把一些东西放在内存当查表，这个思路是最核心的

[kiwi]

版大的配置來拿部署夠好用了！

目前在測試長一點上下文跟併發。

[Sun Lyu]

@kiwi 哈哈，我现在觉得这配置这价格，真香，把Ai本地的价格打下来，哈哈哈哈

[Tzung-Ru Tsai]

最早32G版本的價格可是跟標題現在的16G差不多，前一陣子突然快跟V100 32G看齊了，有Strata 和 LLM搞定後台，真的變成有VRAM便是娘了

[Sun Lyu]

暂时最后一弹：Q3.8 FN参数到顶之后，我换了3.6 35B A3B对比了一下

（2×MI50 16GB + DDR3-1333 的第三轮）

QFN参数扫完，这台机器停在 decode 67.83 tok/s、窗口 160K、专家命中 97–99%。

结论：

• 现役定版跑出 decode 67.83（三个样本 64.94 / 67.83 / 70.27）、长 prefill 716–727 tok/s、18K 冷预填 736.6 tok/s，窗口 160K
• q36 最好成绩是短上下文 98.8 tok/s、18K 冷预填 962 tok/s、18K 复用后 38.2 tok/s，窗口 32K
• 继续用 Flash-Next。q36 赢的那一格现役已经够用，其余都在输

数字是怎么量的

decode 走流式请求，取 usage 自报的计数，同一配置连发三次取中位。长预填用同一份 11,018 token 的 prompt。短 prompt 一组 185 token、一组 515 token，只看首字节。18K 一组是 87 KB 中文材料，先冷跑一轮，再在同一会话里追问一轮，两轮分开记。空转 180 秒，前面带六条预热。

改完配置要核对自检行。/health 在加载期就返回 ok，看不出开关有没有生效。MTP 要看到 draft layer loaded，流水线要看到 two verifiers per stage，env 类只能读进程环境：

docker exec strata-build bash -c 'tr "\0" "\n" < /proc/<引擎pid>/environ | grep CPU_SHARE'

生产上还能动的地方

第 22 个 env 买的是首字节

STRATA_PREFILL_CPU_SHARE=auto 把预填阶段的一部分专家行分给 CPU 线程池，跟 GPU 并行推进。短 prompt 本来就在等 CPU 算专家，收益应该落在这里。

指标	21 个 env	22 个 env（现役）	变化
decode 中位（3×512）	70.49	67.83	−3.8%
长预填 11,018 token	710.3 / 746.6	726.8 / 716.3	持平
18K 冷预填	774.2 tok/s	736.6 tok/s	−4.9%
短 prompt 185 token	1.527 s	1.326 s	−13.2%
短 prompt 515 token	2.226 s	1.861 s	−16.4%

decode 和长预填的差都在噪声带里，同一个臂连发三次就能从 62.32 摆到 71.38。稳的只有首字节，三个样本 1.516 / 1.326 / 1.314，摘掉这个 env 是 1.687 / 1.517 / 1.527。

所以它留在生产里，理由是交互的时候人等的是第一个字。它不提吞吐，这一条不该写成提速。

env 从 21 项变 22 项，备份在 strata-iq3s-prod-pre-cpushare.json，restart-prod.sh 一条命令上去，WARM_DONE 17.8s，冒烟正常。

窗口定在 160K

模型原生 262K。192K 也能跑，同批对比摆在下面。

配置	max-context / kv-resident	启动	短上下文 decode	长预填
192K	196608 / 65536	66.1 s	59.3	803.2 @19K、901.4 @53K
160K（定版）	163840 / 54624	60.3 s	60.6	744.1 @11K

这一组是同一批测的，只读相对差，绝对值看开头那行。

192K 那版还跑过一组长上下文。

prompt	decode	prefill
29 token	59.3	—
19,026 token	42.0	803.2
53,125 token	52.7	901.4

上下文从 29 token 涨到 5.3 万，decode 掉 11%。这笔开销跟 KV 放不放得下无关，是 attention 本身随长度线性涨价。160K 比 192K 短上下文快 1.3 tok/s，启动快 6 秒，多出来的 32K 用不上，生产里最长的一次真实会话 145,985 token，占 160K 的 89%。

那条会话的日志比合成测试更有说服力。

prompt 145,985 tokens = 144,739 reused + 1,246 read in 4.7 s
1,982 generated in 36.8 s (53.9 tok/s)
decode expert cache hit rate: 97.4%

接近满窗还有 54 tok/s。

会话缓存，低的是前缀复用率

「命中率只有七成」这个印象一直挂在专家缓存上，挂错了。专家缓存常年在 0.96–0.98，低的是 --conversation-cache-* 管的前缀复用。

根因很朴素。默认四个停放槽，--conversation-cache-min-tokens 默认 0，任何请求都能占一个。几条二十来 token 的请求就把槽占满，长会话挤不进去。

--conversation-cache-mib 16384
--conversation-cache-slots 8
--conversation-cache-min-tokens 8192
--conversation-cache-park-gain 512

生效后日志里能看到短请求被挡在门外。

strata serve: conversation cache: skip parking (43 tokens, below the 8192 minimum; parked=0)

长会话的复用率跟着上到 99.2%。

代价在冷启动。缓存槽刚改完、引擎还在爬坡的时候，8,592 token 的冷预填跑过 373 秒；稳定之后 18K 冷预填 23.6 秒。重启后必须有预热，否则第一个真实请求可能卡几分钟。warm.py 已经内置进 restart-prod.sh，一次 WARM_DONE 17.7s，首请求 TTFT 从 11.9 秒降到 0.59 秒。

十六个开关臂全落在噪声里

后面又扫了十六个开关，分三批，同一口径。扫描结果折算成相对基线的差。

批次	臂	相对基线（tok/s）
草稿与管道	--spec 4	−4.1
	--spec 5	−13.8
	--spec-min-p 0.5	−0.8
	--prefill-pipe 300	−2.0
	--prefill-pipe 450	−2.3
GDN 五臂	GDN_PIPELINE / GDN_SPLIT / GDN_CHUNKED / GDN_COMMIT_SPLIT / GDN_PP=2	−1.3 / −0.8 / −3.4 / −1.5 / +0.3
KV 与 MTP 五臂	KV_PREFETCH / KV_GROW / MTP_TOP2 / MTP_FULL_HEAD / LOOKAHEAD	−2.1 / −3.4 / −3.4 / −2.3 / −1.7

掉出噪声带的只有 --spec 4 和 --spec 5。--spec 3 是唯一最优点，而且非单调，4 慢，5 更慢，6 也不快，不是越大越好。

结论很干脆：靠调参数榨不出东西了。想再快只剩两条路，改内核或者换模型。

外部路线都试过一轮

replicas 单流掉两成

同平台（E5 + X99 + 128G DDR3 + 双卡）用 replicas 拓扑的报告里，单流明显更高，就想验证一卡一副本能不能绕开跨卡握手。

拓扑	单流 decode	双流聚合	专家命中	启动
层切分（现役）	基准	57.4–61.7	0.952	87 s
replicas（一卡一副本）	−22%	66.1	0.682	120 s

replicas 输在显存。每个副本只看得见自己那张 16GB 卡，专家槽位从约 12,200 掉到 5,500 上下，命中率 0.952 → 0.682，缺口全落到 DDR3 上。同平台那台能赢，推测是同样的机制反过来，副本看得见的显存够放槽位，命中率不掉。

判定是保持层切分。我们的负载是单 agent 长会话顺序推进，单流速度决定一切；replicas 只在多人同用、每人能接受 40 tok/s 的时候有意义。

rocBLAS 解表派发不通

xxDoman/strata-gfx906 那条线里唯一带速度的是 rocBLAS solution table。库版本对不上（我们是 5.5.0，它的表是 5.2.0），只能现场调一份。

• 42 个 shape，32 个有更优解，最大 T=8192 N=1 K=2560 增益 6.36 倍；
• 表落在 /src/models/gfx906-rocblas.txt。

reader 一编进引擎，生产直接不可用。

reading the prompt: 13 of 20 tokens, 143 s
no progress for 60 s ... stopping the engine

带与不带 STRATA_ROCBLAS_TUNING 结果一样，表的问题是次要的。走不通的是 rocBLAS 在这个 shape 上的派发，缺 Tensile 内核或者首次 JIT 极慢，没证实。已回滚，重编后二进制字节数与补丁前完全一致。

单测证明「这档算得快」和「引擎能用这档」是两件事。

Avalon 能加载算不对

fraiadoc/Avalon 是 llama.cpp 的 fork。移植时改了逻辑 warp32、MMQ GCN 的 nthreads/2、host warp_size，编得过；官方 89.72 GiB 的包也能加载，25 个专家 pool 全部 register=1。

坏在 MMQ 批量路径在 gfx906 上 unspecified launch failure：ub≥16 必崩，ub≤4 才正常，加 GGML_CUDA_NO_MMQ 短路也换不来可用吞吐。90 GB 模型已删。

这条留下的教训比速度重要。随机 token 的 bench 只证明跑得动，证明不了算得对。

其余几条判负

• --kv k8v4：−7%
• --ple-io direct 和 mmap：与 ram 无差别，PLE 26.8 GiB 已经整表锁内存，每 token 只取 90 字节每行
• 中文 MTP 草稿词表：默认 106,299 token 的集合里只有 27 个汉字（共 55,328 个），但中文 decode 61.0 tok/s、英文 63.9 tok/s，只差 4%，prompt-lookup 草稿仍能命中；--add cjk 是 +3–5% 的小杠杆
• vLLM 的 gfx906 两个实现（MobyDick #53 / #34）：收益场景是 page-cache 随机 gather，我们的 PLE 已经整表锁内存，没空间
• Basalt：只支持 NVIDIA Blackwell，源码里没有 HIP 后端

转去 Qwen3.6-35B-A3B

参数到顶、窗口定死、内核路径能试的都试过之后，剩下的变量只有模型本身。Qwen3.6-35B-A3B 是官方仓库里的一等模型，22.9 GB 的 UD-Q4_K_XL 塞得进显存，总参 35B、激活 3B，理论上比 125B 总参 / 6B 激活的 Flash-Next 轻。

它按单卡写死

只读调研的证据链很硬：src/qwen35/ 里 ggml_ 零引用，没有 cudaSetDevice 和 cudaGetDeviceCount，所有权重裸 cudaMalloc，--gpu 只由父进程的 HIP_VISIBLE_DEVICES 生效。而且整个目录被 CMakeLists.txt:860 的 if(NOT STRATA_HIP_GFX906) 排除，理由写在同处注释里：它调用 wave64 兼容层没映射的 cuBLAS 入口，kernel 也假设 32-lane warp。

所以这条路得做移植，改配置没有用。按五个阶段推进：能编、单卡跑通、层切分、和 serve 共存、接 MTP。

让它能编译

最小改动两处。

1. include/strata/platform/hip_compat/cublas_v2.h 补 cublasSgemm 和 cublasGemmStridedBatchedEx 两个映射（指向 hipblas）；
2. CMakeLists.txt 加 option(STRATA_BUILD_QWEN35_GFX906 ... OFF)，条件改成 if(NOT STRATA_HIP_GFX906 OR STRATA_BUILD_QWEN35_GFX906)。

默认 OFF 是关键。实验用独立的 build-906-q36，生产构建完全不受影响。结果 167/167 编译通过，--families 从 qwen4exp 变成 qwen4exp qwen35moe。

单卡上撞三道门

门	报错	结局
专家类型白名单	blk.0.experts are IQ2_S/IQ3_S, which the native expert kernels do not take	UD-IQ3_XXS 是逐层混合量化，补表补不完，换档
32-lane warp 假设	MUL_MAT failed 加 unspecified launch failure	逻辑 warp32 改造，四处
MMQ 批量路径	ub=16/64/256 必崩，ub≤8 正常	GCN 配置表按 64 lane 波设计，改 nthreads/2

换成 UD-Q4_K_XL（22,853,663,008 B）后加载成功，25 个专家 pool 全部注册，-ub 1 下中文跑通。

层切分暴露四个跨设备 bug

单卡跑通只是及格线。两卡实现选了两个 Engine 对象各自 init()，没有去手抄七十多个指针的 scratch 快照。这一层暴露的四个 bug，单卡上永远看不到。

1. iq_embed_rows: invalid resource handle，所有入口前置 use_dev()（cudaSetDevice）；
2. k_moe_out 在 GPU0 上 Memory Fault，边界层的 wn 指向了 dev1 的 next->attn_norm，改成 load 末尾复制到 dev0；
3. 残差流没跨设备，k_moe_out 同时写 h 和残差 x，原实现只搬了 h，从边界层起全错；
4. xq_ready 跨 forward 泄漏，边界层让 mkq 保持为真，标志留在 A 上，下一个 token 复用了上一个 token 的 q8 镜像。

修完速度：pp 约 345 → 935（2.7 倍），tg 约 64 → 77（+20%）。

这一层比速度更值钱的收获，是那个不变量闸门。

同一条 prompt、--n 1，只改 --layer-split
  split 20 → gen: 2972   ← 三个切点给出三个答案
  split 22 → gen: 271
  split 23 → gen: 248044 (EOS)
  单卡     → gen: 1

修好后四者全是 token 1；--n 16 四份逐 token 相同；--dump-full-logits 显示 split 20、22 与 split 23 的 max|Δlogit| = 0。

切分类改造必须配这个测试。没有它，速度数字再漂亮也是错的。

和 serve 共存，卡住四个 token

删掉 q36_engine.cu:3980 那个 if (layer_split > 0 && serve_mode) return 2;，补上 select_all、copy_session_all、reserve_bnd，8082 上中文正确，API 与 CLI 逐 token 相同（25/25）。

会话前缀复用恒为 0。根因在模板。生成提示末尾会追加一个空思考块 <think>\n\n</think>\n\n，四个 token，而历史轮渲染默认不回放 think 块，第二轮的 prompt 与第一轮的前缀差这四个 token。match_len() 要求整段 live 是新 prompt 的前缀，引擎又没有 KV 回退，于是恒 0。

把 chat_template.jinja 里那个 preserve_thinking 条件改成 {%- if true %}，历史 assistant 轮永远回放 think 块，立刻变成 REUSE 42 of 62。

顺带一个运维坑：tokenizer 目录是容器 root 建的，宿主机改不动，只能 scp 脚本进 /src/bench/ 再用 docker exec 跑。

接 MTP 改了七处

矛盾在权重的位置上。MTP 权重在 head 引擎（dev1），验证窗口必须从拥有 layer 0 的引擎（dev0）进入。

1. commit() 循环加 if (!owns_layer(l)) continue;，每卡只推进自己那半层的 GDN 状态；
2. spec_step() 的验证窗口改由 peer ? *peer : *this 跑，之后双侧 commit 加双侧 pos 同步，只让一边 commit 会让两卡状态分叉；
3. 新增 spec() 委托到 head_e，覆盖 warmup、serve、CLI 三个调用点；
4. 新增 mtp_prompt()，prompt 的 MTP K/V 预填写到 head 引擎，拷贝全用它的 stream；
5. main() 删掉 split 下的 mtp_d = 0；
6. spec_step 入口补 use_dev()；
7. Model::embd1 给 dev1 一份嵌入副本，否则 MTP 种子行在 dev1 读 dev0 的嵌入会越界（rc=141），代价 0.50 GiB。

验证跑了 split 20 / 22 / 23 乘 mtp 0 / 3 四种配置，两套 prompt 各 16 token 全部 IDENTICAL。

两个模型的成绩

同一份材料、同一条会话、同一台机器。材料被两个 tokenizer 数成 17,408 和 18,376 个 token，不是两份材料。

指标	q36（split 22 + mtp 3）	现役 Flash-Next	谁赢
冷预填 18K	19.1 s（962 tok/s）	23.6 s（736.6 tok/s）	q36，1.31 倍
复用读入	0.125 s	0.381 s	q36
18K 冷 decode（28 token）	50.0（drafts 19/24）	53.9	现役
18K 复用后 decode（128 token）	38.2（drafts 80/144）	58.3	现役，1.53 倍
上下文窗口	32,768	163,840	现役

那 28 个 token 的样本太小，又是在冷启动的爬坡段里取的，波动比差距还大，这一格别太当真。

MTP 在 q36 身上是双刃剑。

场景	无 MTP	--mtp 3	变化
1,242 token 短上下文	77.9	98.8	+26.8%
18,376 token 冷预填	960.9	919.9	−4.3%
18K 冷 decode	44.2	50.0	+13%
18K 复用后 decode	43.7	38.2	−12.6%

--mtp 5 掉到 37.0，--mtp 6 掉到 33.9，草稿越激进越糟。

机制说得清。q36 的 decode attention 按 query 行展开，窗口越大越亏；投机解码每轮要验证多个 token，等于在长上下文里把这份成本乘若干倍。现役 Flash-Next 的 QSA 层本来就是为长窗口设计的，48 层里只有 12 层需要 KV，MTP 验证才吃得下。

两个模型最好成绩并排

维度	Qwen3.8-Flash-Next（现役）	Qwen3.6-35B-A3B（q36）
量化与体积	IQ3_S，两分片 54.8 + 28.8 GB，PLE 嵌入表走内存 mmap	UD-Q4_K_XL，22.85 GB
架构	125B 总参 / 6B 激活 MoE	35B 总参 / 3B 激活 MoE
短上下文 decode	70.27 峰值、67.83 中位	98.8（1.2K 上下文，mtp 3）
中长上下文 decode	58.3（18K，复用后）	38.2（18K，复用后）
冷预填	736.6 tok/s（18K）	962 tok/s（18K）
长 prefill（11K）	716–727 tok/s	962 tok/s（18K 口径）
上下文窗口	163,840	32,768
显存余量	两卡各约 1 GiB	dev0 1.16、dev1 1.51 GiB
代码状态	上游分支加一份配置，没改业务代码	改 5 个文件，其中 q36_engine.cu 803 行 patch，移植两天
生产运行	连跑 30 天	从未上生产

q36 领先的只有短上下文那一格，冷预填领先 31%；长上下文和反复追问都归现役，窗口大五倍。

没换的原因

我们的负载形状长在 Flash-Next 的优势区。

1. 主流程是同一份材料反复追问，是讲一讲、改一改、再问细节，同一份一万到两万 token 的材料在同一条会话里来回推进。这条路吃的是复用后 decode，58.3 对 38.2，现役快 53%。q36 赢的那条是冷预填快 4.5 秒，每份材料只发生一次。
2. 窗口是硬约束。32K 装不下一讲完整材料加多轮对话，我们真实会话已经到 14.6 万 token。q36 的可行带在 40 层里是 split 17..23，往上要么放不下要么爆显存。
3. 稳定性差着数量级。现役是上游分支加一份配置，连跑 3天；q36 是改 5 个文件移植出来的，上游一次更新就可能要重来。
4. 它的收益场景现役能覆盖。超长材料只问一轮，现役多付 4.5 秒预填就能做，窗口还是五倍。为这 4.5 秒去背一个自己维护的移植分支，不划算。
5. 想让它真赢，得改 kernel。唯一能把 q36 的 18K decode 推到现役之上的路，是改 k_attn_part 和 k_attn_prep，让窗口内 T 行共享一次 KV 读取。那是一两天的 kernel 工作，且不保证成功。

教训

切分类改造要配切点不变性测试

同一条 prompt、--n 1、只改切分点，比 argmax。q36 靠它抓出四个 bug。Avalon 那次是反例，ub 调小跑得动、一放大就崩，随机 token 的 bench 把「算不对」伪装成了「调参问题」。

基线必须同批

同一台机器、同一个二进制、同一份 args，隔一天测出来的数能差 16%。所以 A/B 要在同一批里跑，跨天、跨口径的数字不能直接减，否则会得出「越调越慢」这种结论。

默认值是给能跑起来准备的

--adapt-every 0 关掉自适应缓存层，21 个 env 默认全关，--conversation-cache-min-tokens 默认 0 让短请求占满长会话的槽。这几处默认值只管能跑起来，不管跑得好。

别人的增益不能照搬

rocBLAS 那张表在我们机器上确实成立，32/42 个 shape 有更优解，最大 6.36 倍，但派发走不通就是走不通。先核对是不是同一条代码路径、同一个瓶颈，再决定复刻。

还没做的

1. 质量评测没做，速度已经够了。
2. 长跑稳定性。int6 分支版连跑满 24 小时不掉速才算数，目前断续跑了几天，没见异常。
3. 并发没测，官方和社区都说这条路径并发不行。
4. 识图没启用。模型本身是多模态，repo 另发 907 MB 的 mmproj，引擎 0.1.52 也有 --vision，但官方把 gfx906 的图片功能列为未做，/health 仍是 images: false。没启用的能力不写进成绩。
5. 层切分只有 27 和 21 两个值，没穷举。

收尾

第一轮把机器装起来，54 tok/s。第二轮发现开关默认是关的，65。第三轮把剩下的试完，参数到顶、窗口定死、三条外部路线全部判负，于是换了个模型，跑通、跑对、跑快，然后决定不换。

最后那个决定不靠速度。q36 的 98.8 是我这台机器上出现过的最快数字，它只在 32K 窗口里成立。这台机器真正的瓶颈在窗口乘复用，不在算力。

项目地址：Niko1221/Strata
分支：rpcribari: int6
q36 实验归档：strata-research/q36-experiment/（803 行移植 patch 加全部日志）