ggml-org/llama.cpp b12316

on GitHub

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)