model : DeepSeek-V4.1-Flash-1017 pin on upstream deepseek41 (MXFP8 Engram + extras)

common/fit.cpp

@@ -11,6 +11,8 @@
 #include <cassert>
 #include <stdexcept>
 #include <cinttypes>
+#include <memory>
+#include <cstdlib>
 #include <set>
 #include <string>
 #include <vector>
@@ -81,6 +83,26 @@ static std::vector<llama_device_memory_data> common_get_device_memory_data_impl(
         throw std::runtime_error("failed to create llama_context from model");
     }
 
+    // a draft measured next to its target (ctx_other set by common_fit_draft_target): the speculative decoder turns on
+    // nextn embeddings for DFlash/DSpark/Eagle3 drafts after creating them and installs the backend samplers only on
+    // those (not on DFlash2 / LiliCorr), so do the same and reserve again: the breakdown below then covers the compute
+    // graph the draft actually runs with
+    if (cparams->ctx_other != nullptr) {
+        const bool spec_draft = llama_model_target_layer_ids_n(model) > 0;
+        char buf[32] = {};
+        const bool is_lilicorr = llama_model_meta_val_str(model, "lilicorr.candidate_topk", buf, sizeof(buf)) >= 0 && std::atoi(buf) > 0;
+        const bool is_dflash2  = llama_model_dflash_selector_top_k(model) > 0;
+        if (!spec_draft || is_dflash2 || is_lilicorr) {
+            for (size_t i = 0; i < cparams->n_samplers; i++) {
+                llama_set_sampler(ctx, cparams->samplers[i].seq_id, nullptr);
+            }
+        }
+        if (spec_draft) {
+            llama_set_embeddings_nextn(ctx, true, /*masked*/ !is_dflash2 && !is_lilicorr);
+        }
+        llama_sched_reserve(ctx);
+    }
+
     const size_t nd = llama_model_n_devices(model);
     std::vector<llama_device_memory_data> ret(nd + 1);
 
@@ -257,6 +279,67 @@ common_device_memory_data_vec common_get_device_memory_data(
     return ret;
 }
 
