【纯闲着】2700 块的洋垃圾跑 QFN：MI50 16G×2 + 128G DDR3，从装系统到 54 tok/s，这可能是目前价格最低的洋垃圾方案。
-
-
-
1T的NVMe硬盘不计入可不行。这个不便宜。按小黄鱼价加上电源算1000不过分。
加上1000，这也是很有性价比了。 -
這卡有潛力, 已經有大神做strata優化:
https://github.com/Niko1221/Strata/pull/5361 -
這卡有潛力, 已經有大神做strata優化:
https://github.com/Niko1221/Strata/pull/5361 -
，
T terry 固定了此主题
-
-
-
-
-
-
-
调优再进一步
上一篇写到 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
上一篇之后我看到：
這卡有潛力, 已經有大神做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/int631,216,899 字节，解压到
/src/Strata-int6。三个意外的好消息：- ggml 的 pin 没变（还是
3cf03257f2），所以我上一轮为 0.2.9 下的/src/llama.cpp直接复用，不用重下 37MB。 - 这个分支自带了我上一篇手补的那两个宏，而且在
dflash2.cpp:198把 guard 改成了#if defined(STRATA_USE_HIP) || defined(STRATA_HIP_GFX906)——比我"补两个 #define"的土办法干净，说明作者踩过同一个坑。 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 分钟量级"。
九、代价与回滚
诚实列出：
- 引擎版本倒退：0.1.52+PR < 0.1.55 < 0.1.57。如果哪天官方把 PR 的 kernel 合进 main，我立刻回去。
- 回滚成本几乎为零：
/src/models/strata-iq3smtp.json（0.1.55 + MTP + 流水线，54.0）和/src/Strata/build-906/strata都留着，改个exe路径重启就能退回。 - 分支是 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 也能顶住。
十一、待办（更新）
- 质量仍然只有冒烟和风格样本，没有正式评测。 3bit 在中文长文上的失效模式是重复、断段、术语漂移，必须真材料 A/B 才算数——这是我现在唯一的硬闸门。
- 长跑稳定性：分支版连续跑满 24 小时没有掉速/OOM 才算数。
- 没测并发（官方与社区都说这条路径并发不行）。
--prefill-pipe在"专家全驻"场景（换 Coder IQ1_M 那种塞得进 32GB 显存的档）还没复测——PR 里它 +73% @8K 的前提就是全驻。- 层切分仍只用了官方给的 27/21，没穷举。
十二、一句话总结
上一篇的 54 是"把开关打开"，这一篇的 65 是"发现开关根本没打开"。
真正值钱的不是那 11 tok/s，而是这两次都指向同一件事：这套引擎的默认值是给"能跑起来"准备的，不是给"跑得快"准备的。
--adapt-every 0是一处，21 个 env 是另一处，下一处大概率还写在某个没人读的docs/*MI50*.md里。项目地址：Niko1221/Strata
分支：rpcribari: int6
- ggml 的 pin 没变（还是
-
-
我只能说太牛了，这套调优其实就是“把每一块硬件的脾气摸透，让它们各司其职”：
- 贵的显存跑计算，便宜的内存当仓库：把需要高频计算的核心零件死死锁在极速的显存里，把体积庞大但只需要查询的辅助数据塞进便宜的内存里，用最少的钱买到了最大的容量。
- 每次只动一小部分，不搞全员大集合：利用 MoE 架构“激活小部分”的特点，每次生成只调用一小撮专家，把原本几十个人才能干的活压缩到几个人干，直接避开了内存搬运速度慢的死穴。
- 干掉假把戏，开关全拉满：把软件里那些表面开着、实际在偷懒的机制通通揪出来，同时把隐藏在底层的加速开关全部手动打开，不让硬件浪费哪怕一微秒。
- 耐心等机器“热身”好：就像跑马拉松前要热身一样，刚开机数据还没理顺，等它空转跑个几分钟、把缓存全部命中之后，性能就能稳稳冲到最高点。
-
-
我觉得把一些东西放在内存当查表，这个思路是最核心的
-
-
-
最早32G版本的價格可是跟標題現在的16G差不多，前一陣子突然快跟V100 32G看齊了，有Strata 和 LLM搞定後台，真的變成有VRAM便是娘了
-
，系统 取消固定了此主题
-
暂时最后一弹：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、hostwarp_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。
让它能编译
最小改动两处。
include/strata/platform/hip_compat/cublas_v2.h补cublasSgemm和cublasGemmStridedBatchedEx两个映射（指向 hipblas）；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 takeUD-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，单卡上永远看不到。iq_embed_rows: invalid resource handle，所有入口前置use_dev()（cudaSetDevice）；k_moe_out在 GPU0 上 Memory Fault，边界层的wn指向了 dev1 的next->attn_norm，改成 load 末尾复制到 dev0；- 残差流没跨设备，
k_moe_out同时写h和残差x，原实现只搬了h，从边界层起全错； 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）进入。
commit()循环加if (!owns_layer(l)) continue;，每卡只推进自己那半层的 GDN 状态；spec_step()的验证窗口改由peer ? *peer : *this跑，之后双侧 commit 加双侧 pos 同步，只让一边 commit 会让两卡状态分叉；- 新增
spec()委托到head_e，覆盖 warmup、serve、CLI 三个调用点； - 新增
mtp_prompt()，prompt 的 MTP K/V 预填写到 head 引擎，拷贝全用它的 stream； main()删掉 split 下的mtp_d = 0；spec_step入口补use_dev()；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.cu803 行 patch，移植两天生产运行 连跑 30 天 从未上生产 q36 领先的只有短上下文那一格，冷预填领先 31%；长上下文和反复追问都归现役，窗口大五倍。
没换的原因
我们的负载形状长在 Flash-Next 的优势区。
- 主流程是同一份材料反复追问，是讲一讲、改一改、再问细节，同一份一万到两万 token 的材料在同一条会话里来回推进。这条路吃的是复用后 decode，58.3 对 38.2，现役快 53%。q36 赢的那条是冷预填快 4.5 秒，每份材料只发生一次。
- 窗口是硬约束。32K 装不下一讲完整材料加多轮对话，我们真实会话已经到 14.6 万 token。q36 的可行带在 40 层里是
split 17..23，往上要么放不下要么爆显存。 - 稳定性差着数量级。现役是上游分支加一份配置，连跑 3天；q36 是改 5 个文件移植出来的，上游一次更新就可能要重来。
- 它的收益场景现役能覆盖。超长材料只问一轮，现役多付 4.5 秒预填就能做，窗口还是五倍。为这 4.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 倍，但派发走不通就是走不通。先核对是不是同一条代码路径、同一个瓶颈，再决定复刻。
还没做的
- 质量评测没做，速度已经够了。
- 长跑稳定性。int6 分支版连跑满 24 小时不掉速才算数，目前断续跑了几天，没见异常。
- 并发没测，官方和社区都说这条路径并发不行。
- 识图没启用。模型本身是多模态，repo 另发 907 MB 的 mmproj，引擎 0.1.52 也有
--vision，但官方把 gfx906 的图片功能列为未做，/health仍是images: false。没启用的能力不写进成绩。 - 层切分只有 27 和 21 两个值，没穷举。
收尾
第一轮把机器装起来，54 tok/s。第二轮发现开关默认是关的，65。第三轮把剩下的试完，参数到顶、窗口定死、三条外部路线全部判负，于是换了个模型，跑通、跑对、跑快，然后决定不换。
最后那个决定不靠速度。q36 的 98.8 是我这台机器上出现过的最快数字，它只在 32K 窗口里成立。这台机器真正的瓶颈在窗口乘复用，不在算力。
项目地址：Niko1221/Strata
分支：rpcribari: int6
q36 实验归档：strata-research/q36-experiment/（803 行移植 patch 加全部日志）
奖励通知：感谢提供 MI50 / MI60 的优化信息（Strata 分支），对折腾这批卡的同好很有参考价值，奖励 +3 积分。
，这个不用较真
本帖参与讨论人数已达 10+，入选热议，作者 +3 积分奖励！
