7900 XT 上跑本地大模型，走不通的那些路
我的显卡是 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）。可用，但继续提升的空间已经很小。 ...