特调版llama.cpp推荐：给Strix Halo AI MAX 395提升 43% 的Prompt 处理性能

一、Strix Halo：性能表现一直青黄不接

AMD Ryzen AI Max+ 395 “Strix Halo” 是一颗响应AI浪潮的特配 APU：堪比3070性能表现的RDNA 3.5 集显（Radeon 8060S）+ 统一 128GB LPDDR5x 内存。纸面上它有 256GB/s 的内存带宽和近 30 TFLOPS 的 GPU 算力，理论上跑本地大模型应该很有用武之地。

但实际跑起来你会发现一个重要瓶颈：Prompt 处理（prefill）。

长上下文下 prompt 的吞吐量衰减严重。以 Qwen 3.6 35B-A3B（MoE）为例，在 16K 上下文深度时，官方 llama.cpp Vulkan 后端只能跑到 ~1000 t/s，到 64K 更是跌到 ~700 t/s。官方的rocm甚至还不如这个数。相比之下，隔壁同样是128g统一内存的dgx spark，动辄5k~6k的pp让人望而生畏。

相比之下，Token 生成速度（decode）反而不是最关键的，这个东西早就受制于内存带宽让395用户躺平了，毕竟这点和dgx spark半斤八两。

二、llama.cpp 官方的尴尬：只有 Vulkan 能用

llama.cpp 官方仓库对 Windows + AMD GPU 的支持长期只停留在 Vulkan 后端。ROCm/HIP 后端的 Windows 构建一直是空白——官方 release 页面没有 ROCm 的 Windows 包。

实际上，HIP SDK 7.1.1 确实官方支持了 Strix Halo（gfx1151），并且提供了完整的编译器工具链（clang 21.0 + hipcc 7.1）。但问题是：

• ROCm 6.x ~ 7.1.x 存在 已知的 VGPR 计数 bug，会导致 gfx1151 上某些工作负载崩溃。这个 bug 直到 ROCm 7.2 才修复，而 Windows 上的 HIP SDK 还停留在 7.1.1。
• 官方 ROCm 包对 gfx1151 有 broken kernel artifacts（见 ROCm/ROCm#6042），导致 HSA 报错 Invalid Kernel Image。
• 甚至连 llama.cpp 官方推荐的柠檬水（Lemonade）预编译 ROCm 包也存在性能问题（下面会详述）。

所以大多数 Strix Halo 用户只能老老实实用的 Vulkan 后端——能用，但远谈不上好。

三、Lemonade ROCm：官方包反而更慢？

我们分别测试了三个后端在同一台设备（Ryzen AI Max+ 395，128GB RAM）上的性能，使用 Qwen AgentWorld 35B（基于 Qwen 3.6 35B-A3B MoE 的 agent 微调模型）在 Q4_K_XL 量化下：

几个令人震惊的发现：

1. Lemonade ROCm 在大多数场景下比 Vulkan 还慢。在 8K-64K 深度下，Lemonade 的 prefill 全面落后 Vulkan（差异在 5%-12%），decode 也差一截。这跟它使用了通用 ROCm 目标架构（未针对 gfx1151 做 MMQ tile/nwarp 调优）以及 WMMA Flash Attention 路径在 D>128 时有性能退化有关。
2. Strix Halo 特调版在 prefill 上全面碾压：比 Vulkan 快 7%-43%，比 Lemonade 快 20%-25%。这主要归功于多项针对 RDNA 3.5 的 CUDA 内核优化。
3. Decode 速度三者接近，Vulkan 反而略胜。这提示了 UMA（统一内存架构）下 decode 阶段的瓶颈可能不在 GPU 计算本身，而在内存带宽或 kernel launch 开销——Vulkan 的 CPU 端开销更轻，而 ROCm 的 HIP runtime 在统一内存场景下反而有额外损耗。

四、这个 GitHub 项目做了什么

项目地址：justinappler/llama.cpp 的 Strix Halo fork（master 分支），专门针对 gfx1151（RDNA 3.5 iGPU）进行了实验性优化。

核心解决的问题：让 llama.cpp 的 ROCm/HIP 后端在 Strix Halo 上跑出该有的性能，而不是被通用代码路径拖后腿。

从 README 的 findings 表来看，项目经历了多轮「假设→代码→跑分→保留或回滚」的迭代：

总共 4 项保留的优化，2 项失败回滚，2 项放弃或搁置。

五、优化的原理

乍看这些 patch 名字很吓人，但原理可以拆成几个层次理解：

5.1MMQ 矩阵乘法内核调优（Finding #5）

