ggml-cuda: fix cublasSgemm crash with zero-dim matrices in speculative decoding

[Tongas]

Problem

During speculative decoding (DFlash/copyspec), draft verification produces non-consecutive token positions which cause slot management to generate sub-matrix batches where row_diff evaluates to 0.

cuBLAS requires ldc >= max(1, m) where m = row_diff. When row_diff == 0, ldc is also 0 (ldc = row_diff when id != ctx.device), violating the cuBLAS precondition and triggering CUBLAS_STATUS_INVALID_VALUE.

Root cause

No guard exists before the cublasSgemm/cublasGemmEx calls in ggml_cuda_op_mul_mat_cublas() to handle zero-dimension sub-matrices. Unlike OpenBLAS and MKL which define zero-size GEMMs as no-ops, cuBLAS treats them as invalid parameters and aborts.

Fix

Add an early-return guard after row_diff and ldc are computed. When any GEMM dimension (m, n, k) is zero, return immediately. This is safe because dst_dd_i is untouched and the caller does not read partial results.

Tested

Qwen3.6-27B Q4_K_M + DFlash draft, 55k context, RTX 3090 24GB

• Before: crash after ~27k tokens prefill during draft verification
• After: stable generation, multiple sequential requests without crash

Related upstream issues

• [Speculative decoding] feat: add DFlash support ggml-org/llama.cpp#22105 (DFlash PR: hybrid models cannot discard rejected suffixes via seq_rm due to non-decomposable recurrent state)
• Eval bug: llama-server Qwen3.5 vision behave weirdly on certain batchsize settings. ggml-org/llama.cpp#19929 (non-consecutive token position with Qwen3.5 vision models on multi-GPU setups)
• DFlash (Block Diffusion for Flash Speculative Decoding) ggml-org/llama.cpp#21569 (DFlash discussion: MoE hybrid limitations)

Authored-by: Gaston Parravicini

Overview

Fix cublasSgemm/cublasGemmEx crash triggered during DFlash speculative decoding when draft verification produces zero-dimension sub-matrices. Adds an early-return guard in ggml_cuda_op_mul_mat_cublas() before any cuBLAS call.

Additional information

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure:
  -YES — Claude (Anthropic) was used to assist in root cause analysis,
  identifying the zero-dimension guard as the fix, and drafting the commit
  message. The patch was applied, compiled, and tested on hardware by the
  author. All code reviewed line by line before submission.