public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-22 03:48:19 +02:00
Code Issues Packages Projects Releases Wiki Activity

CUDA: Handle compute type for NVFP4 on cublass path (llama/29173)

Browse Source 
* CUDA: Handle compute type for NVFP4 on cublass path

Signed-off-by: ynankani <ynankani@nvidia.com>

* Use BF16 compute type for quantized models if HW allows

Signed-off-by: ynankani <ynankani@nvidia.com>

* Set acc prec to bf16 for nvfp4 as it needs atleast bf16 range

Signed-off-by: ynankani <ynankani@nvidia.com>

* Update ggml/src/ggml-cuda/ggml-cuda.cu

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>

* preserve op_params for per-expert matmul

Signed-off-by: ynankani <ynankani@nvidia.com>

---------

Signed-off-by: ynankani <ynankani@nvidia.com>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
This commit is contained in:
ynankani
2026-10-06 10:38:14 +03:00
committed by Georgi Gerganov
co-authored by Johannes Gäßler
parent 3b68f9014c
commit 10872be7f8
1 changed files with 6 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/ggml-cuda.cu
+6 -1
View File
@@ -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) { 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 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; ggml_type compute_type = src0->type;
if (ggml_is_quantized(compute_type)) { if (ggml_is_quantized(compute_type)) {
compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32; 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; 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; 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; ggml_tensor dst_slice;
memset(&dst_slice, 0, sizeof(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.buffer = dst->buffer;
dst_slice.type = type_dst_sorted; dst_slice.type = type_dst_sorted;
dst_slice.ne[0] = ne0; dst_slice.ne[0] = ne0;