+// drafts that borrow token_embd/output from their target (DFlash, DSpark, Eagle3) cannot create a context on their
+// own: while measuring such a draft, give it an allocation-free target context as ctx_other (its memory is counted with
+// the main model, not here) and the backend top-k samplers the speculative decoder installs on the draft
+struct common_fit_draft_target {
+    llama_context_params * cparams_dft;
+    llama_context_params   cparams_dft_orig;
+    llama_model   * model = nullptr;
+    llama_context * ctx   = nullptr;
+    std::vector<llama_sampler *>         chains;
+    std::vector<llama_sampler_seq_config> configs;
+
+    common_fit_draft_target(const char * path_model, const llama_model_params * mparams, const llama_context_params * cparams,
+            llama_context_params * cparams_dft, bool backend_sampling, ggml_log_level log_level)
+        : cparams_dft(cparams_dft), cparams_dft_orig(*cparams_dft) {
+        struct user_data_t {
+            ggml_log_callback callback;
+            void * user_data;
+            ggml_log_level min_level;
+        } ud;
+        llama_log_get(&ud.callback, &ud.user_data);
+        ud.min_level = log_level;
+        llama_log_set([](ggml_log_level level, const char * text, void * user_data) {
+            const user_data_t * ud = (const user_data_t *) user_data;
+            ud->callback(level >= ud->min_level ? level : GGML_LOG_LEVEL_DEBUG, text, ud->user_data);
+        }, &ud);
+
+        llama_model_params mparams_copy = *mparams;
+        mparams_copy.no_alloc  = true;
+        mparams_copy.load_mode = LLAMA_LOAD_MODE_NONE;
+        model = llama_model_load_from_file(path_model, mparams_copy);
+        if (model != nullptr) {
+            ctx = llama_init_from_model(model, *cparams);
+        }
+        llama_log_set(ud.callback, ud.user_data);
+        if (ctx == nullptr) {
+            return; // the draft is measured on its own, as before
+        }
+
+        cparams_dft->ctx_other = ctx;
+        if (backend_sampling && cparams_dft->n_samplers == 0) {
+            for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) std::max<uint32_t>(1, cparams_dft->n_seq_max); ++seq_id) {
+                llama_sampler * chain = llama_sampler_chain_init(llama_sampler_chain_default_params());
+                llama_sampler_chain_add(chain, llama_sampler_init_top_k(10));
+                chains.push_back(chain);
+                configs.push_back({ seq_id, chain });
+            }
+            cparams_dft->samplers   = configs.data();
+            cparams_dft->n_samplers = configs.size();
+        }
+    }
+
+    ~common_fit_draft_target() {
+        *cparams_dft = cparams_dft_orig;
+        for (llama_sampler * chain : chains) {
+            llama_sampler_free(chain);
+        }
+        llama_free(ctx);
+        llama_model_free(model);
+    }
+};
+
 static void common_params_fit_impl(
         const char * path_model, struct llama_model_params * mparams, struct llama_context_params * cparams,
         float * tensor_split, struct llama_model_tensor_buft_override * tensor_buft_overrides,
@@ -299,6 +382,12 @@ static void common_params_fit_impl(
             LOG_TRC("%s: getting device memory data for the extra model at a context size of %" PRIu32 ":\n",
                 __func__, cparams->n_ctx);
 
+            std::unique_ptr<common_fit_draft_target> draft_target;
+            if (!extra->shares_model) {
+                draft_target.reset(new common_fit_draft_target(
+                    path_model, mparams, cparams, extra->cparams, extra->backend_sampling, log_level));
+            }
+
             dmds_t measured;
             try {
                 measured = common_get_device_memory_data_impl(

common/fit.h

@@ -15,11 +15,13 @@ enum common_params_fit_status {
 // a second model that shares the devices of the main model, e.g. a draft model
 //   - its context follows the context of the main model, so its memory is measured again whenever that context changes
 //   - shares_model tells the fit that the weights are already counted in the main model, as for an MTP context
+//   - backend_sampling adds the backend sampler of the speculative decoder to the measured draft graph
 struct common_fit_extra_model {
     const char * path_model;
     llama_model_params * mparams;
     llama_context_params * cparams;
     bool shares_model;
+    bool backend_sampling = true; // the speculative decoder offloads draft sampling to the backend (its default)
 };
 
 // fits mparams and cparams to free device memory (assumes system memory is unlimited)

common/parsers/deepseek.cpp

@@ -195,6 +195,14 @@ common_chat_params common_chat_params_init_deepseek_v3_2(const common_chat_templ
             foreach_function(tools, [&](size_t tool_index, const json & tool) {
                 const auto & function = tool.at("function");
                 std::string  name     = function.at("name");
+                if (is_v41 && name.find("::") == std::string::npos) {
+                    // Match qualified_name() in the V4.1 template; rendering above validates conflicts.
+                    const auto ns = tool.contains("namespace") && !tool.at("namespace").is_null()
+                        ? tool.at("namespace") : function.value("namespace", json());
+                    if (!ns.is_null()) {
+                        name = (ns.is_object() ? ns.at("name") : ns).get<std::string>() + "::" + name;
+                    }
+                }
 
                 std::vector<common_peg_parser> required_parsers;
                 std::vector<common_peg_parser> optional_parsers;

conversion/deepseek41.py

@@ -2,6 +2,7 @@
 
 import json
 import math
+import os
 import re
 from pathlib import Path
 from typing import Any, Callable, Iterable, TYPE_CHECKING, cast
@@ -117,9 +118,10 @@ def _e8m0_check(scale_u8: np.ndarray, name: str) -> None:
 @ModelBase.register("DeepseekV41ForCausalLM")
 class DeepseekV41Model(DeepseekV4Model):
     """DeepSeek-V4.1 text model, lossless: FP8 (e4m3, 32x32 ue8m0 blocks) linears dequantized to BF16 (F32 with
-    --outtype f32), routed experts copied bit-exact as MXFP4, Engram tables (e4m3 rows, one e8m0 scale per 32) as the
-    exact BF16 lookup (F32 with --outtype f32), router and Engram gate F32. DSpark (mtp.*) is not converted here
-    (--dspark, DeepseekV41DSparkModel); the vision tower goes to an mmproj (--mmproj, DeepseekV41VisionModel).
+    --outtype f32), routed experts copied bit-exact as MXFP4, Engram tables (e4m3 rows, one e8m0 scale per 32) copied
+    bit-exact as MXFP8 (DSV41_ENGRAM_TYPE=bf16: the exact BF16 lookup, F32 with --outtype f32), router and Engram gate
+    F32. DSpark (mtp.*) is not converted here (--dspark, DeepseekV41DSparkModel); the vision tower goes to an mmproj
+    (--mmproj, DeepseekV41VisionModel).
     The Engram tables of all layers are written as one per_layer_token_embd (rows concatenated in layer order) with
     the PLE keys of qwen4exp: absolute head offsets and bucket counts, so the runtime can read it lazily."""
     model_arch = gguf.MODEL_ARCH.DEEPSEEK41
@@ -131,15 +133,18 @@ def __init__(self, *args, **kwargs):
         super().__init__(*args, **kwargs)
         self.hparams.setdefault("num_hash_layers", 0)  # no hash-routed layers in V4.1
         self._f32 = self.ftype == gguf.LlamaFileType.ALL_F32
+        self._engram_mxfp8 = os.environ.get("DSV41_ENGRAM_TYPE", "mxfp8").lower() != "bf16"
         qc = self.hparams.get("quantization_config") or {}
         self._fp8_block = tuple(qc.get("weight_block_size") or (32, 32))
         if qc.get("scale_fmt", "ue8m0") != "ue8m0":
             raise ValueError(f"unsupported FP8 scale format {qc.get('scale_fmt')!r}")
         if self.hparams.get("gate_temp", 1.0) != 1.0:
             raise ValueError("gate_temp != 1 is not supported (the router logits would need dividing by it)")
 
-        # V4.1 does not use the DeepSeek-V4 template; a template shipped with the model is still added by set_vocab
-        self.gguf_writer.remove_key(gguf.Keys.Tokenizer.CHAT_TEMPLATE)
+        template_path = Path(__file__).parent.parent / "models" / "templates" / "deepseek-ai-DeepSeek-V4.1-Flash-1017.jinja"
+        if template_path.is_file():
+            self.gguf_writer.remove_key(gguf.Keys.Tokenizer.CHAT_TEMPLATE)
+            self.gguf_writer.add_chat_template(template_path.read_text(encoding="utf-8"))
 
     def index_tensors(self, remote_hf_model_id: str | None = None) -> dict[str, Callable[[], Tensor]]:
         # num_hidden_layers lives in text_config, which is merged only after indexing; MTP stages are skipped anyway
@@ -221,7 +226,7 @@ def chunks() -> Iterable[np.ndarray]:
         self.gguf_writer.add_tensor(self.format_tensor_name(key, bid), data, raw_dtype=gguf.GGMLQuantizationType.MXFP4)
 
     def _write_engram_tables(self, layer_ids: list[int]) -> None:
-        # official lookup e4m3 * 2^(e8m0 - 127), exact in BF16
+        # official lookup e4m3 * 2^(e8m0 - 127): stored as the source MXFP8 blocks (bit-exact) or as exact BF16
         parts, dim = [], None
         for bid in layer_ids:
             wn, sn = f"layers.{bid}.engram.embed.weight", f"layers.{bid}.engram.embed.scale"
@@ -236,6 +241,22 @@ def _write_engram_tables(self, layer_ids: list[int]) -> None:
         name = self.format_tensor_name(gguf.MODEL_TENSOR.PER_LAYER_TOKEN_EMBD)
         n_rows = sum(p[2] for p in parts)
 
+        if self._engram_mxfp8:
+            def blocks() -> Iterable[np.ndarray]:
+                from safetensors import safe_open
+                for wn, sn, rows, wpart, spart in parts:
+                    with safe_open(wpart, framework="pt") as fw, safe_open(spart, framework="pt") as fs:
+                        ws, ss = fw.get_slice(wn), fs.get_slice(sn)
+                        for r0 in range(0, rows, 1 << 18):
+                            q = ws[r0:r0 + (1 << 18)].view(torch.uint8).reshape(-1, 32).numpy()
+                            e = ss[r0:r0 + (1 << 18)].view(torch.uint8).reshape(-1, 1).numpy()
+                            if (e == 0xFF).any() or ((q & 0x7F) == 0x7F).any():
+                                raise ValueError(f"{wn}: NaN E8M0 scale or E4M3 code")
+                            yield np.concatenate([e, q], axis=1)
+            data = cast(np.ndarray, _StreamedBytes((n_rows, dim // 32 * 33), blocks))
+            self.gguf_writer.add_tensor(name, data, raw_dtype=gguf.GGMLQuantizationType.MXFP8)
+            return
+
         def chunks() -> Iterable[np.ndarray]:
             from safetensors import safe_open
             for wn, sn, rows, wpart, spart in parts:

ggml/include/ggml.h

@@ -432,6 +432,10 @@ extern "C" {
         GGML_TYPE_NVFP4   = 40, // NVFP4 (4 blocks, E4M3 scale)
         GGML_TYPE_Q1_0    = 41,
         GGML_TYPE_Q2_0    = 42,
+        // MXFP8 (E4M3, E8M0 scale), gather-only. Fork-local, top of the reserved range so the
+        // count above the IQ1 types stays put; files written while it was 43 (before upstream
+        // took 43 for F8_E4M3) are told apart by their byte size when read (gguf.cpp).
+        GGML_TYPE_MXFP8   = 63,
         // IQ1_S's 8-element ternary codebook with a narrower index; see ggml-common.h.
         // 42..63 stay reserved for upstream (42 is already Q2_0), so fork-local types start
         // at 64 and a stock build rejects these files instead of misreading them.

ggml/src/ggml-common.h

@@ -225,6 +225,14 @@ typedef struct {
 } block_mxfp4;
 static_assert(sizeof(block_mxfp4) == sizeof(uint8_t) + QK_MXFP4/2, "wrong mxfp4 block size/padding");
 
+// MXFP8 (OCP MX): 32 E4M3 (fn) values sharing one E8M0 scale, value = e4m3(q[j]) * 2^(e - 127)
+#define QK_MXFP8 32
+typedef struct {
+    uint8_t e;          // E8M0
+    uint8_t qs[QK_MXFP8]; // E4M3 (fn) codes
+} block_mxfp8;
+static_assert(sizeof(block_mxfp8) == sizeof(uint8_t) + QK_MXFP8, "wrong mxfp8 block size/padding");
+
 #define QK_NVFP4 64
 #define QK_NVFP4_SUB 16  // sub-block size for per-group scales
 typedef struct {

ggml/src/ggml-cpu/amx/amx.cpp

@@ -155,6 +155,10 @@ class extra_buffer_type : ggml::cpu::extra_buffer_type {
         if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) {
             return false;
         }
+        // 2D weights only: the packed buffer holds a single slice; a batched src1 is folded into M
+        if (ggml_nrows(src0) != src0->ne[1]) {
+            return false;
+        }
         if (!src0->buffer || src0->buffer->buft != ggml_backend_amx_buffer_type()) {
             return false;
         }

ggml/src/ggml-cpu/amx/mmq.cpp

@@ -2349,13 +2349,15 @@ void ggml_backend_amx_mul_mat(const ggml_compute_params * params, struct ggml_te
     // amx kernels will be added once 6th gen xeon is released.
     const bool is_floating_type = TYPE == GGML_TYPE_F16;
 
-    const int M = dst->ne[1];
+    // single 2D src0 slice (supports_op) and contiguous src1/dst: all src1 batches are one gemm
+    GGML_ASSERT(ggml_nrows(src0) == src0->ne[1]);
+    const int M = ggml_nrows(dst);
     const int N = dst->ne[0];
     const int K = src0->ne[0];
     const int ldc = dst->nb[1] / dst->nb[0];
 
-    const int64_t ne2 = dst->ne[2];
-    const int64_t n_batch = ne2 * dst->ne[3];
+    const int64_t ne2 = 1;
+    const int64_t n_batch = 1;
 
     if (is_floating_type) {
         constexpr int BLOCK_M = 4;

ggml/src/ggml-cpu/ggml-cpu.cpp

@@ -429,6 +429,16 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st
         return true;
     }
 
+    // MXFP8 is a gather-only type: GET_ROWS reads it, nothing else does (and nothing writes it)
+    if (op->type == GGML_TYPE_MXFP8) {
+        return false;
+    }
+    for (int i = 0; i < GGML_MAX_SRC; i++) {
+        if (op->src[i] && op->src[i]->type == GGML_TYPE_MXFP8 && !(op->op == GGML_OP_GET_ROWS && i == 0)) {
+            return false;
+        }
+    }
+
     // check extra buffer types
     // note: only the first sources are checked for extra buffer types to reduce overhead, increase if necessary
     for (int i = 0; i < 4; i++) {

ggml/src/ggml-cpu/ops.cpp

@@ -5383,6 +5383,7 @@ void ggml_compute_forward_get_rows(
         case GGML_TYPE_Q8_1:
         case GGML_TYPE_MXFP4:
         case GGML_TYPE_NVFP4:
+        case GGML_TYPE_MXFP8:
         case GGML_TYPE_F8_E4M3:
         case GGML_TYPE_Q2_K:
         case GGML_TYPE_Q3_K:
@@ -6147,6 +6148,7 @@ void ggml_compute_forward_clamp(
         case GGML_TYPE_Q8_1:
         case GGML_TYPE_MXFP4:
         case GGML_TYPE_NVFP4:
+        case GGML_TYPE_MXFP8:
         case GGML_TYPE_F8_E4M3:
         case GGML_TYPE_Q2_K:
         case GGML_TYPE_Q3_K:

ggml/src/ggml-cuda/dequantize.cuh

@@ -119,6 +119,24 @@ static __device__ __forceinline__ void dequantize_q8_0(const void * vx, const in
     v.y *= d;
 }
 
+// E4M3 (fn) decode, exact in fp32; 0x7F/0xFF are NaN
+static __device__ __forceinline__ float ggml_cuda_e4m3_to_fp32(const uint8_t x) {
+    const int exp = (x >> 3) & 0xF;
+    const int man = x & 0x7;
+    float v = exp == 0 ? ldexpf((float) man, -9) : ldexpf((float) (8 + man), exp - 10);
+    v = (exp == 0xF && man == 0x7) ? __int_as_float(0x7FC00000) : v;
+    return (x & 0x80) ? -v : v;
+}
+
+static __device__ __forceinline__ void dequantize_mxfp8(const void * vx, const int64_t ib, const int iqs, float2 & v){
+    const block_mxfp8 * x = (const block_mxfp8 *) vx;
+
+    const float d = x[ib].e == 0xFF ? __int_as_float(0x7FC00000) : ggml_cuda_e8m0_to_fp32(x[ib].e);
+
+    v.x = ggml_cuda_e4m3_to_fp32(x[ib].qs[iqs + 0]) * d;
+    v.y = ggml_cuda_e4m3_to_fp32(x[ib].qs[iqs + 1]) * d;
+}
+
 //================================== k-quants
 
 // Each call dequantizes one super-block of QK_K values into y using the

ggml/src/ggml-cuda/getrows.cu

@@ -344,6 +344,10 @@ static void ggml_cuda_get_rows_switch_src0_type(
             get_rows_cuda_q<QK8_0, QR8_0, dequantize_q8_0>(src0_d, src1_d, dst_d,
                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
             break;
+        case GGML_TYPE_MXFP8:
+            get_rows_cuda_q<QK_MXFP8, 1, dequantize_mxfp8>(src0_d, src1_d, dst_d,
+                ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
+            break;
         case GGML_TYPE_Q2_K:
             get_rows_cuda_kq<64, dst_t, dequantize_q2_K<dst_t>>(src0_d, src1_d, dst_d,
                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -5901,6 +5901,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
                     case GGML_TYPE_IQ1_S:
                     case GGML_TYPE_IQ1_M:
                     case GGML_TYPE_IQ4_XS:
+                    case GGML_TYPE_MXFP8:
                         return true;
                     case GGML_TYPE_IQ4_NL:
                     case GGML_TYPE_MXFP4: