llama + spec: MTP Support by am17an · Pull Request #22673 · ggml-org/llama.cpp

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 (see my PR: #25680)

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