Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
spiritbuun merged 2 commits into from
May 5, 2026
Merged

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

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
3416ff1
ggml-cuda: fix cublasSgemm crash with zero-dim matrices in speculativ…
May 2, 2026
a12fd99
ggml-cuda: extend zero-dim guard to cover ne00 and ldc
May 3, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 10 additions & 0 deletions ggml/src/ggml-cuda/ggml-cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1493,6 +1493,16 @@ static void ggml_cuda_op_mul_mat_cublas(
// ldc == nrows of the matrix that cuBLAS writes into
int64_t ldc = id == ctx.device ? ne0 : row_diff;

// Guard: cuBLAS requires m >= 1, n >= 1, k >= 1 for Sgemm/GemmEx.
// During speculative decoding (DFlash/copyspec), draft verification can
// produce non-consecutive token positions which result in zero-size
// sub-matrices. cuBLAS treats these as invalid parameters and aborts
// with CUBLAS_STATUS_INVALID_VALUE. Zero-size GEMMs are defined as
// no-ops (no output written), matching OpenBLAS and MKL behavior.
if (row_diff == 0 || src1_ncols == 0 || ne10 == 0 || ne00 == 0 || ldc == 0) {
return;
}

const int cc = ggml_cuda_info().devices[id].cc;

const bool supports_bf16 = GGML_CUDA_CC_IS_NVIDIA(cc) || GGML_CUDA_CC_IS_AMD(cc) ||
Expand Down
Loading