Commit d6be06557 for llama.cpp
commit d6be065572cdb7bf64014c8c5771e210c9570e44
Author: Pascal <admin@serveurperso.com>
Date: Sun Nov 8 08:37:46 2026 +0100
llama, mtmd: follow the reference precision of ggml_prec_set_acc calls (#31851)
* llama, mtmd: follow the reference precision of ggml_prec_set_acc calls
Requests that the reference implementation does not keep in F32 move
to GGML_PREC_BF16: the GLM4 family down and output projections, the
MiMoVL vision matmuls and the Qwen3-TTS speaker encoder conv1d. The
DeepSeek V4/V4.1 routers and the MiniMax-M3 indexer stay in F32 with
a link to the reference line that computes them in F32.
* ggml: handle GGML_PREC_BF16 the same way as GGML_PREC_F32 in backends
Backends that reject GGML_PREC_F32 on the src1 of MUL_MAT_ID now also
reject GGML_PREC_BF16 (CUDA, Hexagon, OpenCL, WebGPU, Vulkan), and
backends that accept GGML_PREC_F32 for flash attention now also accept
GGML_PREC_BF16 with the same F32 accumulation (CPU, SpaceMIT, ET).
Vulkan keeps the F32 transfer between devices for both. The CUDA
cuBLAS path already handles GGML_PREC_BF16 and is unchanged.
* tests: cover GGML_PREC_BF16 in the flash attention tests
The flash attention cases run with GGML_PREC_BF16 next to F32 and
DEFAULT, and var_to_str prints bf16 so these cases keep their own name.
* llama: request GGML_PREC_BF16 for the KQ multiplication
* cont : add TODO in clip.cpp
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index 55f625557..750367c76 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -9591,6 +9591,7 @@ void ggml_compute_forward_flash_attn_ext(
switch (dst->op_params[3]) {
case GGML_PREC_DEFAULT:
case GGML_PREC_F32:
+ case GGML_PREC_BF16:
{
// uses F32 accumulators
ggml_compute_forward_flash_attn_ext_f16(params, dst);
diff --git a/ggml/src/ggml-cpu/spacemit/ime.cpp b/ggml/src/ggml-cpu/spacemit/ime.cpp
index 5f114f587..f6d047732 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_DEFAULT);
+ 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_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 9c467ce95..a49c5247a 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -5978,7 +5978,7 @@ 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_F32) {
+ 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)) {
return false;
}
#ifdef GGML_USE_MUSA
diff --git a/ggml/src/ggml-et/ggml-et-ops.cpp b/ggml/src/ggml-et/ggml-et-ops.cpp
index 876513867..6b8005042 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_DEFAULT) {
+ if (prec != GGML_PREC_F32 && prec != GGML_PREC_BF16 && prec != GGML_PREC_DEFAULT) {
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 466e68904..d88f21389 100644
--- a/ggml/src/ggml-et/ggml-et.cpp
+++ b/ggml/src/ggml-et/ggml-et.cpp
@@ -1299,7 +1299,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_DEFAULT) &&
+ supported = me_eligible && mask_ok && (prec == GGML_PREC_F32 || prec == GGML_PREC_BF16 || prec == GGML_PREC_DEFAULT) &&
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 76923d26b..2013f7ae2 100644
--- a/ggml/src/ggml-hexagon/htp-kparams.cpp
+++ b/ggml/src/ggml-hexagon/htp-kparams.cpp
@@ -944,7 +944,7 @@ static bool htp_matmul_check_types(
if (src1->type != GGML_TYPE_F32 || !src2 || src2->type != GGML_TYPE_I32) {
return false;
}
- if (ggml_get_op_params_i32(dst, 3) == GGML_PREC_F32) {
+ 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) {
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 771fe7751..a985e0882 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -12018,7 +12018,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
return false;
case GGML_OP_MUL_MAT_ID:
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
- if (ggml_get_op_params_i32(op, 3) == GGML_PREC_F32 &&
+ if ((ggml_get_op_params_i32(op, 3) == GGML_PREC_BF16 || ggml_get_op_params_i32(op, 3) == GGML_PREC_F32) &&
use_adreno_moe_kernels(backend_ctx, op->src[0])) {
return false;
}
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-comm.cpp b/ggml/src/ggml-vulkan/ggml-vulkan-comm.cpp
index a2370b8cd..90be9d889 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_F32;
+ return ggml_get_op_params_i32(t, 0) == GGML_PREC_BF16 || ggml_get_op_params_i32(t, 0) == 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-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
index 8347cd1b1..e5c90f6a7 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
@@ -15696,7 +15696,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
return false;
}
}
- if (ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
+ if (ggml_prec_get_src(op, 1) == GGML_PREC_BF16 || ggml_prec_get_src(op, 1) == GGML_PREC_F32) {
return false;
}
switch (src0_type) {
diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp
index 31d67fbf8..0d78a1b70 100644
--- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp
+++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp
@@ -4904,7 +4904,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_F32) {
+ if (ggml_get_op_params_i32(op, 3) == GGML_PREC_BF16 || ggml_get_op_params_i32(op, 3) == GGML_PREC_F32) {
supports_op = false;
}
break;
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index 9f99ba83e..b5d4642a2 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -2172,8 +2172,8 @@ ggml_tensor * llm_graph_context::build_ffn(
if (down) {
cur = build_lora_mm(down, cur);
if (arch == LLM_ARCH_GLM4 || arch == LLM_ARCH_GLM4_MOE || arch == LLM_ARCH_JAIS2) {
- // GLM4, GLM4_MOE, and JAIS2 seem to have numerical issues with half-precision accumulators
- ggml_prec_set_acc(cur, GGML_PREC_F32);
+ // GLM4, GLM4_MOE, and JAIS2 have numerical issues with F16 accumulators
+ ggml_prec_set_acc(cur, GGML_PREC_BF16);
}
}
@@ -2278,6 +2278,7 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
if (probs_in == nullptr) {
logits = build_lora_mm(gate_inp, cur); // [n_expert, n_tokens]
if (gating_op == LLAMA_EXPERT_GATING_FUNC_TYPE_SQRT_SOFTPLUS) {
+ // the reference computes the router in F32: https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash/blob/60d8d70770c6776ff598c94bb586a859a38244f1/inference/model.py#L565
ggml_prec_set_acc(logits, GGML_PREC_F32);
}
cb(logits, "ffn_moe_logits", il);
@@ -3112,9 +3113,7 @@ ggml_tensor * llm_graph_context::build_attn_mha(
ggml_tensor * kq = ggml_mul_mat(ctx0, k, q);
cb(kq, "kq", il);
- // note: this op tends to require high floating point range
- // while for some models F16 is enough, for others it is not, so we default to F32 here
- ggml_prec_set_acc(kq, GGML_PREC_F32);
+ ggml_prec_set_acc(kq, GGML_PREC_BF16);
if (model.arch == LLM_ARCH_GROK) {
// need to do the following:
@@ -3365,9 +3364,9 @@ ggml_tensor * llm_graph_context::build_attn(
if (wo) {
if (arch == LLM_ARCH_GLM4 || arch == LLM_ARCH_GLM4_MOE || arch == LLM_ARCH_JAIS2) {
- // GLM4, GLM4_MOE, and JAIS2 seem to have numerical issues with half-precision accumulators
+ // GLM4, GLM4_MOE, and JAIS2 have numerical issues with F16 accumulators
cur = build_lora_mm(wo, cur);
- ggml_prec_set_acc(cur, GGML_PREC_F32);
+ ggml_prec_set_acc(cur, GGML_PREC_BF16);
if (wo_s) {
cur = ggml_mul(ctx0, cur, wo_s);
}
@@ -3454,9 +3453,9 @@ ggml_tensor * llm_graph_context::build_attn(
if (wo) {
if (arch == LLM_ARCH_GLM4 || arch == LLM_ARCH_GLM4_MOE) {
- // GLM4 and GLM4_MOE seem to have numerical issues with half-precision accumulators
+ // GLM4 and GLM4_MOE have numerical issues with F16 accumulators
cur = build_lora_mm(wo, cur);
- ggml_prec_set_acc(cur, GGML_PREC_F32);
+ ggml_prec_set_acc(cur, GGML_PREC_BF16);
if (wo_s) {
cur = ggml_mul(ctx0, cur, wo_s);
}
diff --git a/src/models/deepseek41.cpp b/src/models/deepseek41.cpp
index ea24f6a7a..96de90412 100644
--- a/src/models/deepseek41.cpp
+++ b/src/models/deepseek41.cpp
@@ -875,7 +875,7 @@ llama_model_deepseek41::graph::graph(const llm_graph_params & params) :
cur = build_norm(cur, layer.ffn_norm, nullptr, LLM_NORM_RMS, il);
cb(cur, "ffn_norm", il);
- // F32 router accumulation: expert selection is near-tie sensitive
+ // the reference computes the router in F32: https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash-1017/blob/1bae0e1088993b9b82de08f59de5be5a34e42170/inference/model.py#L813
ggml_tensor * logits = build_lora_mm(layer.ffn_gate_inp, cur);
ggml_prec_set_acc(logits, GGML_PREC_F32);
cb(logits, "ffn_moe_logits", il);
diff --git a/src/models/minimax-m3.cpp b/src/models/minimax-m3.cpp
index 2002d08be..466985cc4 100644
--- a/src/models/minimax-m3.cpp
+++ b/src/models/minimax-m3.cpp
@@ -386,6 +386,7 @@ llama_model_minimax_m3::graph::graph(const llm_graph_params & params) : llm_grap
ggml_tensor * iq4 = ggml_reshape_4d(ctx0, iq, n_idx_dim, Hd, 1, ns);
ggml_tensor * sc = ggml_mul_mat(ctx0,
ggml_reshape_4d(ctx0, ikp, n_idx_dim, n_ps, 1, ns), iq4);
+ // the reference computes the indexer scores in F32: https://github.com/huggingface/transformers/blob/0f53828e83f2212060233d2ea552cc76617b8712/src/transformers/models/minimax_m3_vl/modeling_minimax_m3_vl.py#L562
ggml_prec_set_acc(sc, GGML_PREC_F32);
// unmapped positions come out -inf, so they can never rank into the top-k
sc = ggml_add_inplace(ctx0, sc,
@@ -467,7 +468,7 @@ llama_model_minimax_m3::graph::graph(const llm_graph_params & params) : llm_grap
ggml_tensor * ikp = ggml_get_rows(ctx0, ik_s, psl_s); // [n_idx_dim, n_ps]
ggml_tensor * sc = ggml_mul_mat(ctx0, ikp,
ggml_reshape_2d(ctx0, iq_s, n_idx_dim, Hd*n_tps));
- // indexer scores run in F32
+ // the reference computes the indexer scores in F32: https://github.com/huggingface/transformers/blob/0f53828e83f2212060233d2ea552cc76617b8712/src/transformers/models/minimax_m3_vl/modeling_minimax_m3_vl.py#L562
ggml_prec_set_acc(sc, GGML_PREC_F32);
sc = ggml_reshape_3d(ctx0, sc, n_ps, Hd, n_tps);
// unmapped positions (holes, padding, empty cells) come out -inf
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index cd9d2b894..f3f35329a 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -442,7 +442,11 @@ static std::string var_to_str(ggml_type type) {
}
static std::string var_to_str(ggml_prec prec) {
- return prec == GGML_PREC_F32 ? "f32" : "def";
+ switch (prec) {
+ case GGML_PREC_F32: return "f32";
+ case GGML_PREC_BF16: return "bf16";
+ default: return "def";
+ }
}
static std::string var_to_str(ggml_op_pool pool) {
@@ -12584,8 +12588,8 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
for (int kv : { 113, 512, 1024, }) {
if (nr2 != 1 && kv != 512) continue;
for (int nb : { 1, 3, 32, 75, }) {
- for (ggml_prec prec : {GGML_PREC_F32, GGML_PREC_DEFAULT}) {
- if (hsk != 128 && prec == GGML_PREC_DEFAULT) continue;
+ for (ggml_prec prec : {GGML_PREC_F32, GGML_PREC_BF16, GGML_PREC_DEFAULT}) {
+ if (hsk != 128 && prec != GGML_PREC_F32) continue;
for (ggml_type type_KV : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16, GGML_TYPE_Q8_0, GGML_TYPE_Q5_1, GGML_TYPE_Q5_0, GGML_TYPE_Q4_1, GGML_TYPE_Q4_0, GGML_TYPE_IQ4_NL}) {
if (type_KV != GGML_TYPE_F16 && hsk != 64 && hsk != 72) continue;
// DeepSeek MLA: the V cache is a sub-view of the K cache
diff --git a/tools/mtmd/clip.cpp b/tools/mtmd/clip.cpp
index ac99b762e..0a3b3c98d 100644
--- a/tools/mtmd/clip.cpp
+++ b/tools/mtmd/clip.cpp
@@ -875,8 +875,8 @@ ggml_tensor * clip_graph::build_attn(
v = ggml_cont(ctx0, v);
ggml_tensor * kq = ggml_mul_mat(ctx0, k, q);
- // F32 may not needed for vision encoders?
- // ggml_prec_set_acc(kq, GGML_PREC_F32);
+ // TODO: this most likely requires GGML_PREC_BF16, but will enable it later to avoid scope-creep
+ // ggml_prec_set_acc(kq, GGML_PREC_BF16);
kq = ggml_soft_max_ext(ctx0, kq, kq_mask, kq_scale, 0.0f);
if (sinks != nullptr) {
diff --git a/tools/mtmd/models/mimovl.cpp b/tools/mtmd/models/mimovl.cpp
index 0cde450bf..4843698e6 100644
--- a/tools/mtmd/models/mimovl.cpp
+++ b/tools/mtmd/models/mimovl.cpp
@@ -2,7 +2,7 @@
ggml_tensor * clip_graph_mimovl::build_mm(ggml_tensor * w, ggml_tensor * x) const {
ggml_tensor * cur = ggml_mul_mat(ctx0, w, x);
- ggml_prec_set_acc(cur, GGML_PREC_F32);
+ ggml_prec_set_acc(cur, GGML_PREC_BF16);
return cur;
}
diff --git a/tools/mtmd/models/qwen3tts-spkenc.cpp b/tools/mtmd/models/qwen3tts-spkenc.cpp
index 405fbb9cb..ac6cafd9a 100644
--- a/tools/mtmd/models/qwen3tts-spkenc.cpp
+++ b/tools/mtmd/models/qwen3tts-spkenc.cpp
@@ -27,7 +27,7 @@ ggml_tensor * clip_graph_qwen3tts_spkenc::conv1d_same(ggml_tensor * x, ggml_tens
ggml_tensor * w2d = ggml_reshape_2d(ctx0, w, (int64_t) K * IC, OC);
ggml_tensor * y = ggml_mul_mat(ctx0, w2d, col); // [OC, T_out]
- ggml_prec_set_acc(y, GGML_PREC_F32);
+ ggml_prec_set_acc(y, GGML_PREC_BF16);
ggml_tensor * b2d = ggml_reshape_2d(ctx0, b, OC, 1);
y = ggml_add(ctx0, y, b2d);