From 8c6ca4ab76d48022d8639c0c8b75ac907d2e3c6b Mon Sep 17 00:00:00 2001
From: Konrad Moren <kmoren@nvidia.com>
Date: Mon, 5 Oct 2026 13:43:07 +0200
Subject: [PATCH] CUDA: Optimize accumulation in mmq for NVFP4 type
 (llama/29857)

* ggml_cuda: optimize accumulation in mmq_vec_dot_fp4_fp4_mma for better performance

* remove whitespace

* fix: correct indentation in mma_block_scaled_fp4 loop
---
 ggml/src/ggml-cuda/mmq-vec-dot.cuh | 7 ++-----
 1 file changed, 2 insertions(+), 5 deletions(-)

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 <ggml_type type, int J, bool fallback> 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<tile_C *>(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<type>(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];
-                }
             }
         }
     }