Eval bug: MTP decoding crash on pre-Ampere GPUs (with working patch!)

[theIvanR]

Name and Version

version: 10011 (bf2c86d)
built with MSVC 19.29.30159.0 for Windows AMD64

Operating systems

Windows

GGML backends

CUDA

Hardware

• CPU: Xeon E5-2699 v3
• RAM: 64 GB
• GPU: 4× Tesla K40c (Compute Capability 3.5 / SM35)
• CUDA: 11.8

Models

Any MTP models, for example Qwen 3.6 35B A3B MTP GGUF model (unsloth from HF)

Problem description & steps to reproduce

Steps to reproduce

1. Build the latest llama.cpp from master.
2. Run a Qwen 3.6 35B A3B MTP GGUF model (unsloth from HF)
3. Use CUDA backend.
4. Enable MTP inference.

Run with:

@echo off
setlocal

set CUDA_LAUNCH_BLOCKING=1

set "LLAMA_EXE=C:\Users\Admin\source\llama.cpp\build_gpu_cuda\bin\llama-server.exe"

set "MODEL_DIR=C:\Users\Admin\Documents\LLM Models\unsloth\Qwen3.6-35B-A3B-MTP-GGUF"
set "MODEL=Qwen3.6-35B-A3B-UD-Q4_K_XL.gguf"
set "MMPROJ=mmproj-F16.gguf"

pushd "%MODEL_DIR%" || exit /b 1

"%LLAMA_EXE%" ^
  -m "%MODEL%" ^
  --mmproj "%MMPROJ%" ^
  -ngl 999 ^
  --tensor-split 0.25,0.25,0.25,0.25 ^
  --spec-type draft-mtp --spec-draft-n-max 2

set "ERR=%ERRORLEVEL%"
popd
exit /b %ERR%

Expected Outcome

• Working inference.

Actual Outcome

• Without MTP enabled model inference works, with MTP enabled model inference crashes.
  cuda.cu:104: CUDA error

Cause and Remediation of Problem

• in short, BF16 compute is triggered to eagerly.
• for a detailed (working) solution see my PR: Enhance ggml_cuda_mul_mat_cublas with compute type fallback #25680

First Bad Commit

When this PR was merged: #25679

Relevant log output

Logs

cuda.cu:104: CUDA error

[xxxajk]

broken on Vulkan too, segfaults

[1064210.858450] llama-server[1601344]: segfault at 10 ip 0000766ca6ae51f4 sp 00007ffeea720178 error 4 in libllama.so.0.0.10068[e51f4,766ca6a95000+217000] likely on CPU 1 (core 1, socket 0)
[1064210.858458] Code: 00 c3 0f 1f 40 00 f3 0f 1e fa 8b 47 08 c3 0f 1f 84 00 00 00 00 00 f3 0f 1e fa 8b 47 0c c3 0f 1f 84 00 00 00 00 00 f3 0f 1e fa <8b> 47 10 c3 0f 1f 84 00 00 00 00 00 f3 0f 1e fa 8b 47 14 c3 0f 1f

Tesla P40, AMD Radeon RX 5700 XT, mtp-Qwen_Qwen3.6-27B-Q8_0
:-/

[theIvanR]

@xxxajk did you try running my patched fork?

[xxxajk]

@xxxajk did you try running my patched fork?

Have a prebuild for Vulkan on x86 Linux by any chance?

[stew675]

Leaving this potential fix here. I have no way to reproduce as I do not own an nVidia GPU. Perhaps someone can try it and file a PR with it if it solves the issue?

diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 598f3228c..5effdee59 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -1617,10 +1617,12 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
 }
 
 static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    const int cc = ggml_cuda_info().devices[ctx.device].cc;
+
     ggml_type compute_type = src0->type;
     if (ggml_is_quantized(compute_type)) {
-        compute_type = fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
-    } else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc)) {
+        compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
+    } else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc)) {
         compute_type = GGML_TYPE_F32;
     }
     if (dst->op_params[0] == GGML_PREC_F32) {
@@ -1644,6 +1646,17 @@ static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml
         }
     }
 
+    // cuBLAS does not support BF16 GEMM on all GPUs, fall back to a supported compute type
+    if (compute_type == GGML_TYPE_BF16 && !bf16_mma_hardware_available(cc)) {
+        static bool warned = false;
+        compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
+        if (!warned) {
+            GGML_LOG_WARN("%s: BF16 compute type not supported on device CC=%d; using %s\n",
+                    __func__, cc, compute_type == GGML_TYPE_F16 ? "F16" : "F32");
+            warned = true;
+        }
+    }
+
     switch (compute_type) {
         case GGML_TYPE_F32:
             ggml_cuda_mul_mat_cublas_impl<GGML_TYPE_F32>(ctx, src0, src1, dst);