这是最主要的作用。RDNA 3.5（gfx1151）虽然属于 RDNA 3 家族，但 CU 数量只有 40 个（远少于桌面 RDNA 3 的 96 个），L2 Cache 也只有 8MB。通用的 MMQ（矩阵乘-量化）内核把 tile 大小和 warp 数配置设为 RDNA 3 通用的值，对 Strix Halo 来说太大了——寄存器压力过高导致溢出，实际性能打折。

特调版为 gfx1151 定制了 tile 大小（dense 层用 128，MoE expert dispatch 用 48）和 warp 调度策略。也是+27% prefill 速度的关键。

5.2 量化 KV Cache 不要用（Finding #1）

llama.cpp 默认会对 KV Cache 做量化以节省内存，但在 Strix Halo 上（128GB 统一内存），内存根本不是瓶颈，量化反而引入了额外的解量化计算开销。禁用 KV 量化后，长上下文性能直接翻了十几倍——省内存的操作用在不缺内存的设备上是负优化。

5.3 Flash Attention 走 TILE 不走 MMA（Finding #7/#8）

对于 D>128 的长序列注意力计算，通用的 AMD MMA（矩阵乘累加）内核在 RDNA 3.5 上效率不行。上游 PR #22880 最终选择了 TILE 内核来处理 D>128（commit message 里 JohannesGaessler 说：「I was not able to get better performance than the tile kernel for head sizes > 128」）。特调版把 D=256+ 的路由固定到 TILE 内核，额外赢了 +6.3% pp @ 16K。

5.4 淘汰rocWMMA （Finding #6）

rocWMMA 是 AMD 的 WMMA（Wave Matrix Multiply-Accumulate）API，看起来应该更快。但实测结果恰恰相反：落地时性能持平，到后来随着 ROCm 版本更新，反而成了性能杀手（−34% pp @ 16K）。原因：上游 PR #22880 已经让 D>128 走 TILE 内核，rocWMMA 路径在大深度下不再被调用，但残留的 WMMA 调度逻辑仍然产生了开销。特调版直接在生产环境关掉了它（GGML_HIP_ROCWMMA_FATTN=OFF）。

5.5 无论是rocm主分支还是TheRock 分支均适用，TheRock表现会更好一点

项目实际面向两条 ROCm 线：

官方README 里引以为傲的 ~1350 t/s pp、~917 t/s @ 16K 深度，全都是 TheRock 7.13.0a 的数据。我们的实测用的是 HIP SDK 7.1.1（WIndows 上唯一的官方安装器版本），虽然也拿到了不错的提升（+43% vs Vulkan），但跟 TheRock 7.13 的潜在上限之间还有距离。

由于Windows 用户只能用 HIP SDK 7.1.1，而作者的最佳实践跑在 Linux Docker 里的 TheRock 7.13 上。如果你追求极限性能，可能需要考虑 Linux + Docker（可选） + TheRock 路线。

六、这个项目的缺点

1. 没有预编译包，必须自行编译。

这可能是这个项目Star数稀少的最大门槛。因为需要用户自主安装：

• Visual Studio Build Tools 2022/2026（含C++桌面开发环境、 MSVC v143 和 Windows 11 SDK）
• CMake 4.x
• Ninja
• HIP SDK 7.1.1 for Windows

编译过程中还会遇到一堆坑（见第七章）。

2. 已知的 Strix Halo + ROCm 缺陷：

• decode 速度没有提升：从 benchmark 可以看出，tg 速度三者基本持平，Vulkan 甚至更优。这是因为 decode 阶段的计算密度太低（每次只处理一个 token），瓶颈在 kernel launch 延迟和内存带宽，GPU 内核优化对此无效。
• ROCm 版本约束：目前的构建基于 HIP SDK 7.1.1 + clang 21.0。ROCm 7.2+ 虽然修复了 VGPR bug，但 Windows 上的 HIP SDK 还没跟上。未来 ROCm 版本升级后可能需要重新适配。
• 长期 rocFFT 工作负载可能崩溃：HIP SDK 7.1.1 的已知 bug（官方 release notes 提及），在 Strix Halo 上跑长时间 rocFFT 有间歇性崩溃风险。不过 llama.cpp 推理不重度依赖 rocFFT，影响有限。
• 只适配了 gfx1151：项目中的 tile/nwarp 参数是按 40 CU、8MB L2 调优的，如果 AMD 后续推出不同规格的 RDNA 3.5 APU，可能需要重新调参（理论上说，495可能直接用得上）。

3. 代码是实验性的单分支维护：

