metal : rename all_sum -> sum_all

ggerganov · ggerganov · commit 24a9ea8b44a8 · 2025-03-25T14:34:38.000+02:00
ggml-ci
diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal
@@ -2394,9 +2394,9 @@ void kernel_mul_mv_impl(
                 sumf += (T0) x[i] * (T1) y[i];
             }
 
-            float all_sum = simd_sum(sumf);
+            float sum_all = simd_sum(sumf);
             if (tiisg == 0) {
-                dst_f32[(uint64_t)r1*args.ne0 + r0] = all_sum;
+                dst_f32[(uint64_t)r1*args.ne0 + r0] = sum_all;
             }
         }
     } else {
@@ -2417,10 +2417,10 @@ void kernel_mul_mv_impl(
                 sumf += dot((float4) x4[i], (float4) y4[i]);
             }
 
-            float all_sum = simd_sum(sumf);
+            float sum_all = simd_sum(sumf);
             if (tiisg == 0) {
-                for (int i = 4*(args.ne00/4); i < args.ne00; ++i) all_sum += (float) (x[i] * y[i]);
-                dst_f32[(uint64_t)r1*args.ne0 + r0] = all_sum;
+                for (int i = 4*(args.ne00/4); i < args.ne00; ++i) sum_all += (float) (x[i] * y[i]);
+                dst_f32[(uint64_t)r1*args.ne0 + r0] = sum_all;
             }
         }
     }
@@ -2482,9 +2482,9 @@ kernel void kernel_mul_mv_1row(
         for (int i = tiisg; i < args.ne00; i += 32) {
             sumf += (float) x[i] * (float) y[i];
         }
-        float all_sum = simd_sum(sumf);
+        float sum_all = simd_sum(sumf);
         if (tiisg == 0) {
-            dst_f32[r0] = all_sum;
+            dst_f32[r0] = sum_all;
         }
     } else {
         device const T4     * x4 = (device const T4     *) x;
@@ -2494,11 +2494,11 @@ kernel void kernel_mul_mv_1row(
             sumf += dot((float4) x4[i], y4[i]);
         }
 
-        float all_sum = simd_sum(sumf);
+        float sum_all = simd_sum(sumf);
 
         if (tiisg == 0) {
-            for (int i = 4*(args.ne00/4); i < args.ne00; ++i) all_sum += (float) (x[i] * y[i]);
-            dst_f32[r0] = all_sum;
+            for (int i = 4*(args.ne00/4); i < args.ne00; ++i) sum_all += (float) (x[i] * y[i]);
+            dst_f32[r0] = sum_all;
         }
     }
 }
@@ -2543,9 +2543,9 @@ kernel void kernel_mul_mv_l4(
             sumf += dot((float4) x4[i], y4[i]);
         }
 
-        float all_sum = simd_sum(sumf);
+        float sum_all = simd_sum(sumf);
         if (tiisg == 0) {
-            dst_f32[(uint64_t)r1*args.ne0 + r0] = all_sum;
+            dst_f32[(uint64_t)r1*args.ne0 + r0] = sum_all;
         }
     }
 }
