From cbd950b220ed009e6e545be5ef9c1ad0931d6d70 Mon Sep 17 00:00:00 2001 From: Iwan Kawrakow Date: Fri, 23 Feb 2024 16:20:47 +0200 Subject: [PATCH] iq3_s: make it work on metal for QK_K = 64 Pleasent surprise: the coding was super-block size independent, so all it took was to delete some QK_K == 256 guards. --- ggml-metal.metal | 7 ------- 1 file changed, 7 deletions(-) diff --git a/ggml-metal.metal b/ggml-metal.metal index aa954926b..b3bf40539 100644 --- a/ggml-metal.metal +++ b/ggml-metal.metal @@ -4489,7 +4489,6 @@ void kernel_mul_mv_iq3_s_f32_impl( threadgroup_barrier(mem_flags::mem_threadgroup); } -#if QK_K == 256 const int ix = tiisg; device const float * y4 = y + 32 * ix; @@ -4535,12 +4534,6 @@ void kernel_mul_mv_iq3_s_f32_impl( y4 += 32 * 32; } -#else - (void) x; - (void) y; - (void) yl; - (void) nb32; -#endif for (int row = 0; row < N_DST; ++row) { all_sum = simd_sum(sumf[row]);