没有稳定 release tag，只有一个 master 分支随上游 rebase。每次上游有大更新时可能产生冲突。作者的设计哲学是「每个优化一个 commit + 保留文档」，回滚的代码也留在 docs 里——这既是优点（透明可审计），也是缺点（文档散落在 strix-halo/ 目录下的十几个 md 文件里）。

七、编译全流程与踩坑记录

以下是我们在 Windows 11 + Ryzen AI Max+ 395 上成功编译的全部步骤和踩过的坑：

7.1 环境准备

# 1. 安装 Visual Studio Build Tools 2026
#    下载页: https://visualstudio.microsoft.com/downloads/#build-tools-for-visual-studio-2026
#    勾选: ☑ 使用 C++ 的桌面开发
#          ☑ MSVC v143 C++ 生成工具 (最新)
#          ☑ Windows 11 SDK (10.0.26100.0)

# 2. 安装 CMake 和 Ninja
winget install Kitware.CMake       # → CMake 4.3.3
winget install Ninja-build.Ninja   # → Ninja 1.13.2

# 3. 安装 HIP SDK 7.1.1 for Windows
#    下载: https://www.amd.com/en/developer/resources/rocm-hub/hip-sdk.html
#    安装时选择 "Full" 模式，包含 HIP SDK Core + Libraries (Development) + Runtime Compiler

# 4. 将 HIP SDK 加入 PATH
#    管理员 PowerShell:
[Environment]::SetEnvironmentVariable(
    "Path",
    "C:\Program Files\AMD\ROCm\7.1\bin;" + 
    [Environment]::GetEnvironmentVariable("Path", "User"),
    "User"
)

7.2 坑 #1：clang 21 + MSVC STL 数学头文件冲突

error: __device__ function 'isgreater' cannot overload __host__ __device__ function 'isgreater'

原因：VS 2026 自带的 MSVC 14.51 在 <cmath> 中给 isgreater/isless/isunordered 等函数加了 constexpr，clang 21 将其解释为 __host__ __device__，与 HIP SDK 内部头的 __device__ 声明冲突。

解决：CMake 配置时加 -DCMAKE_CXX_FLAGS="-Xclang -fno-cuda-host-device-constexpr"。

7.3 坑 #2：rocwmma-version.hpp 头文件缺失

HIP SDK 7.1.1 for Windows 不含 rocWMMA 头文件。但特调版本来就应该关掉 rocWMMA（见第六章），所以 -DGGML_HIP_ROCWMMA_FATTN=OFF 即可。注意不要像我一样随手复制 README 里的 cmake 命令把 =ON 也抄上了。

7.4 坑 #3：中文用户名路径导致的编码错误

clang-offload-bundler: error: 在多字节的目标代码页中，没有此 Unicode 字符可以映射到的字符。

原因：C:\Users\申凌峰\AppData\Local\Temp\ 路径中的中文让 clang-offload-bundler 在生成 GPU 代码对象时炸了。

解决：编译前 $env:TEMP = "C:\Temp"，用一个纯 ASCII 路径的临时目录。

7.5 坑 #4：GPU device 端不能用 MSVC STL 的 std::begin()