@@ -4447,7 +4447,7 @@ void kernel_mul_mv_q2_K_f32_impl(
     device const float      * y = (device const float      *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int ix = tiisg/8;  // 0...3
     const int it = tiisg%8;  // 0...7
@@ -4503,9 +4503,9 @@ void kernel_mul_mv_q2_K_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }
@@ -4727,7 +4727,7 @@ void kernel_mul_mv_q4_K_f32_impl(
     float yl[16];
     float yh[16];
 
-    float sumf[N_R0_Q4_K]={0.f}, all_sum;
+    float sumf[N_R0_Q4_K]={0.f};
 
     device const float * y4 = y + ix * QK_K + 64 * iq + 8 * ir;
 
@@ -4793,9 +4793,9 @@ void kernel_mul_mv_q4_K_f32_impl(
     device float * dst_f32 = (device float *) dst + (int64_t)im*args.ne0*args.ne1 + (int64_t)r1*args.ne0;
 
     for (int row = 0; row < N_R0_Q4_K && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }
@@ -4981,7 +4981,6 @@ void kernel_mul_mv_q6_K_f32_impl(
     // TODO: support nr0 > 1
     static_assert(nr0 == 1, "nr0 > 1 not supported");
     float sumf[1] = { 0.f };
-    float all_sum;
 
     const short tid = tiisg/2;
     const short ix  = tiisg%2;
@@ -5020,9 +5019,9 @@ void kernel_mul_mv_q6_K_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }
@@ -5070,7 +5069,7 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
     device const float         * y = (device const float         *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int nb32 = nb * (QK_K / 32);
 
@@ -5130,9 +5129,9 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum * 0.25f;
+            dst_f32[first_row + row] = sum_all * 0.25f;
         }
     }
 }
@@ -5178,7 +5177,7 @@ void kernel_mul_mv_iq2_xs_f32_impl(
     device const float        * y = (device const float        *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int nb32 = nb * (QK_K / 32);
 
@@ -5248,9 +5247,9 @@ void kernel_mul_mv_iq2_xs_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum * 0.25f;
+            dst_f32[first_row + row] = sum_all * 0.25f;
         }
     }
 }
@@ -5297,7 +5296,7 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
     device const float         * y = (device const float         *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int nb32 = nb * (QK_K / 32);
 
@@ -5358,9 +5357,9 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum * 0.5f;
+            dst_f32[first_row + row] = sum_all * 0.5f;
         }
     }
 }
@@ -5407,7 +5406,7 @@ void kernel_mul_mv_iq3_s_f32_impl(
     device const float       * y = (device const float       *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int nb32 = nb * (QK_K / 32);
 
@@ -5470,9 +5469,9 @@ void kernel_mul_mv_iq3_s_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }
@@ -5519,7 +5518,7 @@ void kernel_mul_mv_iq2_s_f32_impl(
     device const float       * y = (device const float       *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int nb32 = nb * (QK_K / 32);
 
@@ -5583,9 +5582,9 @@ void kernel_mul_mv_iq2_s_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum * 0.25f;
+            dst_f32[first_row + row] = sum_all * 0.25f;
         }
     }
 }
