llama.cpp b11847修复RDNA3/RDNA4的HIP bf16 WMMA编译问题
该版本将bf16 WMMA片段的元素类型由 __bf16改为short,以匹配AMD内建函数要求的向量类型,修复gfx1150/1151与gfx1200/1201上的HIP构建;贡献者称改动按位等价,且无需按LLVM/ROCm版本设门槛。
AI解读:llama.cpp发布b11847构建版本,唯一列出的变更是修复HIP后端在gfx11/gfx12上bf16 WMMA操作数的向量类型(PR #30885)。
问题出在 __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32及其 _gfx12版本:这两个内建函数要求A/B操作数是short的向量,而代码把片段声明为ext_vector_type(N) __bf16,导致编译报错。
提交说明称short与 __bf16同为 16 位,片段仅通过reinterpret_cast以原始字节传入内建函数,因此改元素类型按位一致。
提交说明称该内建函数自引入起操作数类型即为short,因此改动可跨LLVM/ROCm版本使用,不需要版本判断。
修复范围是RDNA3(gfx1150/1151)和RDNA4(gfx1200/1201)的HIP构建;补丁署名Pablo Cholaky。
发布页同时提供macOS、Linux、Windows、Android等多平台预编译包,其中macOS Apple Silicon的KleidiAI版本和openEuler构建标注为DISABLED。
llama.cpp发布b11847版本,变更日志中列出的唯一一项是HIP后端的编译修复:bf16 WMMA操作数的向量类型改为short,对应PR #30885,署名Pablo Cholaky。
编译失败的原因与改法
提交说明解释,__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(以及 _gfx12版本)的A/B操作数类型是short的向量(_ExtVector),而不是 __bf16。把片段声明成ext_vector_type(N) __bf16会编译失败,报错为无法用const bf16x16_t的左值初始化short ext_vector_type(16) 参数。
提交者称short与 __bf16都是 16 位,片段只通过reinterpret_cast以原始字节传给内建函数,改变元素类型后按位一致;该内建函数的操作数类型自引入以来一直是short,因此不需要按LLVM/ROCm版本加条件。
适用硬件与发布产物
提交说明将修复针对的构建范围限定为RDNA3(gfx1150/1151)与RDNA4(gfx1200/1201)的HIP构建。发布页同时列出macOS、iOS、Linux、Android、Windows以及UI的预编译产物。
- macOS Apple Silicon的KleidiAI版本标注为DISABLED,链接到PR #23780。
- openEuler的x86/aarch64构建(310p、910b ACL Graph)同样标注为DISABLED,链接到PR #23705。
- Linux侧提供ROCm 10.0的Ubuntu x64包,另有CUDA 12.8、CUDA 13.4、Vulkan、OpenVINO、SYCL等构建。