Commit 5d4c8eeb2 for llama.cpp

commit 5d4c8eeb24db049acf74bd6f354bed6eaeffc7dc
Author: Feiteng <lifeiteng0422@gmail.com>
Date:   Thu Nov 12 11:45:56 2026 +0800

    CUDA: skip empty MoE tiles in MMQ stream-k fixup (#32067)

    * CUDA: skip empty MoE tiles in MMQ stream-k fixup

    mul_mat_q_stream_k_fixup loads ids_dst[col_low + j] for all j < J before
    checking j_sup. mul_mat_q returns early on col_low == col_high, the fixup
    kernel did not. mm_ids_helper emits one extra tile per expert that may be
    empty; for trailing experts col_low == ncols_dst, so the load reads one
    element past the J_best-1 padding added in #29953. The kernel returns at
    j >= j_sup before storing anything, so results are unaffected.

    Return early when j_sup <= 0, like the main kernel. The return is
    block-uniform (tsj is loaded once per block) and precedes __syncthreads().

    Add test_mul_mat_id_unused_tail to test-backend-ops: only the first
    n_mats/2 experts are selected, so the trailing experts have zero rows and
    their tiles are empty; m=33, k=512 keeps tile efficiency low enough that
    stream-k splits along K and launches the fixup kernel.

    * CUDA: drop the extra MUL_MAT_ID test case and the explanatory comment (review)

diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
index 4d67555ef..42a3afd24 100644
--- a/ggml/src/ggml-cuda/mmq.cuh
+++ b/ggml/src/ggml-cuda/mmq.cuh
@@ -1396,6 +1396,10 @@ static __global__ void mul_mat_q_stream_k_fixup(
     const int col_high = tsj.col_high;
     const int j_sup    = col_high - col_low;

+    if (j_sup <= 0) {
+        return;
+    }
+
     for (int j = threadIdx.y*warp_size + threadIdx.x; j < J; j += nwarps*warp_size) {
         ids_dst_shared[j] = ids_dst[col_low + j];
     }