try CI fix

Author: JohannesGaessler

ggml/src/ggml-cuda/mmf.cu

@@ -72,7 +72,11 @@ static __global__ void mul_mat_f(
                     const int j = j0 + itB*tile_B::I;
 
                     const float2 tmp = j < cols_per_block ? y2[j*stride_col_y + col] : make_float2(0.0f, 0.0f);
-                    tile_xy[j0*tile_k_padded + threadIdx.x] = {tmp.x, tmp.y};
+                    if constexpr (std::is_same_v<T, half2>) {
+                        tile_xy[j0*tile_k_padded + threadIdx.x] = make_half2(tmp.x, tmp.y);
+                    } else {
+                        tile_xy[j0*tile_k_padded + threadIdx.x] = make_bfloat162(tmp.x, tmp.y);
+                    }
                 }
             } else {
                 static_assert(std::is_same_v<T, void>, "unsupported type");

ggml/src/ggml-cuda/vendors/hip.h

@@ -199,7 +199,8 @@
     #define __has_builtin(x) 0
 #endif
 
-typedef hip_bfloat16 nv_bfloat16;
+typedef __hip_bfloat16 nv_bfloat16;
+typedef __hip_bfloat162 nv_bfloat162;
 
 typedef int8_t int8x4_t __attribute__((ext_vector_type(4)));
 typedef uint8_t uint8x4_t __attribute__((ext_vector_type(4)));

ggml/src/ggml-cuda/vendors/musa.h

@@ -137,4 +137,5 @@
 #define cudaStreamEndCapture musaStreamEndCapture
 #define cudaOccupancyMaxActiveBlocksPerMultiprocessor musaOccupancyMaxActiveBlocksPerMultiprocessor
 
-typedef mt_bfloat16 nv_bfloat16;
+typedef __mt_bfloat16 nv_bfloat16;
+typedef __mt_bfloat162 nv_bfloat162;