@@ -5632,7 +5631,7 @@ void kernel_mul_mv_iq1_s_f32_impl(
     device const float       * y = (device const float       *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int nb32 = nb * (QK_K / 32);
 
@@ -5683,9 +5682,9 @@ void kernel_mul_mv_iq1_s_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }
@@ -5732,7 +5731,7 @@ void kernel_mul_mv_iq1_m_f32_impl(
     device const float       * y = (device const float       *) (src1 + offset1);
 
     float yl[32];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     const int nb32 = nb * (QK_K / 32);
 
@@ -5792,9 +5791,9 @@ void kernel_mul_mv_iq1_m_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }
@@ -5848,7 +5847,7 @@ void kernel_mul_mv_iq4_nl_f32_impl(
     threadgroup_barrier(mem_flags::mem_threadgroup);
 
     float4 yl[4];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     device const float * yb = y + ix * QK4_NL + it * 8;
 
@@ -5897,9 +5896,9 @@ void kernel_mul_mv_iq4_nl_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }
@@ -5954,7 +5953,7 @@ void kernel_mul_mv_iq4_xs_f32_impl(
     threadgroup_barrier(mem_flags::mem_threadgroup);
 
     float4 yl[4];
-    float sumf[nr0]={0.f}, all_sum;
+    float sumf[nr0]={0.f};
 
     device const float * yb = y + ix * QK_K + ib * 32 + il * 8;
 
@@ -6000,9 +5999,9 @@ void kernel_mul_mv_iq4_xs_f32_impl(
     device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
 
     for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {
-        all_sum = simd_sum(sumf[row]);
+        float sum_all = simd_sum(sumf[row]);
         if (tiisg == 0) {
-            dst_f32[first_row + row] = all_sum;
+            dst_f32[first_row + row] = sum_all;
         }
     }
 }

Original file line number	Diff line number	Diff line change
`@@ -2394,9 +2394,9 @@ void kernel_mul_mv_impl(`
`2394`	`2394`	`sumf += (T0) x[i] * (T1) y[i];`
`2395`	`2395`	`}`
`2396`	`2396`
`2397`		`- float all_sum = simd_sum(sumf);`
	`2397`	`+ float sum_all = simd_sum(sumf);`
`2398`	`2398`	`if (tiisg == 0) {`
`2399`		`- dst_f32[(uint64_t)r1*args.ne0 + r0] = all_sum;`
	`2399`	`+ dst_f32[(uint64_t)r1*args.ne0 + r0] = sum_all;`
`2400`	`2400`	`}`
`2401`	`2401`	`}`
`2402`	`2402`	`} else {`
`@@ -2417,10 +2417,10 @@ void kernel_mul_mv_impl(`
`2417`	`2417`	`sumf += dot((float4) x4[i], (float4) y4[i]);`
`2418`	`2418`	`}`
`2419`	`2419`
`2420`		`- float all_sum = simd_sum(sumf);`
	`2420`	`+ float sum_all = simd_sum(sumf);`
`2421`	`2421`	`if (tiisg == 0) {`
`2422`		`- for (int i = 4(args.ne00/4); i < args.ne00; ++i) all_sum += (float) (x[i] y[i]);`
`2423`		`- dst_f32[(uint64_t)r1*args.ne0 + r0] = all_sum;`
	`2422`	`+ for (int i = 4(args.ne00/4); i < args.ne00; ++i) sum_all += (float) (x[i] y[i]);`
	`2423`	`+ dst_f32[(uint64_t)r1*args.ne0 + r0] = sum_all;`
`2424`	`2424`	`}`
`2425`	`2425`	`}`
`2426`	`2426`	`}`
`@@ -2482,9 +2482,9 @@ kernel void kernel_mul_mv_1row(`
`2482`	`2482`	`for (int i = tiisg; i < args.ne00; i += 32) {`
`2483`	`2483`	`sumf += (float) x[i] * (float) y[i];`
`2484`	`2484`	`}`
`2485`		`- float all_sum = simd_sum(sumf);`
	`2485`	`+ float sum_all = simd_sum(sumf);`
`2486`	`2486`	`if (tiisg == 0) {`
`2487`		`- dst_f32[r0] = all_sum;`
	`2487`	`+ dst_f32[r0] = sum_all;`
`2488`	`2488`	`}`
`2489`	`2489`	`} else {`
`2490`	`2490`	`device const T4 * x4 = (device const T4 *) x;`
`@@ -2494,11 +2494,11 @@ kernel void kernel_mul_mv_1row(`
`2494`	`2494`	`sumf += dot((float4) x4[i], y4[i]);`
`2495`	`2495`	`}`
`2496`	`2496`
`2497`		`- float all_sum = simd_sum(sumf);`
	`2497`	`+ float sum_all = simd_sum(sumf);`
`2498`	`2498`
`2499`	`2499`	`if (tiisg == 0) {`
`2500`		`- for (int i = 4(args.ne00/4); i < args.ne00; ++i) all_sum += (float) (x[i] y[i]);`
`2501`		`- dst_f32[r0] = all_sum;`
	`2500`	`+ for (int i = 4(args.ne00/4); i < args.ne00; ++i) sum_all += (float) (x[i] y[i]);`
	`2501`	`+ dst_f32[r0] = sum_all;`
`2502`	`2502`	`}`
`2503`	`2503`	`}`
`2504`	`2504`	`}`
`@@ -2543,9 +2543,9 @@ kernel void kernel_mul_mv_l4(`
`2543`	`2543`	`sumf += dot((float4) x4[i], y4[i]);`
`2544`	`2544`	`}`
`2545`	`2545`
`2546`		`- float all_sum = simd_sum(sumf);`
	`2546`	`+ float sum_all = simd_sum(sumf);`
`2547`	`2547`	`if (tiisg == 0) {`
`2548`		`- dst_f32[(uint64_t)r1*args.ne0 + r0] = all_sum;`
	`2548`	`+ dst_f32[(uint64_t)r1*args.ne0 + r0] = sum_all;`
`2549`	`2549`	`}`
`2550`	`2550`	`}`
`2551`	`2551`	`}`
`@@ -4447,7 +4447,7 @@ void kernel_mul_mv_q2_K_f32_impl(`
`4447`	`4447`	`device const float * y = (device const float *) (src1 + offset1);`
`4448`	`4448`
`4449`	`4449`	`float yl[32];`
`4450`		`- float sumf[nr0]={0.f}, all_sum;`
	`4450`	`+ float sumf[nr0]={0.f};`
`4451`	`4451`
`4452`	`4452`	`const int ix = tiisg/8; // 0...3`
`4453`	`4453`	`const int it = tiisg%8; // 0...7`
`@@ -4503,9 +4503,9 @@ void kernel_mul_mv_q2_K_f32_impl(`
`4503`	`4503`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`4504`	`4504`
`4505`	`4505`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`4506`		`- all_sum = simd_sum(sumf[row]);`
	`4506`	`+ float sum_all = simd_sum(sumf[row]);`
`4507`	`4507`	`if (tiisg == 0) {`
`4508`		`- dst_f32[first_row + row] = all_sum;`
	`4508`	`+ dst_f32[first_row + row] = sum_all;`
`4509`	`4509`	`}`
`4510`	`4510`	`}`
`4511`	`4511`	`}`
`@@ -4727,7 +4727,7 @@ void kernel_mul_mv_q4_K_f32_impl(`
`4727`	`4727`	`float yl[16];`
`4728`	`4728`	`float yh[16];`
`4729`	`4729`
`4730`		`- float sumf[N_R0_Q4_K]={0.f}, all_sum;`
	`4730`	`+ float sumf[N_R0_Q4_K]={0.f};`
`4731`	`4731`
`4732`	`4732`	`device const float * y4 = y + ix * QK_K + 64 * iq + 8 * ir;`
`4733`	`4733`
`@@ -4793,9 +4793,9 @@ void kernel_mul_mv_q4_K_f32_impl(`
`4793`	`4793`	`device float * dst_f32 = (device float ) dst + (int64_t)imargs.ne0args.ne1 + (int64_t)r1args.ne0;`
`4794`	`4794`
`4795`	`4795`	`for (int row = 0; row < N_R0_Q4_K && first_row + row < args.ne0; ++row) {`
`4796`		`- all_sum = simd_sum(sumf[row]);`
	`4796`	`+ float sum_all = simd_sum(sumf[row]);`
`4797`	`4797`	`if (tiisg == 0) {`
`4798`		`- dst_f32[first_row + row] = all_sum;`
	`4798`	`+ dst_f32[first_row + row] = sum_all;`
`4799`	`4799`	`}`
`4800`	`4800`	`}`
`4801`	`4801`	`}`
`@@ -4981,7 +4981,6 @@ void kernel_mul_mv_q6_K_f32_impl(`
`4981`	`4981`	`// TODO: support nr0 > 1`
`4982`	`4982`	`static_assert(nr0 == 1, "nr0 > 1 not supported");`
`4983`	`4983`	`float sumf[1] = { 0.f };`
`4984`		`- float all_sum;`
`4985`	`4984`
`4986`	`4985`	`const short tid = tiisg/2;`
`4987`	`4986`	`const short ix = tiisg%2;`
`@@ -5020,9 +5019,9 @@ void kernel_mul_mv_q6_K_f32_impl(`
`5020`	`5019`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5021`	`5020`
`5022`	`5021`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5023`		`- all_sum = simd_sum(sumf[row]);`
	`5022`	`+ float sum_all = simd_sum(sumf[row]);`
`5024`	`5023`	`if (tiisg == 0) {`
`5025`		`- dst_f32[first_row + row] = all_sum;`
	`5024`	`+ dst_f32[first_row + row] = sum_all;`
`5026`	`5025`	`}`
`5027`	`5026`	`}`
`5028`	`5027`	`}`
`@@ -5070,7 +5069,7 @@ void kernel_mul_mv_iq2_xxs_f32_impl(`
`5070`	`5069`	`device const float * y = (device const float *) (src1 + offset1);`
`5071`	`5070`
`5072`	`5071`	`float yl[32];`
`5073`		`- float sumf[nr0]={0.f}, all_sum;`
	`5072`	`+ float sumf[nr0]={0.f};`
`5074`	`5073`
`5075`	`5074`	`const int nb32 = nb * (QK_K / 32);`
`5076`	`5075`
`@@ -5130,9 +5129,9 @@ void kernel_mul_mv_iq2_xxs_f32_impl(`
`5130`	`5129`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5131`	`5130`
`5132`	`5131`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5133`		`- all_sum = simd_sum(sumf[row]);`
	`5132`	`+ float sum_all = simd_sum(sumf[row]);`
`5134`	`5133`	`if (tiisg == 0) {`
`5135`		`- dst_f32[first_row + row] = all_sum * 0.25f;`
	`5134`	`+ dst_f32[first_row + row] = sum_all * 0.25f;`
`5136`	`5135`	`}`
`5137`	`5136`	`}`
`5138`	`5137`	`}`
`@@ -5178,7 +5177,7 @@ void kernel_mul_mv_iq2_xs_f32_impl(`
`5178`	`5177`	`device const float * y = (device const float *) (src1 + offset1);`
`5179`	`5178`
`5180`	`5179`	`float yl[32];`
`5181`		`- float sumf[nr0]={0.f}, all_sum;`
	`5180`	`+ float sumf[nr0]={0.f};`
`5182`	`5181`
`5183`	`5182`	`const int nb32 = nb * (QK_K / 32);`
`5184`	`5183`
`@@ -5248,9 +5247,9 @@ void kernel_mul_mv_iq2_xs_f32_impl(`
`5248`	`5247`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5249`	`5248`
`5250`	`5249`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5251`		`- all_sum = simd_sum(sumf[row]);`
	`5250`	`+ float sum_all = simd_sum(sumf[row]);`
`5252`	`5251`	`if (tiisg == 0) {`
`5253`		`- dst_f32[first_row + row] = all_sum * 0.25f;`
	`5252`	`+ dst_f32[first_row + row] = sum_all * 0.25f;`
`5254`	`5253`	`}`
`5255`	`5254`	`}`
`5256`	`5255`	`}`
`@@ -5297,7 +5296,7 @@ void kernel_mul_mv_iq3_xxs_f32_impl(`
`5297`	`5296`	`device const float * y = (device const float *) (src1 + offset1);`
`5298`	`5297`
`5299`	`5298`	`float yl[32];`
`5300`		`- float sumf[nr0]={0.f}, all_sum;`
	`5299`	`+ float sumf[nr0]={0.f};`
`5301`	`5300`
`5302`	`5301`	`const int nb32 = nb * (QK_K / 32);`
`5303`	`5302`
`@@ -5358,9 +5357,9 @@ void kernel_mul_mv_iq3_xxs_f32_impl(`
`5358`	`5357`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5359`	`5358`
`5360`	`5359`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5361`		`- all_sum = simd_sum(sumf[row]);`
	`5360`	`+ float sum_all = simd_sum(sumf[row]);`
`5362`	`5361`	`if (tiisg == 0) {`
`5363`		`- dst_f32[first_row + row] = all_sum * 0.5f;`
	`5362`	`+ dst_f32[first_row + row] = sum_all * 0.5f;`
`5364`	`5363`	`}`
`5365`	`5364`	`}`
`5366`	`5365`	`}`
`@@ -5407,7 +5406,7 @@ void kernel_mul_mv_iq3_s_f32_impl(`
`5407`	`5406`	`device const float * y = (device const float *) (src1 + offset1);`
`5408`	`5407`
`5409`	`5408`	`float yl[32];`
`5410`		`- float sumf[nr0]={0.f}, all_sum;`
	`5409`	`+ float sumf[nr0]={0.f};`
`5411`	`5410`
`5412`	`5411`	`const int nb32 = nb * (QK_K / 32);`
`5413`	`5412`
`@@ -5470,9 +5469,9 @@ void kernel_mul_mv_iq3_s_f32_impl(`
`5470`	`5469`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5471`	`5470`
`5472`	`5471`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5473`		`- all_sum = simd_sum(sumf[row]);`
	`5472`	`+ float sum_all = simd_sum(sumf[row]);`
`5474`	`5473`	`if (tiisg == 0) {`
`5475`		`- dst_f32[first_row + row] = all_sum;`
	`5474`	`+ dst_f32[first_row + row] = sum_all;`
`5476`	`5475`	`}`
`5477`	`5476`	`}`
`5478`	`5477`	`}`
`@@ -5519,7 +5518,7 @@ void kernel_mul_mv_iq2_s_f32_impl(`
`5519`	`5518`	`device const float * y = (device const float *) (src1 + offset1);`
`5520`	`5519`
`5521`	`5520`	`float yl[32];`
`5522`		`- float sumf[nr0]={0.f}, all_sum;`
	`5521`	`+ float sumf[nr0]={0.f};`
`5523`	`5522`
`5524`	`5523`	`const int nb32 = nb * (QK_K / 32);`
`5525`	`5524`
`@@ -5583,9 +5582,9 @@ void kernel_mul_mv_iq2_s_f32_impl(`
`5583`	`5582`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5584`	`5583`
`5585`	`5584`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5586`		`- all_sum = simd_sum(sumf[row]);`
	`5585`	`+ float sum_all = simd_sum(sumf[row]);`
`5587`	`5586`	`if (tiisg == 0) {`
`5588`		`- dst_f32[first_row + row] = all_sum * 0.25f;`
	`5587`	`+ dst_f32[first_row + row] = sum_all * 0.25f;`
`5589`	`5588`	`}`
`5590`	`5589`	`}`
`5591`	`5590`	`}`
`@@ -5632,7 +5631,7 @@ void kernel_mul_mv_iq1_s_f32_impl(`
`5632`	`5631`	`device const float * y = (device const float *) (src1 + offset1);`
`5633`	`5632`
`5634`	`5633`	`float yl[32];`
`5635`		`- float sumf[nr0]={0.f}, all_sum;`
	`5634`	`+ float sumf[nr0]={0.f};`
`5636`	`5635`
`5637`	`5636`	`const int nb32 = nb * (QK_K / 32);`
`5638`	`5637`
`@@ -5683,9 +5682,9 @@ void kernel_mul_mv_iq1_s_f32_impl(`
`5683`	`5682`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5684`	`5683`
`5685`	`5684`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5686`		`- all_sum = simd_sum(sumf[row]);`
	`5685`	`+ float sum_all = simd_sum(sumf[row]);`
`5687`	`5686`	`if (tiisg == 0) {`
`5688`		`- dst_f32[first_row + row] = all_sum;`
	`5687`	`+ dst_f32[first_row + row] = sum_all;`
`5689`	`5688`	`}`
`5690`	`5689`	`}`
`5691`	`5690`	`}`
`@@ -5732,7 +5731,7 @@ void kernel_mul_mv_iq1_m_f32_impl(`
`5732`	`5731`	`device const float * y = (device const float *) (src1 + offset1);`
`5733`	`5732`
`5734`	`5733`	`float yl[32];`
`5735`		`- float sumf[nr0]={0.f}, all_sum;`
	`5734`	`+ float sumf[nr0]={0.f};`
`5736`	`5735`
`5737`	`5736`	`const int nb32 = nb * (QK_K / 32);`
`5738`	`5737`
`@@ -5792,9 +5791,9 @@ void kernel_mul_mv_iq1_m_f32_impl(`
`5792`	`5791`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5793`	`5792`
`5794`	`5793`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5795`		`- all_sum = simd_sum(sumf[row]);`
	`5794`	`+ float sum_all = simd_sum(sumf[row]);`
`5796`	`5795`	`if (tiisg == 0) {`
`5797`		`- dst_f32[first_row + row] = all_sum;`
	`5796`	`+ dst_f32[first_row + row] = sum_all;`
`5798`	`5797`	`}`
`5799`	`5798`	`}`
`5800`	`5799`	`}`
`@@ -5848,7 +5847,7 @@ void kernel_mul_mv_iq4_nl_f32_impl(`
`5848`	`5847`	`threadgroup_barrier(mem_flags::mem_threadgroup);`
`5849`	`5848`
`5850`	`5849`	`float4 yl[4];`
`5851`		`- float sumf[nr0]={0.f}, all_sum;`
	`5850`	`+ float sumf[nr0]={0.f};`
`5852`	`5851`
`5853`	`5852`	`device const float * yb = y + ix * QK4_NL + it * 8;`
`5854`	`5853`
`@@ -5897,9 +5896,9 @@ void kernel_mul_mv_iq4_nl_f32_impl(`
`5897`	`5896`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`5898`	`5897`
`5899`	`5898`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`5900`		`- all_sum = simd_sum(sumf[row]);`
	`5899`	`+ float sum_all = simd_sum(sumf[row]);`
`5901`	`5900`	`if (tiisg == 0) {`
`5902`		`- dst_f32[first_row + row] = all_sum;`
	`5901`	`+ dst_f32[first_row + row] = sum_all;`
`5903`	`5902`	`}`
`5904`	`5903`	`}`
`5905`	`5904`	`}`
`@@ -5954,7 +5953,7 @@ void kernel_mul_mv_iq4_xs_f32_impl(`
`5954`	`5953`	`threadgroup_barrier(mem_flags::mem_threadgroup);`
`5955`	`5954`
`5956`	`5955`	`float4 yl[4];`
`5957`		`- float sumf[nr0]={0.f}, all_sum;`
	`5956`	`+ float sumf[nr0]={0.f};`
`5958`	`5957`
`5959`	`5958`	`device const float * yb = y + ix * QK_K + ib * 32 + il * 8;`
`5960`	`5959`
`@@ -6000,9 +5999,9 @@ void kernel_mul_mv_iq4_xs_f32_impl(`
`6000`	`5999`	`device float * dst_f32 = (device float ) dst + (uint64_t)imargs.ne0args.ne1 + (uint64_t)r1args.ne0;`
`6001`	`6000`
`6002`	`6001`	`for (int row = 0; row < nr0 && first_row + row < args.ne0; ++row) {`
`6003`		`- all_sum = simd_sum(sumf[row]);`
	`6002`	`+ float sum_all = simd_sum(sumf[row]);`
`6004`	`6003`	`if (tiisg == 0) {`
`6005`		`- dst_f32[first_row + row] = all_sum;`
	`6004`	`+ dst_f32[first_row + row] = sum_all;`
`6006`	`6005`	`}`
`6007`	`6006`	`}`
`6008`	`6007`	`}`