HIP: SOLVE_TRI falls back to rocBLAS strsm which fails on gfx906 (MI50) – suggest extending custom kernel

问题场景

用户运行 llama.cpp 加载使用 Gated Delta Networks / linear attention 架构的模型（如 Qwen3.5、Qwen3-Next、Kimi Linear），这些模型会触发 GGML_OP_SOLVE_TRI 运算。在 AMD gfx906（MI50）显卡上执行时崩溃，日志显示 SOLVE_TRI 运算失败。类似问题也出现在 Radeon 780M（gfx1103）上，且失败与否与 -ngl 参数相关。

报错原文

rocBLAS error from hip error code: 'hipErrorInvalidDeviceFunction':98
ggml_cuda_compute_forward: SOLVE_TRI failed

# RDNA3 iGPU 上的附加错误：
ROCm error: invalid device function
  current device: 0, in function ggml_cuda_compute_forward at ggml/src/ggml-cuda/ggml-cuda.cu:2408

原因分析

在 solve_tri.cu 的调度逻辑中（第 264 行附近）：

if (n <= MAX_N_FAST && k <= MAX_K_FAST) {   // MAX_N_FAST=64, MAX_K_FAST=32
    solve_tri_f32_cuda(...)   // 自定义 CUDA/HIP 内核 - 在 gfx906 上正常
} else {
    solve_tri_f32_cublas(...)  // 调用 cublasStrsmBatched -> rocBLAS strsm -> 在 gfx906 上崩溃
}

自定义 solve_tri_f32_fast 内核在 gfx906 上运行正常（使用标准 CUDA/HIP 指令：warp shuffle、共享内存）。只有在维度超过 n=64 或 k=32 时，才会触发 cublasStrsmBatched（rocBLAS Strsm）路径并崩溃。

可能原因：ROCm 7.x 中的 rocBLAS 没有为 gfx906 预编译 Strsm 的 Tensile 内核。AMD 自 ROCm 5.7 起已弃用 gfx906 支持。即使 Arch Linux 的 rocBLAS 包为 gfx906 构建，也不包含 Strsm 系列内核。注意：Issue 最初称 Qwen3.5-35B-A3B 使用 n=64, k=14（在快速内核限制内），但实际测试确认运行时为 n=64, k=64，k 超过 MAX_K_FAST=32，因此也会走 rocBLAS 路径并崩溃。

环境排查

• 确认 ROCm 版本与显卡架构兼容性：ROCm 7.x 已弃用 gfx906（MI50）支持，建议检查是否有针对该架构的 rocBLAS Strsm 内核
• 确认 GGML_OP_SOLVE_TRI 运算的实际维度（n 和 k 值），可在日志中查找 SOLVE_TRI: n= k= 输出
• 检查 llama.cpp 版本（Issue 中为 8179，构建提交 ecbcb7ea9）
• 多卡环境下确认所有 GPU 均为同一架构（如双 MI50 均为 gfx906）
• 在 Windows + RDNA3 环境中，需额外关注 -ngl 参数与显存分配对路径选择的影响

解决步骤

1. 临时方案（已验证）：编辑 ggml/src/ggml-cuda/ggml-cuda.cu，在 ggml_cuda_supports_op() 函数中注释掉 GGML_OP_SOLVE_TRI：

   case GGML_OP_TRI:
   case GGML_OP_DIAG:
   //case GGML_OP_SOLVE_TRI:  // disabled: rocBLAS Strsm lacks gfx906 kernel
       return true;

2. 重新编译 llama.cpp：cmake --build build --config Release -j$(nproc)
3. 重新运行模型验证是否还触发 rocBLAS 崩溃。
4. 长期修复方向（Issue 建议，未验证）：可优先尝试扩展自定义 solve_tri_f32_fast 内核，使其支持更大的维度（k>32, n>64），完全消除对 rocBLAS 的依赖；或在 ggml_cuda_supports_op() 中做运行时架构检测，对已知缺少 rocBLAS Strsm 支持的架构返回 false，自动回退 CPU；或在 solve_tri_f32_cublas() 中捕获 hipErrorInvalidDeviceFunction 错误并优雅降级。

验证方法

使用触发失败的模型重新运行推理（如 ./build/bin/llama-cli -m Qwen3-Coder-Next-MXFP4.gguf -p "Hello" -n 10），确认不再出现 rocBLAS error from hip error code: 'hipErrorInvalidDeviceFunction':98 或 ggml_cuda_compute_forward: SOLVE_TRI failed。如果采用注释掉 GGML_OP_SOLVE_TRI 的方案，可观察模型成功生成输出，且日志中不再出现 rocBLAS 与 SOLVE_TRI 相关的错误。

参考来源

ggml-org/llama.cpp #19972

AI 工具推荐

想把多个 AI 模型放在一个入口？

GamsGo AI 集成 ChatGPT、DeepSeek、Gemini、Claude、Midjourney、Veo 等常用模型，适合写作、绘图、视频和日常 AI 工作流。

了解 GamsGo AI

推广链接：通过此链接购买，我可能获得佣金，不影响你的价格。