llama, mtmd: follow the reference precision of ggml_prec_set_acc calls

@@ -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);

@@ -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);

@@ -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

@@ -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;     }

@@ -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

@@ -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) {

@@ -12002,7 +12002,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;             }

@@ -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:

@@ -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) {

@@ -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;

@@ -2169,8 +2169,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);         }     } @@ -2275,6 +2275,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);@@ -3084,9 +3085,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:@@ -3337,9 +3336,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);             }@@ -3426,9 +3425,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);             }

@@ -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);

@@ -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

@@ -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

@@ -831,8 +831,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) {

@@ -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; } 

@@ -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);