From 10872be7f8fe59570e85d9bab17efc708849f1fb Mon Sep 17 00:00:00 2001 From: ynankani Date: Thu, 1 Oct 2026 11:23:52 +0000 Subject: [PATCH] CUDA: Handle compute type for NVFP4 on cublass path (llama/29173) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * CUDA: Handle compute type for NVFP4 on cublass path Signed-off-by: ynankani * Use BF16 compute type for quantized models if HW allows Signed-off-by: ynankani * Set acc prec to bf16 for nvfp4 as it needs atleast bf16 range Signed-off-by: ynankani * Update ggml/src/ggml-cuda/ggml-cuda.cu Co-authored-by: Johannes Gäßler * preserve op_params for per-expert matmul Signed-off-by: ynankani --------- Signed-off-by: ynankani Co-authored-by: Johannes Gäßler --- ggml/src/ggml-cuda/ggml-cuda.cu | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index ebbb5c106..f9fb46c2a 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1616,6 +1616,7 @@ 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; + const ggml_prec prec = (ggml_prec) ggml_get_op_params_i32(dst, 0); 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; @@ -1629,7 +1630,10 @@ static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml compute_type = GGML_TYPE_F32; } } - if (dst->op_params[0] == GGML_PREC_F32) { + // F16 is the only compute type that can not satisfy a request for BF16 + if (prec == GGML_PREC_BF16 && compute_type == GGML_TYPE_F16) { + compute_type = fast_bf16_hardware_available(cc) ? GGML_TYPE_BF16 : GGML_TYPE_F32; + } else if (prec == GGML_PREC_F32) { compute_type = GGML_TYPE_F32; } @@ -2033,6 +2037,7 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * ggml_tensor dst_slice; memset(&dst_slice, 0, sizeof(dst_slice)); + memcpy(dst_slice.op_params, dst->op_params, sizeof(dst_slice.op_params)); dst_slice.buffer = dst->buffer; dst_slice.type = type_dst_sorted; dst_slice.ne[0] = ne0;