Fix build with LLVM 23

Since https://github.com/llvm/llvm-project/pull/176033 (LLVM 23), AMDGPU
builtins take ext_vector_type operands, which allow no implicit conversion
between element types. Older clang (e. g. 22 in ROCm 7.x) accepts both.

This is a short version of https://github.com/ggml-org/llama.cpp/pull/29536,
which also adds CDNA fix.
--- a/src/ggml-cuda/mma.cuh
+++ b/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]);
@@ -1430,14 +1430,14 @@ namespace ggml_cuda_mma {
         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;
         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;
             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);