From be4b2d1d654d563db92af887b16dae9383c0540d Mon Sep 17 00:00:00 2001
From: cheese-cakee <farzanaman99@gmail.com>
Date: Tue, 14 Jul 2026 04:19:46 +0530
Subject: [PATCH] perf(ggml-cuda): skip shared reduction and barrier for
 single-wave mul_mat_vec_q

Wrap the cross-wave reduction epilogue of mul_mat_vec_q in
if constexpr (nwarps > 1). For single-warp instantiations this removes a
block-wide __syncthreads() and the shared-memory staging that the reduction
would otherwise emit, without changing any computed result.

The reduction is a no-op when there is only one warp: the shared array is only
written under threadIdx.y > 0 (which never holds), the combine loop runs
nwarps-1 == 0 iterations, and the barrier synchronizes a single warp. Guarding
the block on nwarps > 1 makes that explicit so the compiler no longer has to
emit the barrier for the single-wave decode path. For nwarps > 1 the statements
are unchanged, only nested inside the guard.
---
 .../deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu | 63 +++++++++++--------
 1 file changed, 38 insertions(+), 25 deletions(-)

diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu
index ec6288f97..befd0e731 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu
+++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu
@@ -608,32 +608,54 @@ static __global__ void mul_mat_vec_q(
         }
     }
 
-    __shared__ float tmp_shared[nwarps-1 > 0 ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size];
-    __shared__ float tmp_shared_gate[(has_fusion && (nwarps-1 > 0)) ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size];
-    if constexpr (!has_fusion) {
-        (void) tmp_shared_gate;
-    } else if (!use_gate) {
-        (void) tmp_shared_gate;
-    }
+    // A one-wave specialization has no cross-wave partials. Avoid reserving
+    // shared memory and issuing a block barrier in that common decode path.
+    // This does not change the per-lane accumulation or warp reduction order.
+    if constexpr (nwarps > 1) {
+        __shared__ float tmp_shared[nwarps-1][ncols_dst][rows_per_cuda_block][warp_size];
+        __shared__ float tmp_shared_gate[has_fusion ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size];
+        if constexpr (!has_fusion) {
+            (void) tmp_shared_gate;
+        } else if (!use_gate) {
+            (void) tmp_shared_gate;
+        }
+
+        if (threadIdx.y > 0) {
+#pragma unroll
+            for (int j = 0; j < ncols_dst; ++j) {
+#pragma unroll
+                for (int i = 0; i < rows_per_cuda_block; ++i) {
+                    tmp_shared[threadIdx.y-1][j][i][threadIdx.x] = tmp[j][i];
+                    if constexpr (has_fusion) {
+                        if (use_gate) {
+                            tmp_shared_gate[threadIdx.y-1][j][i][threadIdx.x] = tmp_gate[j][i];
+                        }
+                    }
+                }
+            }
+        }
+
+        __syncthreads();
+        if (threadIdx.y > 0) {
+            return;
+        }
 
-    if (threadIdx.y > 0) {
 #pragma unroll
         for (int j = 0; j < ncols_dst; ++j) {
 #pragma unroll
             for (int i = 0; i < rows_per_cuda_block; ++i) {
-                tmp_shared[threadIdx.y-1][j][i][threadIdx.x] = tmp[j][i];
-                if constexpr (has_fusion) {
-                    if (use_gate) {
-                        tmp_shared_gate[threadIdx.y-1][j][i][threadIdx.x] = tmp_gate[j][i];
+#pragma unroll
+                for (int l = 0; l < nwarps-1; ++l) {
+                    tmp[j][i] += tmp_shared[l][j][i][threadIdx.x];
+                    if constexpr (has_fusion) {
+                        if (use_gate) {
+                            tmp_gate[j][i] += tmp_shared_gate[l][j][i][threadIdx.x];
+                        }
                     }
                 }
             }
         }
     }
-    __syncthreads();
-    if (threadIdx.y > 0) {
-        return;
-    }
 
     dst += sample_dst*stride_sample_dst + channel_dst*stride_channel_dst + row0;
 
@@ -642,15 +664,6 @@ static __global__ void mul_mat_vec_q(
     for (int j = 0; j < ncols_dst; ++j) {
 #pragma unroll
         for (int i = 0; i < rows_per_cuda_block; ++i) {
-#pragma unroll
-            for (int l = 0; l < nwarps-1; ++l) {
-                tmp[j][i] += tmp_shared[l][j][i][threadIdx.x];
-                if constexpr (has_fusion) {
-                    if (use_gate) {
-                        tmp_gate[j][i] += tmp_shared_gate[l][j][i][threadIdx.x];
-                    }
-                }
-            }
             tmp[j][i] = warp_reduce_sum<warp_size>(tmp[j][i]);
             if constexpr (has_fusion) {
                 if (use_gate) {