// 原代码：使用 initializer_list
for (int stride_k : {warp_size, warp_size/2, warp_size/4, warp_size/8}) {

MSVC STL 的 std::initializer_list::begin() 被标记为 __host__ only，在 GPU 端调用直接报错。

解决：将 {...}（STL initializer_list）改为 (const int[4]){...}（C 风格复合字面量数组），范围 for 在 C 数组上使用编译器原生指针运算，不需要 STL。

// 修复后
for (int stride_k : (const int[4]){warp_size, warp_size/2, warp_size/4, warp_size/8}) {

只需要改 ggml/src/ggml-cuda/fattn-mma-f16.cuh 中的 2 处。

7.6 最终编译命令

$env:HIP_PATH = "C:\Program Files\AMD\ROCm\7.1"
$env:AMDGPU_TARGETS = "gfx1151"
$env:TEMP = "C:\Temp"
$env:TMP = "C:\Temp"

cmake -S . -B build -G Ninja `
  -DGGML_HIP=ON `
  -DAMDGPU_TARGETS=gfx1151 `
  -DCMAKE_BUILD_TYPE=Release `
  -DGGML_BACKEND_DL=ON `
  -DGGML_CPU_ALL_VARIANTS=ON `
  -DCMAKE_C_FLAGS="-Xclang -fno-cuda-host-device-constexpr" `
  -DCMAKE_CXX_FLAGS="-Xclang -fno-cuda-host-device-constexpr" `
  -DCMAKE_C_COMPILER="C:/Program Files/AMD/ROCm/7.1/bin/clang.exe" `
  -DCMAKE_CXX_COMPILER="C:/Program Files/AMD/ROCm/7.1/bin/clang++.exe"

cmake --build build --config Release -j

八、最终性能对比

回来看我们的三组实测数据（Qwen AgentWorld 35B Q5_K_M，Ryzen AI Max+ 395，128GB LPDDR5x）：

Prefill（Prompt 处理速度，tokens/s，越高越好）

特调版在任何深度下都是最快的。相对于 Vulkan 的优势在短上下文时最明显（+43% @ 4K），随着上下文增长逐渐收敛到 ~7%。相对于 Lemonade 的优势则稳定在 20%-25%——说明 Lemonade 是一个「没有特调的 ROCm 基线」，而特调版的 MMQ tile/nwarp 优化在所有深度下都是稳定收益。

Decode（Token 生成速度，tokens/s，越高越好）

Decode 不是特调版的强项。Vulkan 后端的 decode 速度在三者中最高，特调版和 Lemonade 不相上下。这是因为 agent 场景下的 batch size=1 decode 的计算密度极低，瓶颈完全在 kernel launch 延迟和内存访问，GPU 内核排布优化对此无能为力。这也是项目 README 中把 MMVQ（矩阵-向量量化乘，decode 的主计算路径）优化标为 M 优先级的原因——「Park unless tg becomes the binding constraint」。

如果说35b的表现只是勉强胜出，那么Qwen 3.5 122B 对比数据则趋势更加明显（依然使用Q5_K_M）。

关键发现：122B MoE模型的 prefill 提升幅度（18%–49%）比 35B MoE 模型（7%–43%）更大。这是因为大参数量模型的矩阵计算量更大，MMQ 内核调优的收益更明显。而小参数量模型中 expert dispatch 的稀疏路由已经减少了有效计算量，优化空间相对较小。

Decode 方面，Vulkan 仍然领先 5%–9%，和 35B 上的结论一致。

其实进一步拓展到step3.7 flash，特调版的rocm表现会进一步上升（大约有接近70%的提升），然而考虑到这个模型只能勉强使用Q3缺乏实用价值，因此看到122b的差距即可，文中不再赘述。

细心的读者可能注意到了：Vulkan 后端在 Qwen 35B 的 4K prefill 跑出了 998 t/s，在 122B 的 4K prefill 跑出了 319 t/s——都比它们各自的 8K 深度还低（35B: 1138 t/s，122B: 327 t/s）。按常理，上下文越短 prefill 应该越快才对。
这是首次加载的冷启动效应。benchmark 按深度从小到大依次执行，4K 是第一轮。首次推理时 GPU 需要完成以下一次性开销：
模型权重的 GPU 内存分配与初始化：首次将模型完全加载到 VRAM/UMA 显存区域
Vulkan Shader JIT 编译：Vulkan 后端的计算着色器在第一次调用时触发 GPU 驱动编译，这个开销可达数百毫秒到数秒
KV Cache 首次分配：分配并初始化大块的连续 GPU 内存
第二轮（8K 深度）开始，这些都已完成，所以 8K prefill 的实际计算量虽然更大，但没有了冷启动负担，反而跑出了比 4K 更高的数字。

一句话总结性能：如果你主要做 RAG、长文档分析、代码库检索（prefill-bound），特调版是 No-Brainer 的选择，最高能快 48.6%。如果你只做聊天对话（chatbot），Vulkan 后端就够了，甚至更好。

九、总结

AMD Strix Halo 是一块被软件栈拖累的硬件。官方 ROCm 支持不完善，llama.cpp 的通用代码路径没法充分发挥 RDNA 3.5 的潜力。这个 Strix Halo fork 用工程方法论——一个一个假设地去实验、benchmark、保留或回滚——把 prefill 性能从 Vulkan 的起点推高了 7%–43%。

当然它不完美：decode 没有改善，需要自行编译（踩坑不少），依赖特定的 ROCm 版本，且代码处于实验性单分支维护状态。但对于在 Strix Halo 上跑本地大模型、特别是需要频繁处理长上下文的 agent 编程场景来说，这是目前最好的选择。

如果你也是 Strix Halo 用户，欢迎按第七章的步骤自己编译试试。有任何踩坑经历，欢迎在评论区交流。

数据来源：justinappler/llama.cpp-strix-halo Strix Halo fork，benchmark 在 Ryzen AI Max+ 395 (128GB) 上运行，模型为 Qwen AgentWorld 35B-A3B (Q5_K_M)以及Qwen 3.5 122B-A10B (Q5_K_M)，使用 Feng267/llm_speedtest项目进行测试。