我的显卡是 AMD RX 7900 XT（RDNA3）。前一阵子一直在折腾本地大模型，跑通了不少东西，也留下一些没能继续走下去的方向。这篇文章记录的是后者，也就是试过、但行不通或者不划算的部分。

这块卡没有独立张量核，唯一的矩阵指令是 CU 上的 WMMA，规格只到 16×16×16，支持 fp16/bf16/int8/int4。实测 bf16 峰值约 16–17 TFLOP/s，与 NVIDIA 的专用张量核不在一个量级。fp8、tf32 都没有；TMA、cp.async、mbarrier 这类异步与同步原语也缺失。

预填（prefill）因此是算力受限，在这块卡上很难做快；解码（decode）则是带宽受限，722 GB/s，实测大约只能跑到 56%。

NInfer 移植，正确但代价递增

NInfer 是一个从零实现的 CUDA 单卡推理引擎，原本只支持 NVIDIA sm_120a。我把它移植到了 ROCm / gfx1100。

移植是成功的。能跑，输出正确，PPL 也正常。但越往深处优化，代价越大。

  • MMA 只能用 SIMT 模拟，靠 __shfl 手工实现 m16n8k16。结果正确，速度有限。
  • q8 的 grouped 内核共享内存超过 64 KiB，只能缩小 tile。
  • GDN 的 gating_proj 协作网格超出上限，需要限流。
  • fp8 KV 的预填直接崩溃，共享内存需求超过 64 KiB LDS。
  • q4/q8 的线性测试会在 ~DeviceArena 走到 hipFree 时死锁。问题不在算子，而在 ROCm 7.2.3 的运行时。

还有两条本想推进，评估之后停掉了。

  • 用 RDNA3 的 WMMA 重写。它只支持 16×16×16，fragment 布局与 PTX 的 mma 完全不同，预计需要 15–20 天，收益不划算。
  • 外挂 draft 做投机。需要改动 Engine 架构，涉及双模型、双 KV 与 draft/verify，预计 10–15 天，同样不划算。

最终停在 prefill 249、decode 41.5 tok/s（带 MTP）。可用，但继续提升的空间已经很小。

Bonsai 可用，投机不成立

Bonsai 是 Qwen 基座的三进制模型，2.13 bpw，约 7 GB。

部署本身没有问题，问题都集中在加速方向上。

  • Vulkan 后端没有三进制内核，会静默退回到 CPU，只有 4.4 tok/s，必须改用 ROCm。
  • 外挂 2B draft 做投机，净亏，55.2 降到 44.0，accept 只有 38%。
  • 自己蒸馏 draft。v1 用了 56 万 token，v2 增加到 440 万（八倍数据），accept 仍然是 38%，没有改善。
  • 用 PQ2_0 打包三进制权重后，draft 直接失效，输出乱码，accept 为 0%。三进制模型需要 QAT 或 imatrix，通用 PTQ 不适用。

投机要回本，accept 需要达到 78% 左右，而现实上限只有 38%，差距过大。这条路我放弃了。

MTP 重训，链路完整但指标不变

Qwen 自带一个 MTP 头，原本就用于投机。我想验证重训之后是否还有提升。

于是把整条链路自己搭了一遍。下载 55 GB 原始权重，写流式 int4 加载器（30 GB 内存装不下 27B，只能边读边量化）；还原 MTP 头的语义，输入是 (h_t, emb(t+1))，预测 t+2；生成数据，按 FastMTP 的思路递归训练；重建 GGUF 导出；最后放进 llama.cpp 实测。

accept 没有提升。无论用 int4 的 hidden 还是部署同款的 Q4_K hidden，无论交叉熵还是用 target 的 greedy 做自蒸馏，实际 accept 都没有变化，nmax=1 持平，nmax=2 在噪声范围内波动。

后来想清楚了。accept 取决于 draft 能否匹配 target 的 greedy，而语料中的 ground-truth 并不等于 target 的 greedy。用交叉熵把 head 推向 ground-truth，反而可能偏离 target。预训练的那个头很可能本就由自蒸馏训练得到，已经接近这个架构的合理水平。

这条路不成立。

decode 调参，KV 侧没有空间

KVMem 这条栈上相关的参数，我都扫过一遍。

  • KV 用 q4_0 还是 q8_0，没有差别。
  • KVMem 的 budget，32768 / 16384 / 8192，没有差别。
  • 投机 n_max，2 最好，再往上反而下降，accept 降低。
  • p_min、flash-attn，基本没有差别。

原因也清楚。KVMem 已经把每个 token 需要读取的 KV 限制住，解码与上下文长度、KV 配置关系不大；它只取决于权重带宽和 MTP 的质量，而这两项基本已经到顶。

试过的方向与结果

方向 结果
用 WMMA 重写加速 WMMA 太弱、布局不兼容，暂停
外挂小 draft 投机 accept 上限 38%，净亏
蒸馏 draft、增加数据 没有提升
三进制 PQ2_0 打包 draft 直接失效
重训 MTP 头 链路完整，accept 不变
调 KV 位宽 / budget 全部无效

真正有效的手段，还是那几件并不新奇的。把预填交给成熟的 GEMM 库，使用模型自带的投机头，用 KV 分层把长上下文撑住。其余优化，在这块卡上大多投入产出比很低。

以上就是这一轮的记录。