Repository: ggml-org/llama.cpp
Pull request: CUDA: Hide latency of bias and gate-loading for fused `mul_mat_vec_q`

This PR hides latency of bias and gate-loading for fused `mul_mat_vec_q` kernel by loading them into registers before computation of the dot-product, effectively batching them together with said dot-product. As a lot of threads are alive in this part of the kernel still, the warp scheduler has enough threads available to effectively hide the cost of loading those two single floats.

This gives 3-14% E2E speed-up for gpt-oss models (qwen3moe does not use bias and gate, and I am unaware of any other MoE model that uses bias and gate which I could run E2E perf tests on). The kernel themselves are up to 20% faster for gpt-oss.

| GPU                                                     | Model                   | Test   |   t/s master |   t/s this branch ([babfd19](https://github.com/ggml-org/llama.cpp/pull/16847/commits/babfd190857b5992f42d28e57b953e0b8363541b)) |   Speedup |
|:--------------------------------------------------------|:------------------------|:-------|-------------:|-----------------------------------------------:|----------:|
| RTX 4000 SFF Ada                                 | gpt-oss 20B MXFP4 MoE   | tg128  |        85.41 |                                          88.20 |      1.03 |
| RTX 4000 SFF Ada                                 | gpt-oss 20B MXFP4 MoE   | tg256  |        85.66 |                                          88.57 |      1.03 |
| RTX 4000 SFF Ada                                 | gpt-oss 20B MXFP4 MoE   | tg512  |        84.80 |                                          87.62 |      1.03 |
| RTX 6000 Ada                                     | gpt-oss 20B MXFP4 MoE   | tg128  |       240.55 |                                         248.09 |      1.03 |
| RTX 6000 Ada                                     | gpt-oss 20B MXFP4 MoE   | tg256  |       245.03 |                                         253.23 |      1.03 |
| RTX 6000 Ada                                     | gpt-oss 20B MXFP4 MoE   | tg512  |       242.32 |                                         250.41 |      1.03 |
| RTX PRO 4500 BW                                  | gpt-oss 20B MXFP4 MoE   | tg128  |       212.97 |                                         223.10 |      1.05 |
| RTX PRO 4500 BW                                  | gpt-oss 20B MXFP4 MoE   | tg256  |       220.06 |                                         236.34 |      1.07 |
| RTX PRO 4500 BW                                  | gpt-oss 20B MXFP4 MoE   | tg512  |       218.21 |                                         237.13 |      1.09 |
| RTX PRO 6000 BW Max-Q                            | gpt-oss 120B MXFP4 MoE  | tg128  |       198.64 |                                         208.72 |      1.05 |
| RTX PRO 6000 BW Max-Q                            | gpt-oss 120B MXFP4 MoE  | tg256  |       222.68 |                                         238.11 |      1.07 |
| RTX PRO 6000 BW Max-Q                            | gpt-oss 120B MXFP4 MoE  | tg512  |       224.01 |                                         240.44 |      1.07 |
| RTX PRO 6000 BW Max-Q                            | gpt-oss 20B MXFP4 MoE   | tg128  |       296.58 |                                         338.69 |      1.14 |
| RTX PRO 6000 BW Max-Q                            | gpt-oss 20B MXFP4 MoE   | tg256  |       314.90 |                                         356.62 |      1.13 |
| RTX PRO 6000 BW Max-Q                            | gpt-oss 20B MXFP4 MoE   | tg512  |       329.31 |                                         353.40 |      1.07 |



Changed files:
filename: ggml/src/ggml-cuda/mmvq.cu
patch: @@ -190,12 +190,28 @@ static __global__ void mul_mat_vec_q(
 
     const uint32_t channel_bias = ids ? channel_x : channel_dst;
 
+    float x_biases[ncols_dst][rows_per_cuda_block]    = { { 0.0f } };
+    float gate_biases[ncols_dst][rows_per_cuda_block] = { { 0.0f } };
     if constexpr (has_fusion) {
         if (use_bias) {
             x_bias = x_bias + sample_dst*stride_sample_dst + channel_bias*stride_channel_dst + row0;
+            // 1. Hide latency by prefetching bias and gate here
+            // 2. load only on threads that won't die after partial sum calculation
+            if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 &&
+                (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) {
+                for (int j = 0; j < ncols_dst; ++j) {
+                    x_biases[j][threadIdx.x] = x_bias[j * stride_col_dst + threadIdx.x];
+                }
+            }
         }
         if (use_gate_bias) {
             gate_bias = gate_bias + sample_dst*stride_sample_dst + channel_bias*stride_channel_dst + row0;
+            if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 &&
+                (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) {
+                for (int j = 0; j < ncols_dst; ++j) {
+                    gate_biases[j][threadIdx.x] = gate_bias[j * stride_col_dst + threadIdx.x];
+                }
+            }
         }
     }
 
@@ -283,12 +299,12 @@ static __global__ void mul_mat_vec_q(
             float result = tmp[j][threadIdx.x];
             if constexpr (has_fusion) {
                 if (use_bias) {
-                    result += x_bias[j*stride_col_dst + threadIdx.x];
+                    result += x_biases[j][threadIdx.x];
                 }
                 if (use_gate) {
                     float gate_value = tmp_gate[j][threadIdx.x];
                     if (use_gate_bias) {
-                        gate_value += gate_bias[j*stride_col_dst + threadIdx.x];
+                        gate_value += gate_biases[j][threadIdx.x];
                     }
                     switch (active_glu) {
                         case GGML_GLU_OP_SWIGLU: