perf(ggml-cuda): skip shared reduction and barrier for single-wave mul_mat_vec_q

[cheese-cakee]

Summary

mul_mat_vec_q ends with a cross-wave reduction that combines each warp's
partial sums through shared memory and a block barrier. When the kernel is
instantiated with a single warp (nwarps == 1) that whole step is inert: there
are no other warps to combine. This change wraps the reduction in
if constexpr (nwarps > 1) so the single-wave path no longer carries a
__syncthreads() that synchronizes one warp. The change is byte-for-byte
identical on output; it only affects which instructions the compiler emits for
single-wave instantiations.

Changes

• ggml/src/ggml-cuda/mmvq.cu: guard the shared-memory cross-wave reduction in
  mul_mat_vec_q with if constexpr (nwarps > 1). One file, no API change, no
  new flags.

How it works

For nwarps > 1 the statements are unchanged, only nested inside the guard, so
every multi-warp instantiation behaves exactly as before (same shared array,
same store, same barrier, same combine loop).

For nwarps == 1 the guarded block is provably dead on clean main:

• tmp_shared / tmp_shared_gate are only written under threadIdx.y > 0, and
  with a single warp threadIdx.y is always 0, so the store never runs.
• the combine loop iterates l in [0, nwarps - 1), which is zero iterations.
• the __syncthreads() synchronizes a block that contains one warp.

The shared arrays were already sized with a nwarps-1 > 0 ? nwarps-1 : 1 guard,
so removing them for the single-wave case is not what changes; the barrier is.
ptxas dead-code-eliminates the unused shared array on its own, but it cannot
remove the barrier, because a barrier has cross-thread semantics it must
preserve. Guarding on nwarps > 1 states the intent in source and lets the
single-wave decode path drop the barrier.

Performance

Validated locally on an RTX 4050 Laptop (Ada, sm_89), CUDA 12.6.85.

• Codegen, single-wave: SASS for every nwarps == 1 instantiation of
  mul_mat_vec_q drops from one BAR.SYNC to zero. Multi-warp instantiations
  keep their single BAR.SYNC.
• Codegen, all instantiations: comparing ptxas -v for clean vs patched over
  378 compiled mul_mat_vec_q instantiations, shared-memory usage is identical
  for every one of them (0 differences). Register counts move by small amounts
  in both directions (mean -1.0 registers, i.e. net neutral) as a result of the
  reordering; no instantiation changes occupancy tier.
• Runtime: on this GPU the single-wave path is not reachable. The NVIDIA GENERIC
  parameter table only selects nwarps in {2, 4} for the batch sizes that route
  to mul_mat_vec_q (ncols_dst <= 8), so the removed barrier is never
  executed here and there is no measurable NVIDIA runtime change. An isolated
  microbench that forces the single-wave epilogue measures the removed barrier
  at within +/- 3 percent, i.e. below noise, because a one-warp __syncthreads()
  is nearly free on this architecture.

The intended beneficiary is decode on architectures whose parameter tables
select a single warp for small-batch mat-vec (AMD GCN and RDNA), where this is
the common path. That end to end measurement is not included here because it
needs the AMD hardware.

Limitations

• No end to end token/s numbers: this was validated on a single NVIDIA laptop
  GPU with synthetic inputs, not a full model run.
• The single-wave runtime win is not demonstrated on NVIDIA because NVIDIA does
  not instantiate the single-wave kernel for this op; the AMD measurement is
  deferred.
• Not gated behind an env flag. The change is byte-identical and matches the
  existing if constexpr style in this file, so a runtime toggle would add
  complexity with no upside.

Verification

All local, on RTX 4050 (sm_89), CUDA 12.6.

1. Byte-identity, real kernel. Built the CUDA backend twice from identical
   synthetic inputs, once with clean mmvq.cu and once with the patch, and
   dumped every mul_mat_vec_q result across a matrix of
   {Q4_0, Q8_0, Q4_K, Q5_K, Q6_K} x K in {256, 4096} x N in {1024, 1000} x
   ncols_dst in {1, 2, 4, 5, 8} (100 cases, run with LUCE_MMVQ_MAX_NCOLS=8 so
   batches 1..8 route to mul_mat_vec_q, exercising nwarps=4 for ncols 1..4 and
   nwarps=2 for ncols 5..8). The two dumps are byte-identical (same SHA-256,
   1,619,200 bytes).

2. Single-wave and fusion coverage. NVIDIA never instantiates the
   nwarps == 1 path, so an isolated probe reproduces the exact reduction
   epilogue verbatim for both the clean and patched structure and compares
   output for nwarps in {1, 2, 4} with fusion off and on. All identical.

3. Barrier removal. cuobjdump -sass confirms the BAR.SYNC is present in the
   clean single-wave epilogue and absent in the patched one, and unchanged for
   multi-warp.

[cubic-dev-ai]

No issues found across 1 file

_{Re-trigger cubic}