diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index 4ca6542d8..b39e4d579 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -1225,14 +1225,11 @@ template static __device__ __forceinline_ #pragma unroll for (int n = 0; n < ntx; ++n) { + // accumulate in place into the output sum array + tile_C & C = *reinterpret_cast(sum + (j0 / tile_C::J + n) * tile_C::ne); #pragma unroll for (int frag = 0; frag < nfrags; ++frag) { - tile_C C = {}; mma_block_scaled_fp4(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); -#pragma unroll - for (int l = 0; l < tile_C::ne; ++l) { - sum[(j0 / tile_C::J + n) * tile_C::ne + l] += C.x[l]; - } } } }