Commit 9145ddb95 for llama.cpp

commit 9145ddb9556c1cfa568f3189eb7be44317826e31
Author: WalterCool <waltercool@slash.cl>
Date:   Fri Oct 23 18:31:28 2026 -0300

    hip: fix bf16 WMMA operand vector type for gfx11/gfx12 (#30885)

    The __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 (and _gfx12) builtins
    take their A/B operands as vectors of short (_ExtVector<16, short> /
    _ExtVector<8, short>), not __bf16. Declaring the fragments as
    ext_vector_type(N) __bf16 fails to compile:

      error: cannot initialize a parameter of type 'short
      __attribute__((ext_vector_type(16)))' with an lvalue of type 'const
      bf16x16_t' (vector of 16 '__bf16' values)

    short and __bf16 are both 16-bit and the value is only ever passed to
    the builtin via reinterpret_cast over raw bytes, so changing the element
    type is bit-identical. The operand type has been short since the builtin
    was introduced, so this works across LLVM/ROCm versions without a
    version gate. Fixes HIP builds for RDNA3 (gfx1150/1151) and RDNA4
    (gfx1200/1201).

    Signed-off-by: Pablo Cholaky <waltercool@slash.cl>

diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh
index 1cec98b06..c72f12131 100644
--- a/ggml/src/ggml-cuda/mma.cuh
+++ b/ggml/src/ggml-cuda/mma.cuh
@@ -1409,14 +1409,14 @@ namespace ggml_cuda_mma {
             tile<16, 16, float, dl_d> & D, const tile<16, 8, nv_bfloat162, dl_ab> & A, const tile<16, 8, nv_bfloat162, dl_ab> & B) {
 #if defined(AMD_WMMA_AVAILABLE)
 #if defined(RDNA4)
-        using bf16x8_t = __attribute__((ext_vector_type(8))) __bf16;
+        using bf16x8_t = __attribute__((ext_vector_type(8))) short;
         using floatx8_t = __attribute__((ext_vector_type(8))) float;
         floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
         const bf16x8_t& a_frag = reinterpret_cast<const bf16x8_t&>(A.x[0]);
         const bf16x8_t& b_frag = reinterpret_cast<const bf16x8_t&>(B.x[0]);
         acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, acc_frag);
 #elif defined(RDNA3)
-        using bf16x16_t = __attribute__((ext_vector_type(16))) __bf16;
+        using bf16x16_t = __attribute__((ext_vector_type(16))) short;
         using floatx8_t = __attribute__((ext_vector_type(8))) float;
         floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
         const bf16x16_t& a_frag = reinterpret_cast<const bf16x16_t&>(A.x[0]);