Opt from #880 also for iqk cuda gemv

Iwan Kawrakow · Iwan Kawrakow · commit 5266eeea18cc · 2025-11-11T09:59:56.000+02:00
diff --git a/ggml/src/ggml-cuda/iqk_mmvq_templates.cuh b/ggml/src/ggml-cuda/iqk_mmvq_templates.cuh
@@ -104,7 +104,7 @@ __device__ void iqk_mul_mat_vec_q_kerne(
         }
 
         if (threadIdx.x < rows_per_cuda_block && (rows_per_cuda_block == 1 || row0 + threadIdx.x < nrows_dst)) {
-            dst[j*nrows_dst + row0 + threadIdx.x] = bias ? tmp[j][threadIdx.x] + bias[j*nrows_dst + row0 + threadIdx.x] : tmp[j][threadIdx.x];
+            dst[j*nrows_dst + row0 + threadIdx.x] = bias ? tmp[j][threadIdx.x] + bias[row0 + threadIdx.x] : tmp[j][threadIdx.x];
         }
     }
 }
@@ -211,8 +211,8 @@ __device__ void iqk_fused_mul_mat_vec_q_kernel(
                 default: {
                     constexpr float alpha = 1.702f;
                     constexpr float limit = 7.0f;
-                    g += bias_g[j*nrows_dst + row0 + threadIdx.x];
-                    u += bias_u[j*nrows_dst + row0 + threadIdx.x];
+                    g += bias_g[row0 + threadIdx.x];
+                    u += bias_u[row0 + threadIdx.x];
                     g = fminf(g, limit);
                     u = fmaxf(fminf(u, limit), -limit);
                     r = g / (1.0f + expf(-g * alpha)) * (1.0f + u);

Original file line number	Diff line number	Diff line change
`@@ -104,7 +104,7 @@ __device__ void iqk_mul_mat_vec_q_kerne(`
`104`	`104`	`}`
`105`	`105`
`106`	`106`	`if (threadIdx.x < rows_per_cuda_block && (rows_per_cuda_block == 1 \|\| row0 + threadIdx.x < nrows_dst)) {`
`107`		`- dst[jnrows_dst + row0 + threadIdx.x] = bias ? tmp[j][threadIdx.x] + bias[jnrows_dst + row0 + threadIdx.x] : tmp[j][threadIdx.x];`
	`107`	`+ dst[j*nrows_dst + row0 + threadIdx.x] = bias ? tmp[j][threadIdx.x] + bias[row0 + threadIdx.x] : tmp[j][threadIdx.x];`
`108`	`108`	`}`
`109`	`109`	`}`
`110`	`110`	`}`
`@@ -211,8 +211,8 @@ __device__ void iqk_fused_mul_mat_vec_q_kernel(`
`211`	`211`	`default: {`
`212`	`212`	`constexpr float alpha = 1.702f;`
`213`	`213`	`constexpr float limit = 7.0f;`
`214`		`- g += bias_g[j*nrows_dst + row0 + threadIdx.x];`
`215`		`- u += bias_u[j*nrows_dst + row0 + threadIdx.x];`
	`214`	`+ g += bias_g[row0 + threadIdx.x];`
	`215`	`+ u += bias_u[row0 + threadIdx.x];`
`216`	`216`	`g = fminf(g, limit);`
`217`	`217`	`u = fmaxf(fminf(u, limit), -limit);`
`218`	`218`	`r = g / (1.0f + expf(-g * alpha)) * (1.0f + u);`