Enhance ggml_cuda_mul_mat_cublas with compute type fallback

[theIvanR]

Added graceful CUDA compute type fallback for MTP on older GPUs

Overview

Adds hardware capability checks for CUDA compute type selection to prevent crashes on older NVIDIA architectures when using MTP models.

The original MTP implementation could select BF16 compute paths on GPUs without BF16 support (e.g. Kepler), resulting in cuBLAS failures. This patch adds a fallback mechanism that automatically selects a supported compute type:

• BF16 → FP16 when fast FP16 hardware is available
• BF16 → FP32 otherwise

Modern GPUs with BF16 support remain unaffected. In reference to the original MTP PR #22673

Additional information

Tested with Qwen3.6 35B-A3B Q4XL on Tesla K40 (Kepler):

• Stock: ~22 tok/s
• MTP disabled: ~17.5 tok/s
• MTP enabled (2 draft tokens): ~23.5 tok/s

The fix allows MTP inference to run correctly on older CUDA architectures without requiring BF16 support.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: I used Deepseek newest model and Chatgpt for debugging ideas and helping me with my english. Implementation is my own.

[ggml-gh-bot]

Hi @theIvanR, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

• AI-generated content: This project does not accept PRs, descriptions or commit messages that are fully or predominantly AI-generated. If you have used AI to assist you in writing code, please make sure to disclose that explicitly.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[theIvanR]

CUDA: Fix BF16 compute type selection on unsupported architectures

Problem

When running MTP-enabled models on older NVIDIA GPUs (for example Kepler / compute capability 3.5), llama.cpp could crash inside the CUDA backend.

The issue was caused by ggml-cuda.cu selecting GGML_TYPE_BF16 as the cuBLAS compute type without verifying that the active GPU actually supported BF16 execution.

This was especially noticeable with MTP models because the additional MTP projection layers could request a BF16 compute path that normal inference did not normally hit.

On unsupported architectures, this resulted in an invalid cuBLAS execution path instead of a graceful fallback.

Affected hardware includes:

• Kepler (sm_35)
• Maxwell
• Pascal
• Other GPUs without BF16 MMA support

---

Patch

The fix adds a hardware capability validation step after compute type selection and environment override handling.

The previous flow was:

select default compute type
        |
        v
apply GGML_CUDA_CUBLAS_COMPUTE_TYPE override
        |
        v
dispatch cuBLAS kernel

The patched flow becomes:

select default compute type
        |
        v
apply GGML_CUDA_CUBLAS_COMPUTE_TYPE override
        |
        v
validate compute type against GPU capabilities
        |
        +--> BF16 supported:
        |       use BF16
        |
        +--> BF16 unsupported:
                use FP16 if fast FP16 is available
                otherwise use FP32
        |
        v
dispatch cuBLAS kernel

The patch uses the existing CUDA capability checks:

bf16_mma_hardware_available(cc)
fast_fp16_hardware_available(cc)

instead of hardcoding architecture-specific checks.

Resulting behavior:

• Ampere and newer:

  • BF16 remains enabled

• Volta/Turing:

  • BF16 falls back to FP16 where supported

• Pascal/Maxwell/Kepler:

  • BF16 falls back to FP32

This prevents unsupported GPUs from entering invalid BF16 execution paths while preserving BF16 acceleration on modern hardware.

---

Expected performance

Tested with:

• Model: Qwen 3.6 35B A3B
• Quantization: Q4XL
• Backend: llama.cpp CUDA
• Hardware: Tesla K40 (Kepler)

Observed approximate token generation speeds:

Stock llama.cpp:
    ~22 tokens/sec

MTP disabled:
    ~17.5 tokens/sec

MTP enabled (2 token prediction):
    ~23.5 tokens/sec

The patch restores MTP functionality on older CUDA architectures without requiring BF16 support.

Interestingly, on Kepler hardware forcing FP32 compute can outperform FP16 paths. This is likely because Kepler does not have fast FP16 arithmetic, and cuBLAS can select a more optimized FP32 GEMM path instead.

---

Notes

This patch does not disable BF16 globally.

It only validates that the requested compute type is supported by the active GPU before dispatching the cuBLAS operation.

The intended behavior is:

"Use the fastest available compute type, but never select a compute type that the hardware cannot execute."

This keeps modern GPUs on BF16/FP16 paths while maintaining compatibility with older CUDA architectures.

Patch to ggml-cuda.cu:

// patched version with generic compute type capability fallback
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(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) {
        compute_type = GGML_TYPE_F32;
    }

    const char * env_c = getenv("GGML_CUDA_CUBLAS_COMPUTE_TYPE");
    if (env_c != nullptr) {
        std::string env_cpp = env_c;

        for (char & c : env_cpp) {
            c = std::tolower(c);
        }

        if (env_cpp == "f32" || env_cpp == "fp32") {
            compute_type = GGML_TYPE_F32;
        } else if (env_cpp == "f16" || env_cpp == "fp16") {
            compute_type = GGML_TYPE_F16;
        } else if (env_cpp == "bf16") {
            compute_type = GGML_TYPE_BF16;
        } else if (env_cpp != "auto") {
            GGML_LOG_WARN(
                "%s: unknown value for GGML_CUDA_CUBLAS_COMPUTE_TYPE: %s",
                __func__,
                env_cpp.c_str());
        }
    }

    // Validate requested compute type against hardware capabilities.
    if (compute_type == GGML_TYPE_BF16 && !bf16_mma_hardware_available(cc)) {
        static bool warned = false;

        const bool can_use_fp16 = fast_fp16_hardware_available(cc);

        if (!warned) {
            GGML_LOG_WARN(
                "BF16 compute type not supported on device CC=%d; "
                "falling back to %s.\n",
                cc,
                can_use_fp16 ? "F16" : "F32");

            warned = true;
        }

        compute_type = can_use_fp16
            ? GGML_TYPE_F16
            : GGML_TYPE_F32;
    }

    switch (compute_type) {
        case GGML_TYPE_F32:
            ggml_cuda_mul_mat_cublas_impl<GGML_TYPE_F32>(ctx, src0, src1, dst);
            break;

        case GGML_TYPE_BF16:
            ggml_cuda_mul_mat_cublas_impl<GGML_TYPE_BF16>(ctx, src0, src1, dst);
            break;

        case GGML_TYPE_F16:
            ggml_cuda_mul_mat_cublas_impl<GGML_TYPE_F16>(ctx, src0, src1, dst);
            break;

        default:
            GGML_ABORT("fatal error");
    }
}
// end of patch

[cb88]

My MI50s definitely do not support BF16 and have for the most part always worked with MTP for Q4 and Q8 where they have the most OP/s available.

if you are running a full BF16 model then... you seem to be out of compute anyway and should probably prefer a model that uses a native compute type on a GPU with less compute anyway???? Note how MTP hardly speeds anything up at all only about 5-10% when it should be a 50-100% speedup etc...

[theIvanR]

@cb88 MI50 is an AMD GPU and entirely unrelated to the matter as it is not using Cuda methods. As for speeds, I was using the lowest speed achieved in testing and in some cases speedups can indeed approach 50 percent.

EDIT, which one specifically are you using of the MI50's the 16GB or the 32GB? I have heard mixed reviews on these and would be interested in performances you achieve with the same LLM on Vulkan (does it support ROCM?)

[am17an]

This is not a valid problem/fix. Please create an issue with a reproduction

[theIvanR]

@am17an I think there was a miscommunication here, as it is both a problem and a valid fix as well as a decent writeup, however as you have asked I reformatted my pull request to be an issue. Link here: #25713