Commit 494d0f47b for llama.cpp
commit 494d0f47bde61fd45018acab8334cdfd0d8b3ca0
Author: Georgi Gerganov <ggerganov@gmail.com>
Date: Tue Nov 10 17:03:27 2026 +0200
ggml : add MUL_MAT src1 precision tests and gating (#31896)
diff --git a/ggml/src/ggml-cann/ggml-cann.cpp b/ggml/src/ggml-cann/ggml-cann.cpp
index 8dc8759aa..b49aee73d 100644
--- a/ggml/src/ggml-cann/ggml-cann.cpp
+++ b/ggml/src/ggml-cann/ggml-cann.cpp
@@ -2447,6 +2447,9 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
break;
case GGML_OP_MUL_MAT:
{
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ return false;
+ }
switch (op->src[0]->type) {
#ifndef ASCEND_310P
case GGML_TYPE_BF16:
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index 1884600e6..00a814051 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -9569,8 +9569,8 @@ static void ggml_compute_forward_flash_attn_ext_f16(
void ggml_compute_forward_flash_attn_ext(
const ggml_compute_params * params,
ggml_tensor * dst) {
- switch (dst->op_params[3]) {
- case GGML_PREC_DEFAULT:
+ switch (ggml_prec_get_acc(dst)) {
+ case GGML_PREC_UNDEFINED:
case GGML_PREC_F32:
case GGML_PREC_BF16:
{
diff --git a/ggml/src/ggml-cpu/spacemit/ime.cpp b/ggml/src/ggml-cpu/spacemit/ime.cpp
index d655cdafc..6cf243058 100644
--- a/ggml/src/ggml-cpu/spacemit/ime.cpp
+++ b/ggml/src/ggml-cpu/spacemit/ime.cpp
@@ -1180,7 +1180,7 @@ class tensor_traits_common : public tensor_traits_base {
const int64_t DK = nek0;
const int64_t DV = nev0;
- const bool supported_prec = (dst->op_params[3] == GGML_PREC_F32 || dst->op_params[3] == GGML_PREC_BF16 || dst->op_params[3] == GGML_PREC_DEFAULT);
+ const bool supported_prec = (ggml_prec_get_acc(dst) == GGML_PREC_F32 || ggml_prec_get_acc(dst) == GGML_PREC_BF16 || ggml_prec_get_acc(dst) == GGML_PREC_UNDEFINED);
const bool supported_types = (q->type == GGML_TYPE_F32 && k->type == GGML_TYPE_F16 && v->type == GGML_TYPE_F16);
const bool supported_shape = (DK > 0 && DK <= 128 && DV > 0 && DV <= 128);
const bool supported_vlen = (__riscv_vlenb() == 128);
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 8826d3372..848565460 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -1998,7 +1998,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);
+ const ggml_prec prec = ggml_prec_get_acc(dst);
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;
@@ -6109,7 +6109,13 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
if (b->type == GGML_TYPE_F16 && a->type != GGML_TYPE_F16 && !ggml_cuda_op_mul_mat_use_fwht(op)) {
return false;
}
- if (op->op == GGML_OP_MUL_MAT_ID && (ggml_get_op_params_i32(op, 3) == GGML_PREC_BF16 || ggml_get_op_params_i32(op, 3) == GGML_PREC_F32)) {
+ const int cc_dev = ggml_cuda_info().devices[dev_ctx->device].cc;
+ const bool f32_src1_ok = GGML_CUDA_CC_IS_NVIDIA(cc_dev) && cc_dev < GGML_CUDA_CC_AMPERE;
+ const enum ggml_prec prec_src1 = ggml_prec_get_src(op, 1);
+ if (prec_src1 == GGML_PREC_F32 && (op->op == GGML_OP_MUL_MAT_ID || !f32_src1_ok || a->type != GGML_TYPE_F32)) {
+ return false;
+ }
+ if (prec_src1 == GGML_PREC_BF16 && (op->op == GGML_OP_MUL_MAT_ID || a->type != GGML_TYPE_F32)) {
return false;
}
#ifdef GGML_USE_MUSA
diff --git a/ggml/src/ggml-cuda/mmvf.cu b/ggml/src/ggml-cuda/mmvf.cu
index a62bf53e8..6008305ef 100644
--- a/ggml/src/ggml-cuda/mmvf.cu
+++ b/ggml/src/ggml-cuda/mmvf.cu
@@ -616,7 +616,7 @@ static void mul_mat_vec_f_cuda(
const int64_t ids_stride, enum ggml_prec prec, cudaStream_t stream) {
if constexpr(std::is_same_v<T, half>) {
- if (prec == GGML_PREC_DEFAULT) {
+ if (prec == GGML_PREC_UNDEFINED) {
mul_mat_vec_f_cuda_switch_ncols_dst<T, half>
(x, y, ids, fusion, dst, ncols, nrows, ncols_dst, stride_row, stride_col_y, stride_col_dst,
nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y,
@@ -651,7 +651,8 @@ void ggml_cuda_mul_mat_vec_f(ggml_backend_cuda_context & ctx, const ggml_tensor
GGML_ASSERT( nb0 == ts_dst);
const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
- const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec(dst->op_params[0]) : GGML_PREC_F32;
+ // the fused GLU path gets the GLU node as dst, which carries no matmul precision, so it accumulates in F32
+ const enum ggml_prec prec = fast_fp16_available(cc) && !(fusion && fusion->gate) ? ggml_prec_get_acc(dst) : GGML_PREC_F32;
const float * src1_d = (const float *) src1->data;
const int32_t * ids_d = ids ? (const int32_t *) ids->data : nullptr;
@@ -743,7 +744,7 @@ void ggml_cuda_op_mul_mat_vec_f(
const int id = ggml_cuda_get_device();
const int cc = ggml_cuda_info().devices[id].cc;
- const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec(dst->op_params[0]) : GGML_PREC_F32;
+ const enum ggml_prec prec = fast_fp16_available(cc) ? ggml_prec_get_acc(dst) : GGML_PREC_F32;
// ggml_cuda_op provides single, contiguous matrices
const int64_t stride_row = ne00;
diff --git a/ggml/src/ggml-et/ggml-et-ops.cpp b/ggml/src/ggml-et/ggml-et-ops.cpp
index 6b8005042..c51d16f82 100644
--- a/ggml/src/ggml-et/ggml-et-ops.cpp
+++ b/ggml/src/ggml-et/ggml-et-ops.cpp
@@ -1523,7 +1523,7 @@ bool ggml_et_op_flash_attn_ext(ggml_backend_et_device_context * dev_ctx, const g
}
const ggml_prec prec = ggml_flash_attn_ext_get_prec(node);
- if (prec != GGML_PREC_F32 && prec != GGML_PREC_BF16 && prec != GGML_PREC_DEFAULT) {
+ if (prec != GGML_PREC_F32 && prec != GGML_PREC_BF16 && prec != GGML_PREC_UNDEFINED) {
GGML_LOG_ERROR("ET: FLASH_ATTN_EXT baseline kernel only supports F32 precision\n");
return false;
}
diff --git a/ggml/src/ggml-et/ggml-et.cpp b/ggml/src/ggml-et/ggml-et.cpp
index d88f21389..53e64b323 100644
--- a/ggml/src/ggml-et/ggml-et.cpp
+++ b/ggml/src/ggml-et/ggml-et.cpp
@@ -941,6 +941,9 @@ static bool ggml_backend_et_device_supports_op(ggml_backend_dev_t dev, const ggm
op->nb[1] == op->ne[0] * sizeof(float);
break;
case GGML_OP_MUL_MAT:
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ break;
+ }
// Support Q8_0 x F32 -> F32, F16 x F32 -> F32, F16 x F16 -> F32, and F32 x F32 -> F32 matrix multiplication
// Stride requirements: first dimension must be contiguous for all tensors
if (op->type == GGML_TYPE_F32 &&
@@ -1299,7 +1302,7 @@ static bool ggml_backend_et_device_supports_op(ggml_backend_dev_t dev, const ggm
const bool me_eligible = op->src[1]->type == GGML_TYPE_F16 && op->src[2]->type == GGML_TYPE_F16 &&
(op->src[0]->ne[0] % 32) == 0;
- supported = me_eligible && mask_ok && (prec == GGML_PREC_F32 || prec == GGML_PREC_BF16 || prec == GGML_PREC_DEFAULT) &&
+ supported = me_eligible && mask_ok && (prec == GGML_PREC_F32 || prec == GGML_PREC_BF16 || prec == GGML_PREC_UNDEFINED) &&
max_bias == 0.0f && logit_softcap == 0.0f && op->src[0]->nb[0] == sizeof(float) &&
op->src[1]->nb[0] == k_elem && op->src[2]->nb[0] == v_elem && op->nb[0] == sizeof(float) &&
op->src[0]->ne[0] == op->src[1]->ne[0] && // dk matches
diff --git a/ggml/src/ggml-hexagon/htp-kparams.cpp b/ggml/src/ggml-hexagon/htp-kparams.cpp
index 2013f7ae2..2fbd419ac 100644
--- a/ggml/src/ggml-hexagon/htp-kparams.cpp
+++ b/ggml/src/ggml-hexagon/htp-kparams.cpp
@@ -939,14 +939,15 @@ static bool htp_matmul_check_types(
return false;
}
+ if (ggml_prec_get_src(dst, 1) == GGML_PREC_BF16 || ggml_prec_get_src(dst, 1) == GGML_PREC_F32) {
+ return false;
+ }
+
if (is_matmul_id) {
const struct ggml_tensor * src2 = dst->src[2];
if (src1->type != GGML_TYPE_F32 || !src2 || src2->type != GGML_TYPE_I32) {
return false;
}
- if (ggml_get_op_params_i32(dst, 3) == GGML_PREC_BF16 || ggml_get_op_params_i32(dst, 3) == GGML_PREC_F32) {
- return false;
- }
switch (src0->type) {
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q2_0:
diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m
index 5e6099ce1..0a07a76bf 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.m
+++ b/ggml/src/ggml-metal/ggml-metal-device.m
@@ -1510,6 +1510,18 @@ void ggml_metal_device_get_memory(ggml_metal_device_t dev, size_t * free, size_t
}
}
+static bool ggml_metal_supports_mul_mat_src1_prec(const struct ggml_tensor * op, bool has_tensor) {
+ const enum ggml_prec prec = ggml_prec_get_src(op, 1);
+
+ if (prec == GGML_PREC_F32) {
+ return false;
+ }
+ if (prec == GGML_PREC_BF16) {
+ return has_tensor && op->op == GGML_OP_MUL_MAT && op->src[0]->type != GGML_TYPE_F32;
+ }
+ return true;
+}
+
static bool ggml_metal_supports_mul_mat_op(
bool has_simdgroup_reduction,
const struct ggml_tensor * op,
@@ -1894,6 +1906,9 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
case GGML_OP_SOLVE_TRI:
return has_simdgroup_reduction && op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_MUL_MAT:
+ if (!ggml_metal_supports_mul_mat_src1_prec(op, dev->props.has_tensor)) {
+ return false;
+ }
// the FWHT kernels read an F16 source directly; every other F16 src1 path
// still goes through ggml_metal_supports_mul_mat_op
if (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F16 &&
@@ -1904,6 +1919,9 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
has_simdgroup_reduction, op, true,
ggml_metal_op_mul_mat_use_mm(op, has_simdgroup_mm));
case GGML_OP_MUL_MAT_ID:
+ if (!ggml_metal_supports_mul_mat_src1_prec(op, dev->props.has_tensor)) {
+ return false;
+ }
return ggml_metal_supports_mul_mat_op(
has_simdgroup_reduction, op, false,
ggml_metal_op_mul_mat_id_use_mm(op, has_simdgroup_mm));
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index a4f9ae272..c25a13dc6 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -11937,6 +11937,9 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
case GGML_OP_GROUP_NORM:
return ggml_is_contiguous(op->src[0]);
case GGML_OP_MUL_MAT:
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ return false;
+ }
if (ggml_cl_mul_mat_can_use_fwht(backend_ctx, op)) {
return true;
}
@@ -12056,7 +12059,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
op->src[1]->ne[1] <= INT_MAX && op->ne[1] <= INT_MAX && op->ne[2] <= INT_MAX - 64;
}
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
- if ((ggml_get_op_params_i32(op, 3) == GGML_PREC_BF16 || ggml_get_op_params_i32(op, 3) == GGML_PREC_F32) &&
+ if ((ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) &&
use_adreno_moe_kernels(backend_ctx, op->src[0])) {
return false;
}
diff --git a/ggml/src/ggml-openvino/ggml-openvino.cpp b/ggml/src/ggml-openvino/ggml-openvino.cpp
index 88634557c..3aff66724 100644
--- a/ggml/src/ggml-openvino/ggml-openvino.cpp
+++ b/ggml/src/ggml-openvino/ggml-openvino.cpp
@@ -1581,6 +1581,9 @@ static ggml_openvino_op_support is_op_supported_case(const ggml_tensor * op) {
break;
}
case GGML_OP_MUL_MAT: {
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ return {false, "MUL_MAT with a BF16 or F32 src1 precision is not supported"};
+ }
if (ggml_openvino_is_gpu() && op->src[0] != nullptr && op->src[1] != nullptr &&
ggml_is_quantized(op->src[0]->type) && strcmp(op->src[0]->name, "a") == 0 &&
strcmp(op->src[1]->name, "b") == 0 && op->src[0]->ne[1] == 1 && op->src[1]->ne[1] == 64 &&
diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index a572dd7fe..99da9713e 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -129,7 +129,7 @@ extern int g_ggml_sycl_dynamic_required_precision;
// [TAG_GGML_PREC] src1 precision request of the MUL_MAT/MUL_MAT_ID op dst
static inline int32_t ggml_sycl_src1_prec(const ggml_tensor * dst) {
return g_ggml_sycl_dynamic_required_precision == GGML_SYCL_DYNAMIC_PRECISION_F16 ? GGML_PREC_UNDEFINED :
- dst->op_params[3];
+ ggml_prec_get_src(dst, 1);
}
// [TAG_GGML_PREC] the library GEMM and dmmv may convert src1 of the MUL_MAT/MUL_MAT_ID op dst to f16
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 7e6df8fdd..f634a7f1c 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -3315,7 +3315,7 @@ inline void ggml_sycl_op_mul_mat_sycl(
// the f16 route converts src1 to f16 [TAG_GGML_PREC]
use_fp16 = use_fp16 && ggml_sycl_src1_f16_ok(dst);
if ((src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && use_fp16 && ggml_is_contiguous(src0) &&
- row_diff == src0->ne[1] && dst->op_params[0] == GGML_PREC_DEFAULT) {
+ row_diff == src0->ne[1] && ggml_prec_get_acc(dst) == GGML_PREC_UNDEFINED) {
ggml_sycl_pool_alloc<sycl::half> src1_as_f16(ctx.pool());
if (src1->type != GGML_TYPE_F16) {
scope_op_debug_print scope_dbg_print(__func__, "/to_fp16_sycl", dst, /*num_src=*/2,
@@ -6929,7 +6929,7 @@ static bool mul_mat_uses_onednn(int device, ggml_tensor * dst) {
#endif
#ifdef GGML_SYCL_F16
if ((src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && ggml_is_contiguous(src0) &&
- dst->op_params[0] == GGML_PREC_DEFAULT) {
+ ggml_prec_get_acc(dst) == GGML_PREC_UNDEFINED) {
return true;
}
#endif
@@ -7429,6 +7429,11 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
return false;
}
+ if (op->op == GGML_OP_MUL_MAT &&
+ (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32)) {
+ return false;
+ }
+
return true;
}
case GGML_OP_OUT_PROD:
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-comm.cpp b/ggml/src/ggml-vulkan/ggml-vulkan-comm.cpp
index 90be9d889..84af0661a 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan-comm.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan-comm.cpp
@@ -18,7 +18,7 @@ static bool ggml_vk_comm_needs_f32(const ggml_tensor * t, int depth = 0) {
switch (t->op) {
case GGML_OP_MUL_MAT:
case GGML_OP_MUL_MAT_ID:
- return ggml_get_op_params_i32(t, 0) == GGML_PREC_BF16 || ggml_get_op_params_i32(t, 0) == GGML_PREC_F32;
+ return ggml_prec_get_acc(t) == GGML_PREC_BF16 || ggml_prec_get_acc(t) == GGML_PREC_F32;
case GGML_OP_ADD:
return ggml_vk_comm_needs_f32(t->src[0], depth + 1) || ggml_vk_comm_needs_f32(t->src[1], depth + 1);
case GGML_OP_MUL:
diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp
index bfc0b4c08..b827e3d42 100644
--- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp
+++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp
@@ -4910,6 +4910,9 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
default:
break;
}
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ supports_op = false;
+ }
break;
}
case GGML_OP_MUL_MAT_ID:
@@ -4953,7 +4956,7 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
default:
break;
}
- if (ggml_get_op_params_i32(op, 3) == GGML_PREC_BF16 || ggml_get_op_params_i32(op, 3) == GGML_PREC_F32) {
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
supports_op = false;
}
break;
diff --git a/ggml/src/ggml-zdnn/ggml-zdnn.cpp b/ggml/src/ggml-zdnn/ggml-zdnn.cpp
index 533aefd9c..ccdd4f382 100644
--- a/ggml/src/ggml-zdnn/ggml-zdnn.cpp
+++ b/ggml/src/ggml-zdnn/ggml-zdnn.cpp
@@ -119,6 +119,10 @@ static bool ggml_zdnn_supports_op(const ggml_backend_zdnn_device_context * ctx_d
const int64_t max_batch = ctx_dev->max_size;
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ return false;
+ }
+
if (!ggml_is_matrix(weights) || !ggml_is_matrix(inputs) ||
!ggml_is_contiguous(weights) || !ggml_is_contiguous(inputs) ||
weights->view_src != nullptr ||
diff --git a/ggml/src/ggml-zendnn/ggml-zendnn.cpp b/ggml/src/ggml-zendnn/ggml-zendnn.cpp
index 357cfcd99..85f1263fb 100644
--- a/ggml/src/ggml-zendnn/ggml-zendnn.cpp
+++ b/ggml/src/ggml-zendnn/ggml-zendnn.cpp
@@ -720,6 +720,10 @@ static bool ggml_backend_zendnn_device_supports_op(ggml_backend_dev_t dev, const
return false;
}
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ return false;
+ }
+
if (ggml_zendnn_adaptive_fallback_enabled()) {
const int64_t K = inputs->ne[0];
const int64_t N = (inputs->ne[1]*inputs->ne[2]*inputs->ne[3]);
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index b1a538a31..676aa9ed2 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -5677,6 +5677,91 @@ struct test_mul_mat_w4a8 : public test_mul_mat {
}
};
+// src1 GGML_PREC_BF16: the implementation may keep src1 in F32 or BF16, but must not convert it down to F16,
+// which cannot hold the src1 values that the BF16 range allows
+struct test_mul_mat_prec_src1_bf16 : public test_mul_mat {
+ test_mul_mat_prec_src1_bf16(ggml_type type_a = GGML_TYPE_F32, int64_t m = 64, int64_t n = 32, int64_t k = 256,
+ std::array<int64_t, 2> bs = {1, 1},
+ std::array<int64_t, 2> nr = {1, 1})
+ : test_mul_mat(type_a, GGML_TYPE_F32, m, n, k, bs, nr) {}
+ ggml_tensor * build_graph(ggml_context * ctx) override {
+ ggml_tensor * out = test_mul_mat::build_graph(ctx);
+ for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+ if (t->op == GGML_OP_MUL_MAT) {
+ ggml_prec_set_src(t, GGML_PREC_BF16, 1);
+ }
+ }
+ return out;
+ }
+ void initialize_tensors(ggml_context * ctx) override {
+ for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+ if (strcmp(t->name, "b") == 0) {
+ init_tensor_uniform(t, -1e6f, 1e6f);
+ } else {
+ init_tensor_uniform(t);
+ }
+ }
+ }
+ std::string op_desc(ggml_tensor * t) override {
+ GGML_UNUSED(t);
+ return "MUL_MAT_PREC_SRC1_BF16";
+ }
+};
+
+// src1 GGML_PREC_F32: the implementation can not convert src1 or the F32 src0 tiles down to F16 or TF32
+struct test_mul_mat_prec_src1_f32 : public test_mul_mat {
+ test_mul_mat_prec_src1_f32(ggml_type type_a = GGML_TYPE_F32, int64_t m = 64, int64_t n = 32, int64_t k = 256,
+ std::array<int64_t, 2> bs = {1, 1},
+ std::array<int64_t, 2> nr = {1, 1})
+ : test_mul_mat(type_a, GGML_TYPE_F32, m, n, k, bs, nr) {}
+ double max_nmse_err() override {
+ return 1e-10;
+ }
+ ggml_tensor * build_graph(ggml_context * ctx) override {
+ ggml_tensor * out = test_mul_mat::build_graph(ctx);
+ for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+ if (t->op == GGML_OP_MUL_MAT) {
+ ggml_prec_set_src(t, GGML_PREC_F32, 1);
+ }
+ }
+ return out;
+ }
+ void initialize_tensors(ggml_context * ctx) override {
+ for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+ if (strcmp(t->name, "b") == 0) {
+ init_tensor_uniform(t, -1e6f, 1e6f);
+ } else {
+ init_tensor_uniform(t);
+ }
+ }
+ }
+ std::string op_desc(ggml_tensor * t) override {
+ GGML_UNUSED(t);
+ return "MUL_MAT_PREC_SRC1_F32";
+ }
+};
+
+// src1 GGML_PREC_BF16 with src0 values outside the F16 range: the F32 src0 tiles must not be downcast to F16
+struct test_mul_mat_prec_src1_bf16_src0 : public test_mul_mat_prec_src1_bf16 {
+ test_mul_mat_prec_src1_bf16_src0(ggml_type type_a = GGML_TYPE_F32, int64_t m = 64, int64_t n = 32, int64_t k = 256,
+ std::array<int64_t, 2> bs = {1, 1},
+ std::array<int64_t, 2> nr = {1, 1})
+ : test_mul_mat_prec_src1_bf16(type_a, m, n, k, bs, nr) {}
+ void initialize_tensors(ggml_context * ctx) override {
+ for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+ if (strcmp(t->name, "a") == 0) {
+ init_tensor_uniform(t, -1e6f, 1e6f);
+ } else {
+ init_tensor_uniform(t);
+ }
+ }
+ }
+ std::string op_desc(ggml_tensor * t) override {
+ GGML_UNUSED(t);
+ return "MUL_MAT_PREC_SRC1_BF16_SRC0";
+ }
+};
+
// FP4 native W4A4 path (default precision)
struct test_mul_mat_w4a4 : public test_mul_mat {
test_mul_mat_w4a4(ggml_type type_a = GGML_TYPE_NVFP4, ggml_type type_b = GGML_TYPE_F32,
@@ -11634,6 +11719,36 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 8192, 1, 8192));
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F16, 1024, 7, 1024)); // many rows
+ // src1 precision: few-row paths, aligned tiles, partial tiles (K and output), broadcast batches
+ for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q8_0}) {
+ // Q8_0 requires K to be a multiple of its block size (32)
+ const int64_t k_partial = type_a == GGML_TYPE_Q8_0 ? 96 : 100;
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16(type_a, 64, 1, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16(type_a, 64, 8, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16(type_a, 64, 32, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16(type_a, 128, 512, 512));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16(type_a, 67, 33, k_partial));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16(type_a, 64, 32, 256, {2, 3}, {2, 1}));
+ }
+
+ // src1 precision: same shapes with GGML_PREC_F32, which the backends must not round down to F16 or TF32
+ for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16}) {
+ test_cases.emplace_back(new test_mul_mat_prec_src1_f32(type_a, 64, 1, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_f32(type_a, 64, 8, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_f32(type_a, 64, 32, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_f32(type_a, 128, 512, 512));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_f32(type_a, 67, 33, 100));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_f32(type_a, 64, 32, 256, {2, 3}, {2, 1}));
+ }
+
+ // src0 range: an F32 src0 must not be downcast to F16 when src1 requests BF16
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16_src0(GGML_TYPE_F32, 64, 1, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16_src0(GGML_TYPE_F32, 64, 8, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16_src0(GGML_TYPE_F32, 64, 32, 256));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16_src0(GGML_TYPE_F32, 128, 512, 512));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16_src0(GGML_TYPE_F32, 67, 33, 100));
+ test_cases.emplace_back(new test_mul_mat_prec_src1_bf16_src0(GGML_TYPE_F32, 64, 32, 256, {2, 3}, {2, 1}));
+
// FP4 activation precision (default = native W4A4, src1 GGML_PREC_Q8 = W4A8)
test_cases.emplace_back(new test_mul_mat_w4a8(GGML_TYPE_NVFP4, GGML_TYPE_F32, 32, 1, 256));
test_cases.emplace_back(new test_mul_mat_w4a8(GGML_TYPE_NVFP4, GGML_TYPE_F32, 32, 32, 256));
diff --git a/tools/mtmd/mtmd-subcontext.cpp b/tools/mtmd/mtmd-subcontext.cpp
index f363cba95..ae6db71d1 100644
--- a/tools/mtmd/mtmd-subcontext.cpp
+++ b/tools/mtmd/mtmd-subcontext.cpp
@@ -271,46 +271,6 @@ mtmd_subcontext::support_info_graph mtmd_subcontext::reserve_compute_meta(const
return res;
}
-// TODO: remove this and use ggml_mul_mat with GGML_PREC_F32 once all backends honor it for F32 x F32
-// ggml_mul_mat for 2D F32 tensors, always in F32 on CPU: some GPU backends round the inputs to F16 or TF32, which loses the quiet bins of a spectrum
-static void mtmd_mul_mat_prec_op(ggml_tensor * dst, int ith, int nth, void * /*userdata*/) {
- const ggml_tensor * a = dst->src[0];
- const ggml_tensor * b = dst->src[1];
- const int64_t n_k = a->ne[0];
-
- for (int64_t n = ith; n < b->ne[1]; n += nth) {
- const float * b_row = (const float *) ((const char *) b->data + n * b->nb[1]);
- float * dst_row = (float *) ((char *) dst->data + n * dst->nb[1]);
- for (int64_t m = 0; m < a->ne[1]; m++) {
- const float * a_row = (const float *) ((const char *) a->data + m * a->nb[1]);
- // independent partial sums, so that the compiler can vectorize the loop
- float acc[16] = {};
- int64_t k = 0;
- for (; k + 16 <= n_k; k += 16) {
- for (int l = 0; l < 16; l++) {
- acc[l] += a_row[k + l] * b_row[k + l];
- }
- }
- float sum = 0.0f;
- for (; k < n_k; k++) {
- sum += a_row[k] * b_row[k];
- }
- for (int l = 0; l < 16; l++) {
- sum += acc[l];
- }
- dst_row[m] = sum;
- }
- }
-}
-
-static ggml_tensor * mtmd_mul_mat_prec(ggml_context * ctx, ggml_tensor * a, ggml_tensor * b) {
- GGML_ASSERT(a->type == GGML_TYPE_F32 && b->type == GGML_TYPE_F32);
- GGML_ASSERT(ggml_n_dims(a) <= 2 && ggml_n_dims(b) <= 2 && a->ne[0] == b->ne[0]);
- GGML_ASSERT(a->nb[0] == sizeof(float) && b->nb[0] == sizeof(float));
- ggml_tensor * args[] = { a, b };
- return ggml_custom_4d(ctx, GGML_TYPE_F32, a->ne[1], b->ne[1], 1, 1, args, 2, mtmd_mul_mat_prec_op, GGML_N_TASKS_MAX, nullptr);
-}
-
bool mtmd_audio_log_mel(const mtmd_audio_mel_config & config, const float * samples, const float * window, int64_t n_frames, float * out) {
ggml_backend_sched_t sched = config.sched;
ggml_backend_t backend_cpu = ggml_backend_sched_get_backend(sched, ggml_backend_sched_get_n_backends(sched) - 1);
@@ -356,7 +316,8 @@ bool mtmd_audio_log_mel(const mtmd_audio_mel_config & config, const float * samp
ggml_set_input(inp_filters);
ggml_tensor * cur = ggml_cont(ctx0, ggml_view_2d(ctx0, inp_samples, frame_size, n_cur, config.hop_length * ggml_element_size(inp_samples), 0));
- cur = mtmd_mul_mat_prec(ctx0, inp_dft, cur); // [2 * n_fft_bins, n_cur]
+ cur = ggml_mul_mat(ctx0, inp_dft, cur); // [2 * n_fft_bins, n_cur]
+ ggml_prec_set_src(cur, GGML_PREC_F32, 1);
cur = ggml_sqr(ctx0, cur);
cur = ggml_add(ctx0,
@@ -366,7 +327,8 @@ bool mtmd_audio_log_mel(const mtmd_audio_mel_config & config, const float * samp
cur = ggml_sqrt(ctx0, cur);
}
- cur = mtmd_mul_mat_prec(ctx0, cur, inp_filters); // [n_cur, n_mel]
+ cur = ggml_mul_mat(ctx0, cur, inp_filters); // [n_cur, n_mel]
+ ggml_prec_set_src(cur, GGML_PREC_F32, 1);
cur = config.mel_floor_add
? ggml_scale_bias(ctx0, cur, 1.0f, config.mel_floor)
: ggml_clamp(ctx0, cur, config.mel_floor, FLT_MAX);