[NickF93]

Overview

Mistral Small 4 requests GGML_PREC_F32 for its expert down-projections because their activations can exceed FP16 range. CUDA currently rejects these MUL_MAT_ID operations, so the scheduler executes them on CPU even when the loader reports all model layers offloaded to GPU.

This change handles that precision requirement through the existing sorted-expert fallback. It keeps the gathered activations in FP32 and calls the existing FP32 cuBLAS implementation for each expert slice.

Fusion paths that cannot preserve this requirement are excluded, and synchronization reporting matches the fallback. Support is limited to NVIDIA devices and the input types and layouts handled by that implementation. HIP and MUSA retain their existing behavior.

The regression tests cover Q4_K, Q8_0 and BF16 weights with activations above FP16 range, including single-token, multi token and fusion graphs. The CUDA operation support table is updated accordingly.

Additional information

Fixes #30224.

Validation used patched commit 1e29ebb against its unpatched parent, 609290b.

Correctness

• The focused regression ran on an RTX 3090, RTX PRO 4500 Blackwell and RTX PRO 4000 Blackwell. Before the patch, all 36 selected cases were unsupported. After the patch, all 36 executed and passed against CPU references.
• All 36 cases also passed with GGML_CUDA_CUBLAS_COMPUTE_TYPE=f16, confirming that this override does not narrow the required FP32 activations.
• All 18 selected ordinary-operation cases passed.
• Native CPU CI completed: Debug/Release CTest passed 50/50 and 51/51, model-dependent tests passed 5/5 in each build, and 21,885 supported backend-operation cases passed.
• Native CUDA CI on an RTX A5000 completed the same CTest groups and passed 18,515 supported backend-operation cases, including 15 high-amplitude FP32 expert/fusion cases. Unsupported cases are excluded from these pass counts.

The CUDA CI run initially stopped at the final model tests because the temporary environment lacked external GNU time. After installing that prerequisite, the two unfinished model-test groups passed using the retained builds.

Real generation with Mistral Small 4 UD-Q4_K_M also succeeded. Scheduler logs placed all 36 expert down-projections on CUDA. The short canary produced the correct arithmetic answer with separate reasoning and final content.

A focused regression selector is:

./build/bin/test-backend-ops test \
    -b CUDA0 \
    -o MUL_MAT_ID,MUL_MAT_ID_FUSION,MUL_MAT_VEC_FUSION \
    -p 'n_mats=4.*n_used=2.*amax=100000' \
    --output csv

Performance and perplexity

Matched Release builds used GCC 13 and CUDA 13.3.73. The GPU comparison used Ling 3.0 Flash Q4_K_M on the three-GPU configuration above, with unchanged settings and runtime CUDA graphs disabled.

Each benchmark used three repetitions, 512-token prefill and 128-token decode:

Build	Prefill tok/s, mean ± SD	Decode tok/s, mean ± SD
Baseline before	714.98 ± 49.87	73.16 ± 0.62
Patched	710.36 ± 70.57	73.12 ± 1.15
Baseline after	719.03 ± 60.68	72.87 ± 0.16

The patched result differs from the bracketed baseline mean by -0.93% for prefill and +0.14% for decode. This bounded comparison did not demonstrate a meaningful slowdown.

WikiText-2 perplexity, using two 1024-token chunks, was identical at printed precision:

• CUDA, Ling Q4_K_M: 5.2846 ± 0.41607 before and after.
• CPU, the Qwen3 0.6B F16 CI fixture: 9.4179 ± 0.80961 before and after.

CPU benchmarks were also recorded, but the first baseline differed substantially from its repeat. They do not support a claim of a patch-induced CPU speedup.

The patch applied cleanly in memory against upstream commit 10a60cf303566e10d6a7a2774c17d2085503d87b. That current-master build was not runtime-tested; the results above apply to the tested patch commit.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES. AI assistance was used for investigation, review, validation analysis, CI/CD. I personally reviewed the complete diff and take responsibility for the submitted changes.

[agourakis82]

This describes a local follow-up patch based on the head of #30260 (1e29ebb). The code has not been pushed yet, so the mmid-f32.cu implementation described below is not part of this PR’s current diff.

The patch adds a narrowly scoped FP32 path for MUL_MAT_ID with GGML_PREC_F32 and Q8_0 or Q5_0 weights. It is NVIDIA CUDA only; HIP and MUSA retain their existing paths.

The same support predicate controls dispatch, synchronization requirements, and backend support reporting. When it accepts the operation, ggml_cuda_mul_mat_id calls ggml_cuda_mul_mat_id_f32 and returns before MMVQ, MMQ, MMF, or the fallback that sorts experts. needs_sync returns false for this path, avoiding the synchronization requirement that would otherwise prevent CUDA Graph use. supports_op also accepts it before reaching the existing rejection of F32 src1.

The supported layout is deliberately limited:

• K and M are at most 4096, with K divisible by the quantization block size.
• There are at most 128 experts and 64 selected IDs.
• ids->ne[1..3] must all be 1, so this handles a single token.
• b->ne[1] must be either 1 for broadcast or equal to the number of selected IDs.
• The output shape is [M, n_used], with the checked expert and row strides capped at 8192.

Operations outside those bounds continue through the existing path from #30260.

Execution is per selected ID slot. The selected expert is dequantized into a reusable F32 buffer containing K*M elements, followed by cublasSgemm(OP_T, OP_N, m, 1, k). SGEMM’s N is fixed at 1; this is not a batched GEMM. The activation row is selected with slot % b->ne[1].

Invalid expert IDs have different implementations in the two weight formats. For Q8_0, the dequantization kernel fills the buffer with NaN (0x7fc00000) before SGEMM. For Q5_0, dequantization safely reads expert 0, SGEMM runs, and mask_invalid_q5 then overwrites the affected output with NaN.

[agourakis82]

I’ve pushed the follow-up I described above:

agourakis82@2212996

It is based on 7790a92 and adds NVIDIA CUDA support for an explicit GGML_PREC_F32 requirement in single-token MUL_MAT_ID, limited to Q5_0 and Q8_0 weights and the layouts described in my previous comment.

Validation on an NVIDIA L4 passed all 12 supported regression cases in each configuration: default, F16 override, and CUDA Graphs disabled. Compute Sanitizer reported zero errors with graphs enabled and disabled.

For the full-model check, I compared 197 prediction positions in Mellum against an independent CUDA reference using the upstream decoder and matching SGEMM dimensions. The explicit-F32 experiment showed zero observed differences in logits and NLL, with no greedy disagreements.

The stock Mellum graph does not use this path after the upstream rescaling changes. The scope here is explicit F32 backend support. I also recorded operator latencies on the L4 as descriptive measurements, without claiming a speedup.

Would this fit as a follow-up to #30260?