Pass the bf16 operands of the AMD MFMA builtins as short vectors.

Clang declares these builtins with short-vector operands that carry raw
bf16 bits. llama.cpp hands them __bf16 vectors. AMD's ROCm LLVM fork
accepts that by reinterpreting the bits, but stock LLVM 23 and later
reject it, so a HIP build against a stock-LLVM hip fails in mma.cuh.

Upstream fixed the RDNA (WMMA) half in 9145ddb955 ("hip: fix bf16 WMMA
operand vector type for gfx11/gfx12", #30885), which b11906 includes;
this is the CDNA (MFMA) half it left, the second hunk of
llama-cpp-hip-bf16-builtin-args.patch.

https://github.com/istitov/stuff/issues/300

--- a/ggml/src/ggml-cuda/mma.cuh
+++ b/ggml/src/ggml-cuda/mma.cuh
@@ -1430,14 +1430,14 @@
         using floatx4_t = __attribute__((ext_vector_type(4))) float;
         floatx4_t& acc_frag = reinterpret_cast<floatx4_t&>(D.x[0]);
 #if defined(CDNA4) || defined(CDNA3) || defined(CDNA2)
-        using bf16x4_t = __attribute__((ext_vector_type(4))) __bf16;
+        using bf16x4_t = __attribute__((ext_vector_type(4))) short; // bf16 bits
         const bf16x4_t& a_frag = reinterpret_cast<const bf16x4_t&>(A.x[0]);
         const bf16x4_t& b_frag = reinterpret_cast<const bf16x4_t&>(B.x[0]);
         acc_frag = __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(a_frag, b_frag, acc_frag, 0, 0, 0);
 #elif defined(CDNA1)
 #pragma unroll
         for (int i = 0; i < 2; ++i) {
-            using bf16x2_t = __attribute__((ext_vector_type(2))) __bf16;
+            using bf16x2_t = __attribute__((ext_vector_type(2))) short; // bf16 bits
             const bf16x2_t& a_frag = reinterpret_cast<const bf16x2_t&>(A.x[i]);
             const bf16x2_t& b_frag = reinterpret_cast<const bf16x2_t&>(B.x[i]);
             acc_frag = __builtin_amdgcn_mfma_f32_16x16x8bf16(a_frag, b_frag, acc_frag, 0, 0, 0);