char-fork/llama.cpp
git clone https://git.t4t.associates/char-fork/llama.cpp
5bda51bfb
char/aliceai-t5-moe
1#include "common.h" 2#include "dequantize.h" 3 4constant bool FC_mul_mm_bc_inp [[function_constant(FC_MUL_MM + 0)]]; 5constant bool FC_mul_mm_bc_out [[function_constant(FC_MUL_MM + 1)]]; 6constant short FC_mul_mm_ne12 [[function_constant(FC_MUL_MM + 2)]]; 7constant short FC_mul_mm_ne13 [[function_constant(FC_MUL_MM + 3)]]; 8constant short FC_mul_mm_r2 [[function_constant(FC_MUL_MM + 4)]]; 9constant short FC_mul_mm_r3 [[function_constant(FC_MUL_MM + 5)]]; 10 11// each block_q contains 16*nl weights 12#ifdef GGML_METAL_HAS_TENSOR 13template< 14 typename SA, typename SA_4x4, typename SA_8x8, 15 typename SB, typename SB_2x4, typename SB_8x8, 16 typename block_q, short nl, void (*dequantize_func)(device const block_q *, short, thread SA_4x4 &), 17 typename T0, typename T0_4x4, typename T1, typename T1_2x4> 18kernel void kernel_mul_mm( 19 constant ggml_metal_kargs_mul_mm & args, 20 device const char * srcA, 21 device const char * srcB, 22 device char * dst, 23 threadgroup char * shmem [[threadgroup(0)]], 24 uint3 tgpig [[threadgroup_position_in_grid]], 25 ushort tiitg [[thread_index_in_threadgroup]], 26 ushort sgitg [[simdgroup_index_in_threadgroup]]) { 27 (void) sgitg; 28 29 // Matrix dimensions: A(M,K) x B(K,N) -> C(M,N) 30 const int K = args.ne00; 31 const int M = args.ne0; 32 const int N = args.ne1; 33 34 // Batch dimension handling 35 const int im = tgpig.z; 36 const int i12 = im % FC_mul_mm_ne12; 37 const int i13 = im / FC_mul_mm_ne12; 38 39 // Batch offsets for srcA and srcB 40 const uint64_t offset0 = (i12/FC_mul_mm_r2)*args.nb02 + (i13/FC_mul_mm_r3)*args.nb03; 41 42 // Tile dimensions 43 constexpr int NRB = SZ_SIMDGROUP * N_MM_BLOCK_X * N_MM_SIMD_GROUP_X; 44 constexpr int NRA = SZ_SIMDGROUP * N_MM_BLOCK_Y * N_MM_SIMD_GROUP_Y; 45 46 // Tile offsets in output matrix 47 const int ra = tgpig.y * NRA; 48 const int rb = tgpig.x * NRB; 49 50 // Threadgroup memory for dequantized A tile only 51 threadgroup SA * sa = (threadgroup SA *)(shmem); 52 53 // Work-item count for A loading 54 constexpr int A_WORK_ITEMS = NRA * N_MM_NK; 55 constexpr int NUM_THREADS = N_SIMDWIDTH * N_MM_SIMD_GROUP_X * N_MM_SIMD_GROUP_Y; 56 57 // tA wraps threadgroup memory 58 auto tA = tensor(sa, dextents<int32_t, 2>(N_MM_NK_TOTAL, NRA)); 59 60 // tB wraps device memory directly 61 device T1 * ptrB = (device T1 *)(srcB + args.nb12*i12 + args.nb13*i13); 62 const int strideB = args.nb11 / sizeof(T1); 63 auto tB = tensor(ptrB, dextents<int32_t, 2>(K, N), array<int, 2>({1, strideB})); 64 65 // Configure matmul operation 66 // note: K is dynamic_extent (clamped to the valid range in PHASE 2), since a static 67 // N_MM_NK_TOTAL K tile would read src1 out of bounds when K % N_MM_NK_TOTAL != 0 68 // ref: https://github.com/ggml-org/llama.cpp/pull/27064 69 mpp::tensor_ops::matmul2d< 70 mpp::tensor_ops::matmul2d_descriptor( 71 NRB, NRA, static_cast<int>(dynamic_extent), false, true, true, 72 mpp::tensor_ops::matmul2d_descriptor::mode::multiply_accumulate), 73 execution_simdgroups<N_MM_SIMD_GROUP_X * N_MM_SIMD_GROUP_Y>> mm; 74 75 auto cT = mm.get_destination_cooperative_tensor<decltype(tB), decltype(tA), float>(); 76 77 // Accumulate partial results over K dimension 78 for (int loop_k = 0; loop_k < K; loop_k += N_MM_NK_TOTAL) { 79 // === PHASE 1: Dequantization of A into threadgroup memory === 80 for (int work = tiitg; work < A_WORK_ITEMS; work += NUM_THREADS) { 81 const int row = work / N_MM_NK; 82 const int k_chunk = work % N_MM_NK; 83 const int k_pos = loop_k + k_chunk * 16; 84 const short k_base = k_chunk * 16; 85 86 // Bounds check: skip device read if row is out of matrix bounds 87 if (ra + row < M) { 88 if (is_same<T0_4x4, block_q>::value && FC_mul_mm_bc_inp) { 89 // Element-wise reads when K is not aligned (nb01 not aligned for half4x4/float4x4). 90 // MSL spec Table 2.5: half4x4 requires 8-byte alignment. When K is odd, 91 // nb01 = K*2 is not 8-byte aligned, so odd-row pointers are misaligned. 92 // Mirrors the legacy kernel's existing guard. 93 device const T0 * row_ptr = (device const T0 *)(srcA + args.nb01 * (ra + row) + offset0); 94 95 FOR_UNROLL (short i = 0; i < 16; i++) { 96 sa[row * N_MM_NK_TOTAL + (k_base + i)] = (k_pos + i < K) ? (SA) row_ptr[k_pos + i] : (SA)0; 97 } 98 } else { 99 const int block_idx = k_pos / (16 * nl); 100 const short il = (k_pos / 16) % nl; 101 102 device const block_q * row_ptr = (device const block_q *)(srcA + args.nb01 * (ra + row) + offset0); 103 104 SA_4x4 temp_a; 105 dequantize_func(row_ptr + block_idx, il, temp_a); 106 107 FOR_UNROLL (short i = 0; i < 16; i++) { 108 // Zero-pad A for K positions beyond valid range (handles partial K iterations) 109 sa[row * N_MM_NK_TOTAL + (k_base + i)] = (k_pos + i < K) ? temp_a[i/4][i%4] : (SA)0; 110 } 111 } 112 } else { 113 // Zero-pad rows beyond matrix bounds 114 FOR_UNROLL (short i = 0; i < 16; i++) { 115 sa[row * N_MM_NK_TOTAL + (k_base + i)] = (SA)0; 116 } 117 } 118 } 119 120 threadgroup_barrier(mem_flags::mem_threadgroup); 121 122 // === PHASE 2: Tensor matmul === 123 // Clamp the K extent of both operand tensors to the remaining valid K range so 124 // the dynamic-K op never reads past the K extent of src1 (or the staged A tile). 125 const int kExt = min(N_MM_NK_TOTAL, K - loop_k); 126 127 auto tAv = tensor(sa, dextents<int32_t, 2>(kExt, NRA), array<int, 2>({1, N_MM_NK_TOTAL})); 128 auto tBv = tensor(ptrB + loop_k + rb * strideB, dextents<int32_t, 2>(kExt, N - rb), array<int, 2>({1, strideB})); 129 130 mm.run(tBv, tAv, cT); 131 132 threadgroup_barrier(mem_flags::mem_threadgroup); 133 } 134 135 // Store result tile to output matrix (with batch offset) 136 // cT.store handles bounds checking via tD's extents (M, N) 137 device float * dstBatch = (device float *)dst + im * N * M; 138 139 auto tD = tensor(dstBatch, dextents<int32_t, 2>(M, N), array<int, 2>({1, M})); 140 cT.store(tD.slice(ra, rb)); 141} 142 143#else 144 145template< 146 typename S0, typename S0_4x4, typename S0_8x8, 147 typename S1, typename S1_2x4, typename S1_8x8, 148 typename block_q, short nl, void (*dequantize_func)(device const block_q *, short, thread S0_4x4 &), 149 typename T0, typename T0_4x4, typename T1, typename T1_2x4> 150kernel void kernel_mul_mm( 151 constant ggml_metal_kargs_mul_mm & args, 152 device const char * src0, 153 device const char * src1, 154 device char * dst, 155 threadgroup char * shmem [[threadgroup(0)]], 156 uint3 tgpig[[threadgroup_position_in_grid]], 157 ushort tiitg[[thread_index_in_threadgroup]], 158 ushort sgitg[[simdgroup_index_in_threadgroup]]) { 159 160 threadgroup S0 * sa = (threadgroup S0 *)(shmem); 161 threadgroup S1 * sb = (threadgroup S1 *)(shmem + 4096); 162 163 constexpr int NR0 = 64; 164 constexpr int NR1 = 32; 165 166 constexpr int NK = 32; 167 constexpr int NL0 = NK/16; 168 constexpr int NL1 = NK/8; 169 170 const int im = tgpig.z; 171 const int r0 = tgpig.y*NR0; 172 const int r1 = tgpig.x*NR1; 173 174 // if this block is of 64x32 shape or smaller 175 const short nr0 = (args.ne0 - r0 < NR0) ? (args.ne0 - r0) : NR0; 176 const short nr1 = (args.ne1 - r1 < NR1) ? (args.ne1 - r1) : NR1; 177 178 // a thread shouldn't load data outside of the matrix 179 const short lr0 = ((short)tiitg/NL0) < nr0 ? ((short)tiitg/NL0) : nr0 - 1; // 0 .. 63 180 const short lr1 = ((short)tiitg/NL1) < nr1 ? ((short)tiitg/NL1) : nr1 - 1; // 0 .. 31 181 182 const short il0 = (tiitg % NL0); 183 184 short il = il0; 185 186 const int i12 = im % FC_mul_mm_ne12; 187 const int i13 = im / FC_mul_mm_ne12; 188 189 const uint64_t offset0 = (i12/FC_mul_mm_r2)*args.nb02 + (i13/FC_mul_mm_r3)*args.nb03; 190 const short offset1 = il0/nl; 191 192 device const block_q * x = (device const block_q *)(src0 + args.nb01*(r0 + lr0) + offset0) + offset1; 193 194 const short iy = 8*(tiitg % NL1); 195 196 device const T1 * y = (device const T1 *)(src1 197 + args.nb13*i13 198 + args.nb12*i12 199 + args.nb11*(r1 + lr1) 200 + args.nb10*iy); 201 202 S0_8x8 ma[4]; 203 S1_8x8 mb[2]; 204 205 simdgroup_float8x8 mc[8]; 206 207 for (short i = 0; i < 8; i++){ 208 mc[i] = make_filled_simdgroup_matrix<float, 8>(0.f); 209 } 210 211 for (int loop_k = 0; loop_k < args.ne00; loop_k += NK) { 212 // load data and store to threadgroup memory 213 if (is_same<T0_4x4, block_q>::value && FC_mul_mm_bc_inp) { 214 threadgroup_barrier(mem_flags::mem_threadgroup); 215 216 // no need for dequantization 217 for (short i = 0; i < 16; i++) { 218 const short sx = 2*il0 + i/8; 219 const short sy = (tiitg/NL0)/8; 220 221 //const short lx = i%8; 222 //const short ly = (tiitg/NL0)%8; 223 const short lx = (tiitg/NL0)%8; 224 const short ly = i%8; 225 226 const short ib = 8*sx + sy; 227 228 *(sa + 64*ib + 8*ly + lx) = loop_k + 16*il + i < args.ne00 ? *((device T0 *) x + i) : 0; 229 } 230 } else { 231 S0_4x4 temp_a; 232 dequantize_func(x, il, temp_a); 233 234 threadgroup_barrier(mem_flags::mem_threadgroup); 235 236 FOR_UNROLL (short i = 0; i < 16; i++) { 237 const short sx = 2*il0 + i/8; 238 const short sy = (tiitg/NL0)/8; 239 240 //const short lx = i%8; 241 //const short ly = (tiitg/NL0)%8; 242 const short lx = (tiitg/NL0)%8; 243 const short ly = i%8; 244 245 const short ib = 8*sx + sy; 246 247 // NOTE: this is massively slower.. WTF? 248 //sa[64*ib + 8*ly + lx] = temp_a[i/4][i%4]; 249 250 *(sa + 64*ib + 8*ly + lx) = temp_a[i/4][i%4]; 251 } 252 } 253 254 if (FC_mul_mm_bc_inp) { 255 for (short i = 0; i < 8; ++i) { 256 const short sx = (tiitg%NL1); 257 const short sy = (tiitg/NL1)/8; 258 259 const short lx = i; 260 const short ly = (tiitg/NL1)%8; 261 //const short lx = (tiitg/NL1)%8; 262 //const short ly = i; 263 264 const short ib = 4*sx + sy; 265 266 *(sb + 64*ib + 8*ly + lx) = loop_k + iy + i < args.ne00 ? (S1) *((device T1 *) y + i) : 0; 267 } 268 } else { 269 const short sx = (tiitg%NL1); 270 const short sy = (tiitg/NL1)/8; 271 272 //const short dx = sx; 273 //const short dy = sy; 274 275 const short ly = (tiitg/NL1)%8; 276 277 const short ib = 4*sx + sy; 278 279 *(threadgroup S1_2x4 *)(sb + 64*ib + 8*ly) = (S1_2x4)(*((device T1_2x4 *) y)); 280 } 281 282 il = (il + 2 < nl) ? il + 2 : il % 2; 283 x = (il < 2) ? x + (2 + nl - 1)/nl : x; 284 285 y += NK; 286 287 threadgroup_barrier(mem_flags::mem_threadgroup); 288 289 // load matrices from threadgroup memory and conduct outer products 290 threadgroup const S0 * lsma = (sa + 4*64*(sgitg%2)); 291 threadgroup const S1 * lsmb = (sb + 2*64*(sgitg/2)); 292 293 FOR_UNROLL (short ik = 0; ik < NK/8; ik++) { 294 simdgroup_barrier(mem_flags::mem_none); 295 296 FOR_UNROLL (short i = 0; i < 4; i++) { 297 simdgroup_load(ma[i], lsma + 64*i, 8, 0, false); 298 } 299 300 simdgroup_barrier(mem_flags::mem_none); 301 302 FOR_UNROLL (short i = 0; i < 2; i++) { 303 simdgroup_load(mb[i], lsmb + 64*i, 8, 0, false); 304 } 305 306 simdgroup_barrier(mem_flags::mem_none); 307 308 FOR_UNROLL (short i = 0; i < 8; i++){ 309 simdgroup_multiply_accumulate(mc[i], mb[i/4], ma[i%4], mc[i]); 310 } 311 312 lsma += 8*64; 313 lsmb += 4*64; 314 } 315 } 316 317 if (!FC_mul_mm_bc_out || (r0 + NR0 <= args.ne0 && r1 + NR1 <= args.ne1)) { 318 // if no bounds checks on the output are needed, we can directly write to device memory 319 device float * C = (device float *) dst + 320 (r0 + 32*(sgitg & 1)) + \ 321 (r1 + 16*(sgitg >> 1)) * args.ne0 + im*args.ne1*args.ne0; 322 323 for (short i = 0; i < 8; i++) { 324 simdgroup_store(mc[i], C + 8*(i%4) + 8*args.ne0*(i/4), args.ne0, 0, false); 325 } 326 } else { 327 // block is smaller than 64x32, we should avoid writing data outside of the matrix 328 threadgroup_barrier(mem_flags::mem_threadgroup); 329 330 threadgroup float * temp_str = ((threadgroup float *) shmem) + 32*(sgitg&1) + (16*(sgitg >> 1))*NR0; 331 332 for (short i = 0; i < 8; i++) { 333 simdgroup_store(mc[i], temp_str + 8*(i%4) + 8*NR0*(i/4), NR0, 0, false); 334 } 335 336 threadgroup_barrier(mem_flags::mem_threadgroup); 337 338 if (sgitg == 0) { 339 for (int j = tiitg; j < nr1; j += NR1) { 340 device float * D = (device float *) dst + r0 + (r1 + j)*args.ne0 + im*args.ne1*args.ne0; 341 device float4 * D4 = (device float4 *) D; 342 343 threadgroup float * C = temp_str + (j*NR0); 344 threadgroup float4 * C4 = (threadgroup float4 *) C; 345 346 int i = 0; 347 for (; i < nr0/4; i++) { 348 *(D4 + i) = *(C4 + i); 349 } 350 351 i *= 4; 352 for (; i < nr0; i++) { 353 *(D + i) = *(C + i); 354 } 355 } 356 } 357 } 358} 359 360#endif // GGML_METAL_HAS_TENSOR 361 362template<short ne20> // n_expert_used 363kernel void kernel_mul_mm_id_map0( 364 constant ggml_metal_kargs_mul_mm_id_map0 & args, 365 device const char * src2, 366 device char * htpe, 367 device char * hids, 368 threadgroup char * shmem [[threadgroup(0)]], 369 ushort tpitg[[thread_position_in_threadgroup]], 370 ushort ntg[[threads_per_threadgroup]]) { 371 const short ide = tpitg; // expert id 372 373 uint32_t n_all = 0; 374 375 device int32_t * ids_i32 = (device int32_t *) hids + ide*args.ne21; 376 377 for (int i21 = 0; i21 < args.ne21; i21 += ntg) { // n_tokens 378 if (i21 + tpitg < args.ne21) { 379 device const int32_t * src2_i32 = (device const int32_t *) (src2 + (i21 + tpitg)*args.nb21); 380 381 threadgroup uint16_t * sids = (threadgroup uint16_t *) shmem + tpitg*ne20; 382 383 #pragma unroll(ne20) 384 for (short i20 = 0; i20 < ne20; i20++) { 385 sids[i20] = src2_i32[i20]; 386 } 387 } 388 389 threadgroup_barrier(mem_flags::mem_threadgroup); 390 391 for (short t = 0; t < ntg; t++) { 392 if (i21 + t >= args.ne21) { 393 break; 394 } 395 396 threadgroup const uint16_t * sids = (threadgroup const uint16_t *) shmem + t*ne20; 397 398 short sel = 0; 399 #pragma unroll(ne20) 400 for (short i20 = 0; i20 < ne20; i20++) { 401 sel += (sids[i20] == ide)*(i20 + 1); 402 } 403 404 ids_i32[n_all] = (i21 + t)*ne20 + sel - 1; 405 406 n_all += sel > 0; 407 } 408 409 threadgroup_barrier(mem_flags::mem_threadgroup); 410 } 411 412 device uint32_t * tpe_u32 = (device uint32_t *) (htpe); 413 tpe_u32[ide] = n_all; 414} 415 416typedef decltype(kernel_mul_mm_id_map0<1>) kernel_mul_mm_id_map0_t; 417 418template [[host_name("kernel_mul_mm_id_map0_ne20_1" )]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<1>; 419template [[host_name("kernel_mul_mm_id_map0_ne20_2" )]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<2>; 420template [[host_name("kernel_mul_mm_id_map0_ne20_4" )]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<4>; 421template [[host_name("kernel_mul_mm_id_map0_ne20_5" )]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<5>; 422template [[host_name("kernel_mul_mm_id_map0_ne20_6" )]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<6>; 423template [[host_name("kernel_mul_mm_id_map0_ne20_8" )]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<8>; 424template [[host_name("kernel_mul_mm_id_map0_ne20_10")]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<10>; 425template [[host_name("kernel_mul_mm_id_map0_ne20_16")]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<16>; 426template [[host_name("kernel_mul_mm_id_map0_ne20_22")]] kernel kernel_mul_mm_id_map0_t kernel_mul_mm_id_map0<22>; 427 428template<typename S0, typename S0_4x4, typename S0_8x8, typename S1, typename S1_2x4, typename S1_8x8, typename block_q, short nl, void (*dequantize_func)(device const block_q *, short, thread S0_4x4 &), typename T0, typename T0_4x4, typename T1, typename T1_2x4> 429kernel void kernel_mul_mm_id( 430 constant ggml_metal_kargs_mul_mm_id & args, 431 device const char * src0, 432 device const char * src1, 433 device const char * htpe, 434 device const char * hids, 435 device char * dst, 436 threadgroup char * shmem [[threadgroup(0)]], 437 uint3 tgpig[[threadgroup_position_in_grid]], 438 ushort tiitg[[thread_index_in_threadgroup]], 439 ushort tiisg[[thread_index_in_simdgroup]], 440 ushort sgitg[[simdgroup_index_in_threadgroup]]) { 441 threadgroup S0 * sa = (threadgroup S0 *)(shmem); 442 threadgroup S1 * sb = (threadgroup S1 *)(shmem + 4096); 443 444#ifdef GGML_METAL_HAS_TENSOR 445 threadgroup float * sc = (threadgroup float *)(shmem); 446#endif 447 448 constexpr int NR0 = 64; 449 constexpr int NR1 = 32; 450 451 constexpr int NK = 32; 452 constexpr int NL0 = NK/16; 453 constexpr int NL1 = NK/8; 454 455 const int im = tgpig.z; // expert 456 const int r0 = tgpig.y*NR0; 457 const int r1 = tgpig.x*NR1; 458 459 device const uint32_t * tpe_u32 = (device const uint32_t *) (htpe); 460 device const int32_t * ids_i32 = (device const int32_t *) (hids); 461 462 const int32_t neh1 = tpe_u32[im]; 463 464 if (r1 >= neh1) { 465 return; 466 } 467 468 // if this block is of 64x32 shape or smaller 469 const short nr0 = (args.ne0 - r0 < NR0) ? (args.ne0 - r0) : NR0; 470 const short nr1 = ( neh1 - r1 < NR1) ? ( neh1 - r1) : NR1; 471 472 // a thread shouldn't load data outside of the matrix 473 const short lr0 = ((short)tiitg/NL0) < nr0 ? ((short)tiitg/NL0) : nr0 - 1; // 0 .. 63 474 const short lr1 = ((short)tiitg/NL1) < nr1 ? ((short)tiitg/NL1) : nr1 - 1; // 0 .. 31 475 476 const short il0 = (tiitg % NL0); 477 478 short il = il0; 479 480 const int id = ids_i32[im*args.ne21 + r1 + lr1]; 481 482 const short i11 = (id % args.ne20) % args.ne11; 483 const short i12 = (id / args.ne20); 484 const short i13 = 0; 485 486 const uint64_t offset0 = im*args.nb02 + i13*args.nb03; 487 const short offset1 = il0/nl; 488 489 device const block_q * x = (device const block_q *)(src0 + args.nb01*(r0 + lr0) + offset0) + offset1; 490 491 const short iy = 8*(tiitg % NL1); 492 493 device const T1 * y = (device const T1 *)(src1 494 + args.nb13*i13 495 + args.nb12*i12 496 + args.nb11*i11 497 + args.nb10*iy); 498 499 // skip the upper half of the token tile when the expert did not fill it 500 constexpr short NR1H = NR1/2; 501 502 const bool has_hi = nr1 > NR1H; 503 504 const short lb1 = (short) tiitg/NL1; // 0 .. NR1-1, this thread's row of the B tile 505 506#ifndef GGML_METAL_HAS_TENSOR 507 S0_8x8 ma[4]; 508 S1_8x8 mb[2]; 509 510 simdgroup_float8x8 mc[8]; 511 512 for (short i = 0; i < 8; i++){ 513 mc[i] = make_filled_simdgroup_matrix<float, 8>(0.f); 514 } 515 516 // simdgroups 2,3 own rows NR1H..NR1-1 517 const bool sg_active = has_hi || sgitg < 2; 518#else 519 auto tA = tensor<threadgroup S0, dextents<int32_t, 2>, tensor_inline>(sa, dextents<int32_t, 2>(NK, NR0)); 520 521 // sb is [NR1][NK] row-major 522 auto tB0 = tensor<threadgroup S1, dextents<int32_t, 2>, tensor_inline>(sb, dextents<int32_t, 2>(NK, NR1H)); 523 auto tB1 = tensor<threadgroup S1, dextents<int32_t, 2>, tensor_inline>(sb + NR1H*NK, dextents<int32_t, 2>(NK, NR1H)); 524 525 mpp::tensor_ops::matmul2d< 526 mpp::tensor_ops::matmul2d_descriptor(NR1H, NR0, NK, false, true, false, mpp::tensor_ops::matmul2d_descriptor::mode::multiply_accumulate), 527 execution_simdgroups<4>> mm; 528 529 auto cT0 = mm.get_destination_cooperative_tensor<decltype(tA), decltype(tB0), float>(); 530 auto cT1 = mm.get_destination_cooperative_tensor<decltype(tA), decltype(tB1), float>(); 531#endif 532 533 for (int loop_k = 0; loop_k < args.ne00; loop_k += NK) { 534#ifndef GGML_METAL_HAS_TENSOR 535 // load data and store to threadgroup memory 536 if (is_same<T0_4x4, block_q>::value && FC_mul_mm_bc_inp) { 537 threadgroup_barrier(mem_flags::mem_threadgroup); 538 539 // no need for dequantization 540 for (short i = 0; i < 16; i++) { 541 const short sx = 2*il0 + i/8; 542 const short sy = (tiitg/NL0)/8; 543 544 //const short lx = i%8; 545 //const short ly = (tiitg/NL0)%8; 546 const short lx = (tiitg/NL0)%8; 547 const short ly = i%8; 548 549 const short ib = 8*sx + sy; 550 551 *(sa + 64*ib + 8*ly + lx) = loop_k + 16*il + i < args.ne00 ? (S0) *((device T0 *) x + i) : (S0) 0; 552 } 553 } else { 554 S0_4x4 temp_a; 555 dequantize_func(x, il, temp_a); 556 557 threadgroup_barrier(mem_flags::mem_threadgroup); 558 559 FOR_UNROLL (short i = 0; i < 16; i++) { 560 const short sx = 2*il0 + i/8; 561 const short sy = (tiitg/NL0)/8; 562 563 //const short lx = i%8; 564 //const short ly = (tiitg/NL0)%8; 565 const short lx = (tiitg/NL0)%8; 566 const short ly = i%8; 567 568 const short ib = 8*sx + sy; 569 570 // NOTE: this is massively slower.. WTF? 571 //sa[64*ib + 8*ly + lx] = temp_a[i/4][i%4]; 572 573 *(sa + 64*ib + 8*ly + lx) = temp_a[i/4][i%4]; 574 } 575 } 576 577 if (FC_mul_mm_bc_inp) { 578 for (short i = 0; i < 8; ++i) { 579 const short sx = (tiitg%NL1); 580 const short sy = (tiitg/NL1)/8; 581 582 const short lx = i; 583 const short ly = (tiitg/NL1)%8; 584 //const short lx = (tiitg/NL1)%8; 585 //const short ly = i; 586 587 const short ib = 4*sx + sy; 588 589 *(sb + 64*ib + 8*ly + lx) = loop_k + iy + i < args.ne00 ? (S1) *((device T1 *) y + i) : 0; 590 } 591 } else { 592 const short sx = (tiitg%NL1); 593 const short sy = (tiitg/NL1)/8; 594 595 //const short dx = sx; 596 //const short dy = sy; 597 598 const short ly = (tiitg/NL1)%8; 599 600 const short ib = 4*sx + sy; 601 602 *(threadgroup S1_2x4 *)(sb + 64*ib + 8*ly) = (S1_2x4)(*((device T1_2x4 *) y)); 603 } 604#else 605 // load data and store to threadgroup memory 606 if (is_same<T0_4x4, block_q>::value && FC_mul_mm_bc_inp) { 607 threadgroup_barrier(mem_flags::mem_threadgroup); 608 609 // no need for dequantization 610 for (short i = 0; i < 16; i++) { 611 const short sx = 2*il0 + i/8; 612 const short sy = (tiitg/NL0)/8; 613 614 const short lx = i%8; 615 const short ly = (tiitg/NL0)%8; 616 //const short lx = (tiitg/NL0)%8; 617 //const short ly = i%8; 618 619 *(sa + NK*(8*sy + ly) + 8*sx + lx) = loop_k + 16*il + i < args.ne00 ? *((device T0 *) x + i) : 0; 620 } 621 } else { 622 S0_4x4 temp_a; 623 dequantize_func(x, il, temp_a); 624 625 threadgroup_barrier(mem_flags::mem_threadgroup); 626 627 FOR_UNROLL (short i = 0; i < 16; i++) { 628 const short sx = 2*il0 + i/8; 629 const short sy = (tiitg/NL0)/8; 630 631 const short lx = i%8; 632 const short ly = (tiitg/NL0)%8; 633 //const short lx = (tiitg/NL0)%8; 634 //const short ly = i%8; 635 636 *(sa + NK*(8*sy + ly) + 8*sx + lx) = temp_a[i/4][i%4]; 637 } 638 } 639 640 if (FC_mul_mm_bc_inp) { 641 for (short i = 0; i < 8; ++i) { 642 const short sx = (tiitg%NL1); 643 const short sy = (tiitg/NL1)/8; 644 645 const short lx = i; 646 const short ly = (tiitg/NL1)%8; 647 //const short lx = (tiitg/NL1)%8; 648 //const short ly = i; 649 650 *(sb + NK*(8*sy + ly) + 8*sx + lx) = loop_k + iy + i < args.ne00 ? (S1) *((device T1 *) y + i) : 0; 651 } 652 } else { 653 const short sx = (tiitg%NL1); 654 const short sy = (tiitg/NL1)/8; 655 656 //const short lx = i; 657 const short ly = (tiitg/NL1)%8; 658 //const short lx = (tiitg/NL1)%8; 659 //const short ly = i; 660 661 *(threadgroup S1_2x4 *)(sb + NK*(8*sy + ly) + 8*sx) = (S1_2x4)(*((device T1_2x4 *) y)); 662 } 663#endif 664 665 il = (il + 2 < nl) ? il + 2 : il % 2; 666 x = (il < 2) ? x + (2 + nl - 1)/nl : x; 667 668 y += NK; 669 670 threadgroup_barrier(mem_flags::mem_threadgroup); 671 672#ifndef GGML_METAL_HAS_TENSOR 673 if (sg_active) { 674 // load matrices from threadgroup memory and conduct outer products 675 threadgroup const S0 * lsma = (sa + 4*64*(sgitg%2)); 676 threadgroup const S1 * lsmb = (sb + 2*64*(sgitg/2)); 677 678 FOR_UNROLL (short ik = 0; ik < NK/8; ik++) { 679 simdgroup_barrier(mem_flags::mem_none); 680 681 FOR_UNROLL (short i = 0; i < 4; i++) { 682 simdgroup_load(ma[i], lsma + 64*i, 8, 0, false); 683 } 684 685 simdgroup_barrier(mem_flags::mem_none); 686 687 FOR_UNROLL (short i = 0; i < 2; i++) { 688 simdgroup_load(mb[i], lsmb + 64*i, 8, 0, false); 689 } 690 691 simdgroup_barrier(mem_flags::mem_none); 692 693 FOR_UNROLL (short i = 0; i < 8; i++){ 694 simdgroup_multiply_accumulate(mc[i], mb[i/4], ma[i%4], mc[i]); 695 } 696 697 lsma += 8*64; 698 lsmb += 4*64; 699 } 700 } 701#else 702 auto sA = tA.slice(0, 0); 703 auto sB0 = tB0.slice(0, 0); 704 705 mm.run(sB0, sA, cT0); 706 707 if (has_hi) { 708 auto sB1 = tB1.slice(0, 0); 709 710 mm.run(sB1, sA, cT1); 711 } 712#endif 713 } 714 715 // block is smaller than 64x32, we should avoid writing data outside of the matrix 716 threadgroup_barrier(mem_flags::mem_threadgroup); 717 718#ifdef GGML_METAL_HAS_TENSOR 719 auto tC0 = tensor<threadgroup float, dextents<int32_t, 2>, tensor_inline>(sc, dextents<int32_t, 2>(NR0, NR1H)); 720 cT0.store(tC0); 721 722 if (has_hi) { 723 auto tC1 = tensor<threadgroup float, dextents<int32_t, 2>, tensor_inline>(sc + NR1H*NR0, dextents<int32_t, 2>(NR0, NR1H)); 724 cT1.store(tC1); 725 } 726#else 727 if (sg_active) { 728 threadgroup float * temp_str = ((threadgroup float *) shmem) + 32*(sgitg&1) + (16*(sgitg >> 1))*NR0; 729 730 for (short i = 0; i < 8; i++) { 731 simdgroup_store(mc[i], temp_str + 8*(i%4) + 8*NR0*(i/4), NR0, 0, false); 732 } 733 } 734#endif 735 736 threadgroup_barrier(mem_flags::mem_threadgroup); 737 738 for (short j = sgitg; j < nr1; j += 4) { 739 const int id = ids_i32[im*args.ne21 + r1 + j]; 740 741 const short ide = id % args.ne20; 742 const short idt = id / args.ne20; 743 744 device float * D = (device float *) dst + r0 + ide*args.ne0 + idt*args.ne1*args.ne0; 745 device float4 * D4 = (device float4 *) D; 746 747 threadgroup float * C = (threadgroup float *) shmem + j*NR0; 748 threadgroup float4 * C4 = (threadgroup float4 *) C; 749 750 int i = tiisg; 751 for (; i < nr0/4; i += 32) { 752 *(D4 + i) = *(C4 + i); 753 } 754 755 i = (4*(nr0/4)) + tiisg; 756 for (; i < nr0; i += 32) { 757 *(D + i) = *(C + i); 758 } 759 } 760} 761 762// 763// matrix-matrix multiplication 764// 765 766typedef decltype(kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, float4x4, 1, dequantize_f32, float, float4x4, float, float2x4>) mul_mm_t; 767 768template [[host_name("kernel_mul_mm_f32_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, float4x4, 1, dequantize_f32, float, float4x4, float, float2x4>; 769template [[host_name("kernel_mul_mm_f16_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, half4x4, 1, dequantize_f16, half, half4x4, float, float2x4>; 770#if defined(GGML_METAL_HAS_BF16) 771template [[host_name("kernel_mul_mm_bf16_f32")]] kernel mul_mm_t kernel_mul_mm<bfloat, bfloat4x4, simdgroup_bfloat8x8, bfloat, bfloat2x4, simdgroup_bfloat8x8, bfloat4x4, 1, dequantize_bf16, bfloat, bfloat4x4, float, float2x4>; 772#endif 773template [[host_name("kernel_mul_mm_q1_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q1_0, 8, dequantize_q1_0, float, float4x4, float, float2x4>; 774template [[host_name("kernel_mul_mm_q2_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_0, 4, dequantize_q2_0, float, float4x4, float, float2x4>; 775template [[host_name("kernel_mul_mm_q4_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_0, 2, dequantize_q4_0, float, float4x4, float, float2x4>; 776template [[host_name("kernel_mul_mm_q4_1_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_1, 2, dequantize_q4_1, float, float4x4, float, float2x4>; 777template [[host_name("kernel_mul_mm_q5_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_0, 2, dequantize_q5_0, float, float4x4, float, float2x4>; 778template [[host_name("kernel_mul_mm_q5_1_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_1, 2, dequantize_q5_1, float, float4x4, float, float2x4>; 779template [[host_name("kernel_mul_mm_q8_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q8_0, 2, dequantize_q8_0, float, float4x4, float, float2x4>; 780template [[host_name("kernel_mul_mm_mxfp4_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, float, float2x4>; 781template [[host_name("kernel_mul_mm_q2_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, float, float2x4>; 782template [[host_name("kernel_mul_mm_q3_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, float, float2x4>; 783template [[host_name("kernel_mul_mm_q4_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, float, float2x4>; 784template [[host_name("kernel_mul_mm_q5_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_K, QK_NL, dequantize_q5_K, float, float4x4, float, float2x4>; 785template [[host_name("kernel_mul_mm_q6_K_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q6_K, QK_NL, dequantize_q6_K, float, float4x4, float, float2x4>; 786template [[host_name("kernel_mul_mm_iq2_xxs_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xxs, QK_NL, dequantize_iq2_xxs, float, float4x4, float, float2x4>; 787template [[host_name("kernel_mul_mm_iq2_xs_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xs, QK_NL, dequantize_iq2_xs, float, float4x4, float, float2x4>; 788template [[host_name("kernel_mul_mm_iq3_xxs_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_xxs, QK_NL, dequantize_iq3_xxs, float, float4x4, float, float2x4>; 789template [[host_name("kernel_mul_mm_iq3_s_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_s, QK_NL, dequantize_iq3_s, float, float4x4, float, float2x4>; 790template [[host_name("kernel_mul_mm_iq2_s_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_s, QK_NL, dequantize_iq2_s, float, float4x4, float, float2x4>; 791template [[host_name("kernel_mul_mm_iq1_s_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_s, QK_NL, dequantize_iq1_s, float, float4x4, float, float2x4>; 792template [[host_name("kernel_mul_mm_iq1_m_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_m, QK_NL, dequantize_iq1_m, float, float4x4, float, float2x4>; 793template [[host_name("kernel_mul_mm_iq4_nl_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_nl, 2, dequantize_iq4_nl, float, float4x4, float, float2x4>; 794template [[host_name("kernel_mul_mm_iq4_xs_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_xs, QK_NL, dequantize_iq4_xs, float, float4x4, float, float2x4>; 795template [[host_name("kernel_mul_mm_tq2_0_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_tq2_0, QK_NL, dequantize_tq2_0, float, float4x4, float, float2x4>; 796 797template [[host_name("kernel_mul_mm_f32_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, float4x4, 1, dequantize_f32, float, float4x4, half, half2x4>; 798template [[host_name("kernel_mul_mm_f16_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, half4x4, 1, dequantize_f16, half, half4x4, half, half2x4>; 799template [[host_name("kernel_mul_mm_q1_0_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q1_0, 8, dequantize_q1_0, float, float4x4, half, half2x4>; 800template [[host_name("kernel_mul_mm_q2_0_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_0, 4, dequantize_q2_0, float, float4x4, half, half2x4>; 801template [[host_name("kernel_mul_mm_q4_0_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_0, 2, dequantize_q4_0, float, float4x4, half, half2x4>; 802template [[host_name("kernel_mul_mm_q4_1_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_1, 2, dequantize_q4_1, float, float4x4, half, half2x4>; 803template [[host_name("kernel_mul_mm_q5_0_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_0, 2, dequantize_q5_0, float, float4x4, half, half2x4>; 804template [[host_name("kernel_mul_mm_q5_1_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_1, 2, dequantize_q5_1, float, float4x4, half, half2x4>; 805template [[host_name("kernel_mul_mm_q8_0_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q8_0, 2, dequantize_q8_0, float, float4x4, half, half2x4>; 806template [[host_name("kernel_mul_mm_mxfp4_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, half, half2x4>; 807template [[host_name("kernel_mul_mm_q2_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, half, half2x4>; 808template [[host_name("kernel_mul_mm_q3_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, half, half2x4>; 809template [[host_name("kernel_mul_mm_q4_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, half, half2x4>; 810template [[host_name("kernel_mul_mm_q5_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_K, QK_NL, dequantize_q5_K, float, float4x4, half, half2x4>; 811template [[host_name("kernel_mul_mm_q6_K_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q6_K, QK_NL, dequantize_q6_K, float, float4x4, half, half2x4>; 812template [[host_name("kernel_mul_mm_iq2_xxs_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xxs, QK_NL, dequantize_iq2_xxs, float, float4x4, half, half2x4>; 813template [[host_name("kernel_mul_mm_iq2_xs_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xs, QK_NL, dequantize_iq2_xs, float, float4x4, half, half2x4>; 814template [[host_name("kernel_mul_mm_iq3_xxs_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_xxs, QK_NL, dequantize_iq3_xxs, float, float4x4, half, half2x4>; 815template [[host_name("kernel_mul_mm_iq3_s_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_s, QK_NL, dequantize_iq3_s, float, float4x4, half, half2x4>; 816template [[host_name("kernel_mul_mm_iq2_s_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_s, QK_NL, dequantize_iq2_s, float, float4x4, half, half2x4>; 817template [[host_name("kernel_mul_mm_iq1_s_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_s, QK_NL, dequantize_iq1_s, float, float4x4, half, half2x4>; 818template [[host_name("kernel_mul_mm_iq1_m_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_m, QK_NL, dequantize_iq1_m, float, float4x4, half, half2x4>; 819template [[host_name("kernel_mul_mm_iq4_nl_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_nl, 2, dequantize_iq4_nl, float, float4x4, half, half2x4>; 820template [[host_name("kernel_mul_mm_iq4_xs_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_xs, QK_NL, dequantize_iq4_xs, float, float4x4, half, half2x4>; 821template [[host_name("kernel_mul_mm_tq2_0_f16")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_tq2_0, QK_NL, dequantize_tq2_0, float, float4x4, half, half2x4>; 822 823// 824// indirect matrix-matrix multiplication 825// 826 827typedef decltype(kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, float4x4, 1, dequantize_f32, float, float4x4, float, float2x4>) mul_mm_id; 828 829template [[host_name("kernel_mul_mm_id_f32_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, float4x4, 1, dequantize_f32, float, float4x4, float, float2x4>; 830template [[host_name("kernel_mul_mm_id_f16_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, half4x4, 1, dequantize_f16, half, half4x4, float, float2x4>; 831#if defined(GGML_METAL_HAS_BF16) 832template [[host_name("kernel_mul_mm_id_bf16_f32")]] kernel mul_mm_id kernel_mul_mm_id<bfloat, bfloat4x4, simdgroup_bfloat8x8, bfloat, bfloat2x4, simdgroup_bfloat8x8, bfloat4x4, 1, dequantize_bf16, bfloat, bfloat4x4, float, float2x4>; 833#endif 834template [[host_name("kernel_mul_mm_id_q1_0_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q1_0, 8, dequantize_q1_0, float, float4x4, float, float2x4>; 835template [[host_name("kernel_mul_mm_id_q2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_0, 4, dequantize_q2_0, float, float4x4, float, float2x4>; 836template [[host_name("kernel_mul_mm_id_q4_0_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_0, 2, dequantize_q4_0, float, float4x4, float, float2x4>; 837template [[host_name("kernel_mul_mm_id_q4_1_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_1, 2, dequantize_q4_1, float, float4x4, float, float2x4>; 838template [[host_name("kernel_mul_mm_id_q5_0_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_0, 2, dequantize_q5_0, float, float4x4, float, float2x4>; 839template [[host_name("kernel_mul_mm_id_q5_1_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_1, 2, dequantize_q5_1, float, float4x4, float, float2x4>; 840template [[host_name("kernel_mul_mm_id_q8_0_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q8_0, 2, dequantize_q8_0, float, float4x4, float, float2x4>; 841template [[host_name("kernel_mul_mm_id_mxfp4_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, float, float2x4>; 842template [[host_name("kernel_mul_mm_id_q2_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, float, float2x4>; 843template [[host_name("kernel_mul_mm_id_q3_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, float, float2x4>; 844template [[host_name("kernel_mul_mm_id_q4_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, float, float2x4>; 845template [[host_name("kernel_mul_mm_id_q5_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_K, QK_NL, dequantize_q5_K, float, float4x4, float, float2x4>; 846template [[host_name("kernel_mul_mm_id_q6_K_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q6_K, QK_NL, dequantize_q6_K, float, float4x4, float, float2x4>; 847template [[host_name("kernel_mul_mm_id_iq2_xxs_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xxs, QK_NL, dequantize_iq2_xxs, float, float4x4, float, float2x4>; 848template [[host_name("kernel_mul_mm_id_iq2_xs_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xs, QK_NL, dequantize_iq2_xs, float, float4x4, float, float2x4>; 849template [[host_name("kernel_mul_mm_id_iq3_xxs_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_xxs, QK_NL, dequantize_iq3_xxs, float, float4x4, float, float2x4>; 850template [[host_name("kernel_mul_mm_id_iq3_s_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_s, QK_NL, dequantize_iq3_s, float, float4x4, float, float2x4>; 851template [[host_name("kernel_mul_mm_id_iq2_s_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_s, QK_NL, dequantize_iq2_s, float, float4x4, float, float2x4>; 852template [[host_name("kernel_mul_mm_id_iq1_s_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_s, QK_NL, dequantize_iq1_s, float, float4x4, float, float2x4>; 853template [[host_name("kernel_mul_mm_id_iq1_m_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_m, QK_NL, dequantize_iq1_m, float, float4x4, float, float2x4>; 854template [[host_name("kernel_mul_mm_id_iq4_nl_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_nl, 2, dequantize_iq4_nl, float, float4x4, float, float2x4>; 855template [[host_name("kernel_mul_mm_id_iq4_xs_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_xs, QK_NL, dequantize_iq4_xs, float, float4x4, float, float2x4>; 856template [[host_name("kernel_mul_mm_id_tq2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_tq2_0, QK_NL, dequantize_tq2_0, float, float4x4, float, float2x4>; 857 858template [[host_name("kernel_mul_mm_id_f32_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, float4x4, 1, dequantize_f32, float, float4x4, half, half2x4>; 859template [[host_name("kernel_mul_mm_id_f16_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, half4x4, 1, dequantize_f16, half, half4x4, half, half2x4>; 860template [[host_name("kernel_mul_mm_id_q1_0_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q1_0, 8, dequantize_q1_0, float, float4x4, half, half2x4>; 861template [[host_name("kernel_mul_mm_id_q2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_0, 4, dequantize_q2_0, float, float4x4, half, half2x4>; 862template [[host_name("kernel_mul_mm_id_q4_0_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_0, 2, dequantize_q4_0, float, float4x4, half, half2x4>; 863template [[host_name("kernel_mul_mm_id_q4_1_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_1, 2, dequantize_q4_1, float, float4x4, half, half2x4>; 864template [[host_name("kernel_mul_mm_id_q5_0_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_0, 2, dequantize_q5_0, float, float4x4, half, half2x4>; 865template [[host_name("kernel_mul_mm_id_q5_1_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_1, 2, dequantize_q5_1, float, float4x4, half, half2x4>; 866template [[host_name("kernel_mul_mm_id_q8_0_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q8_0, 2, dequantize_q8_0, float, float4x4, half, half2x4>; 867template [[host_name("kernel_mul_mm_id_mxfp4_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_mxfp4, 2, dequantize_mxfp4, float, float4x4, half, half2x4>; 868template [[host_name("kernel_mul_mm_id_q2_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q2_K, QK_NL, dequantize_q2_K, float, float4x4, half, half2x4>; 869template [[host_name("kernel_mul_mm_id_q3_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q3_K, QK_NL, dequantize_q3_K, float, float4x4, half, half2x4>; 870template [[host_name("kernel_mul_mm_id_q4_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q4_K, QK_NL, dequantize_q4_K, float, float4x4, half, half2x4>; 871template [[host_name("kernel_mul_mm_id_q5_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q5_K, QK_NL, dequantize_q5_K, float, float4x4, half, half2x4>; 872template [[host_name("kernel_mul_mm_id_q6_K_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_q6_K, QK_NL, dequantize_q6_K, float, float4x4, half, half2x4>; 873template [[host_name("kernel_mul_mm_id_iq2_xxs_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xxs, QK_NL, dequantize_iq2_xxs, float, float4x4, half, half2x4>; 874template [[host_name("kernel_mul_mm_id_iq2_xs_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_xs, QK_NL, dequantize_iq2_xs, float, float4x4, half, half2x4>; 875template [[host_name("kernel_mul_mm_id_iq3_xxs_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_xxs, QK_NL, dequantize_iq3_xxs, float, float4x4, half, half2x4>; 876template [[host_name("kernel_mul_mm_id_iq3_s_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq3_s, QK_NL, dequantize_iq3_s, float, float4x4, half, half2x4>; 877template [[host_name("kernel_mul_mm_id_iq2_s_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq2_s, QK_NL, dequantize_iq2_s, float, float4x4, half, half2x4>; 878template [[host_name("kernel_mul_mm_id_iq1_s_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_s, QK_NL, dequantize_iq1_s, float, float4x4, half, half2x4>; 879template [[host_name("kernel_mul_mm_id_iq1_m_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq1_m, QK_NL, dequantize_iq1_m, float, float4x4, half, half2x4>; 880template [[host_name("kernel_mul_mm_id_iq4_nl_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_nl, 2, dequantize_iq4_nl, float, float4x4, half, half2x4>; 881template [[host_name("kernel_mul_mm_id_iq4_xs_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_iq4_xs, QK_NL, dequantize_iq4_xs, float, float4x4, half, half2x4>; 882template [[host_name("kernel_mul_mm_id_tq2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, block_tq2_0, QK_NL, dequantize_tq2_0, float, float4x4, half, half2x4>;