From 142cea8bde2b3de0b0f2c4008dec81037db1e3fd Mon Sep 17 00:00:00 2001 From: Cursor Agent Date: Tue, 15 Sep 2026 12:09:02 +0000 Subject: [PATCH] spec : add DFlash2 support and IQ1 pairing Assisted-by: Cursor Grok 4.6 Co-authored-by: Jackson --- common/speculative.cpp | 101 ++++++---- conversion/__init__.py | 1 + conversion/muse_glimmer.py | 17 +- conversion/qwen.py | 67 ++++++- docs/speculative.md | 20 ++ ggml/src/ggml-cpu/arch-fallback.h | 2 - ggml/src/ggml-cpu/arch/x86/quants.c | 97 +++++++++ gguf-py/gguf/constants.py | 25 +++ gguf-py/gguf/gguf_writer.py | 12 ++ gguf-py/gguf/tensor_mapping.py | 28 +++ src/llama-arch.cpp | 19 ++ src/llama-arch.h | 12 ++ src/llama-context.cpp | 7 +- src/llama-ext.h | 2 + src/llama-hparams.h | 6 + src/llama-model.cpp | 8 + src/llama-model.h | 9 + src/llama-quant.cpp | 12 ++ src/models/dflash.cpp | 301 +++++++++++++++++++++++++--- tests/test-quantize-fns.cpp | 39 ++-- 20 files changed, 700 insertions(+), 85 deletions(-) diff --git a/common/speculative.cpp b/common/speculative.cpp index 2ee1e6b84818..816e28d2a824 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -13,8 +13,10 @@ #include #include +#include #include #include +#include #include #include @@ -929,15 +931,16 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { int32_t block_size = 0; llama_token mask_token_id = 0; + bool is_dflash2 = false; + bool is_mrope = false; + int32_t selector_top_k = 0; + // draft-dspark: the draft carries a Markov head and uses an anchor-first block layout const bool is_dspark; const int32_t * target_layer_ids = nullptr; // model_dft's extract layer indices uint32_t target_layer_ids_n = 0; - // scratch buffer for concatenated target features [n_tokens, n_embd_enc] - std::vector features_buf; - common_speculative_impl_draft_dflash(const common_params_speculative & params, uint32_t n_seq, common_speculative_type type = COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH) : common_speculative_impl(type, n_seq) @@ -967,11 +970,15 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { block_size = std::atoi(buf); } } + selector_top_k = llama_model_dflash_selector_top_k(model_dft); + is_dflash2 = selector_top_k > 0; mask_token_id = llama_vocab_mask(llama_model_get_vocab(model_dft)); LOG_INF("%s: adding speculative implementation '%s'\n", __func__, common_speculative_type_to_str(type).c_str()); LOG_INF("%s: - n_max=%d, n_min=%d, p_min=%.2f\n", __func__, this->params.n_max, this->params.n_min, this->params.p_min); - LOG_INF("%s: - block_size=%d, mask_token_id=%d, n_extract=%u\n", __func__, block_size, mask_token_id, target_layer_ids_n); + LOG_INF("%s: - block_size=%d, mask_token_id=%d, n_extract=%u%s\n", __func__, + block_size, mask_token_id, target_layer_ids_n, + is_dflash2 ? ", dflash2=true" : ""); // DFlash input is [id_last, * (block_size-1)]: in-place denoising yields at most // block_size-1 draft tokens, DSpark yield a full block_size draft tokens @@ -984,7 +991,13 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { } batch = llama_batch_init(llama_n_batch(ctx_dft), 0, n_seq); - batch_inject = llama_batch_init(llama_n_batch(ctx_dft), n_embd_dec, n_seq); + batch_inject = llama_batch_init(llama_n_ubatch(ctx_dft), n_embd_enc, n_seq); + + is_mrope = llama_model_rope_type(model_dft) == LLAMA_ROPE_TYPE_MROPE; + if (is_mrope) { + free(batch_inject.pos); + batch_inject.pos = (llama_pos *) malloc(sizeof(llama_pos) * 4 * llama_n_batch(ctx_dft)); + } smpls.resize(n_seq); for (auto & s : smpls) { @@ -1000,7 +1013,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { llama_set_embeddings_layer_inp(ctx_tgt, (uint32_t) target_layer_ids[k], true); } - llama_set_embeddings_nextn(ctx_dft, true, /*masked*/ true); + llama_set_embeddings_nextn(ctx_dft, true, /*masked*/ !is_dflash2); llama_set_causal_attn(ctx_dft, false); // DFlash needs non-causal attention } @@ -1071,55 +1084,40 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { } const int32_t n_rows = i_batch_end[seq_id] - i_batch_beg[seq_id] + 1; + const bool pos_pinned = batch_in.pos[i_batch_beg[seq_id]] == batch_in.pos[i_batch_end[seq_id]]; + if (has_embeddings && n_rows > 1 && pos_pinned) { + continue; + } + for (int32_t offset = 0; offset < n_rows; offset += n_ubatch) { const int32_t n_chunk = std::min(n_ubatch, n_rows - offset); - // gather this chunk's target features, interleaved by extract layer - features_buf.resize((size_t) n_chunk * n_embd_enc); + batch_inject.n_tokens = n_chunk; for (uint32_t k = 0; k < target_layer_ids_n; ++k) { const float * layer = llama_get_embeddings_layer_inp(ctx_tgt, (uint32_t) target_layer_ids[k]); if (!layer) { GGML_ABORT("DFlash: target layer %d input not extracted.", target_layer_ids[k]); } for (int32_t i = 0; i < n_chunk; ++i) { - float * dst = features_buf.data() + (size_t) i * n_embd_enc + k * (size_t) n_embd_tgt; + float * dst = batch_inject.embd + (size_t) i * n_embd_enc + k * (size_t) n_embd_tgt; const float * src = layer + (size_t) (i_batch_beg[seq_id] + offset + i) * n_embd_tgt; std::memcpy(dst, src, (size_t) n_embd_tgt * sizeof(float)); } } - // fuse extracted features through DFlash encoder - llama_batch enc_batch = { - /*.n_tokens =*/ n_chunk, - /*.token =*/ nullptr, - /*.embd =*/ features_buf.data(), - /*.pos =*/ nullptr, - /*.n_seq_id =*/ nullptr, - /*.seq_id =*/ nullptr, - /*.logits =*/ nullptr, - }; - - int32_t rc = llama_encode(ctx_dft, enc_batch); - if (rc != 0) { - LOG_ERR("%s: llama_encode(ctx_dft) failed rc=%d (n_tokens=%d, offset=%d)\n", - __func__, rc, (int) n_chunk, (int) offset); - return false; - } - - const float * inp_g = llama_get_embeddings_nextn(ctx_dft); - GGML_ASSERT(inp_g && "DFlash encoder produced no output."); - - // inject the DFlash decoder K/V cache at the tokens' target positions - batch_inject.n_tokens = n_chunk; - std::memcpy(batch_inject.embd, inp_g, (size_t) n_chunk * n_embd_dec * sizeof(float)); - for (int32_t i = 0; i < n_chunk; ++i) { - batch_inject.pos[i] = batch_in.pos[i_batch_beg[seq_id] + offset + i]; + const llama_pos p = batch_in.pos[i_batch_beg[seq_id] + offset + i]; + batch_inject.pos[i] = p; + if (is_mrope) { + batch_inject.pos[1 * n_chunk + i] = p; + batch_inject.pos[2 * n_chunk + i] = p; + batch_inject.pos[3 * n_chunk + i] = 0; + } batch_inject.n_seq_id[i] = 1; batch_inject.seq_id[i][0] = seq_id; batch_inject.logits[i] = false; } - rc = llama_decode(ctx_dft, batch_inject); + const int32_t rc = llama_decode(ctx_dft, batch_inject); if (rc != 0) { LOG_ERR("%s: llama_decode(ctx_dft) failed rc=%d (n_tokens=%d, offset=%d)\n", __func__, rc, (int) n_chunk, (int) offset); @@ -1157,7 +1155,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { i_block_beg[seq_id] = batch.n_tokens; n_block [seq_id] = n_block_tokens; for (int32_t i = 0; i < n_block_tokens; ++i) { - common_batch_add(batch, i == 0 ? dp.id_last : mask_token_id, n + i, { seq_id }, true); + common_batch_add(batch, i == 0 ? dp.id_last : mask_token_id, n + i, { seq_id }, !is_dflash2); } } @@ -1185,6 +1183,35 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { auto & result = *dp.result; + if (is_dflash2) { + const float * lattice = llama_get_embeddings_nextn(ctx_dft); + GGML_ASSERT(lattice && "DFlash2 selector produced no lattice"); + + int32_t predecessor = 0; + for (int32_t i = 1; i < n_block_tokens; ++i) { + const float * row = lattice + (size_t) (beg + i) * n_embd_dec; + const float * scores = row + selector_top_k + (size_t) predecessor * selector_top_k; + + predecessor = (int32_t) std::distance(scores, + std::max_element(scores, scores + selector_top_k)); + if (params.p_min > 0.0f) { + float sum = 0.0f; + for (int32_t k = 0; k < selector_top_k; ++k) { + sum += std::exp(scores[k] - scores[predecessor]); + } + if (1.0f / sum < params.p_min) { + break; + } + } + result.push_back((llama_token) row[predecessor]); + } + + if (result.size() < (size_t) params.n_min) { + result.clear(); + } + continue; + } + if (is_dspark) { // DSpark predicts the next token from position 0 and optionally truncates // at the first position below the confidence threshold. diff --git a/conversion/__init__.py b/conversion/__init__.py index c7d8046c495f..8900dc80b76a 100644 --- a/conversion/__init__.py +++ b/conversion/__init__.py @@ -53,6 +53,7 @@ "DeepseekV3ForCausalLM": "deepseek", "DeepseekV32ForCausalLM": "deepseek", "DFlashDraftModel": "qwen", + "DFlash2DraftModel": "qwen", "Qwen3DSparkModel": "qwen", "DeepseekV4ForCausalLM": "deepseek", "DeepseekV4DSparkModel": "deepseek", diff --git a/conversion/muse_glimmer.py b/conversion/muse_glimmer.py index cc588e8321bd..1a45455533b3 100644 --- a/conversion/muse_glimmer.py +++ b/conversion/muse_glimmer.py @@ -163,11 +163,19 @@ def set_gguf_parameters(self): super().set_gguf_parameters() h = self.hparams - self.gguf_writer.add_block_size(int(h["block_size"])) + self.gguf_writer.add_block_size(int(h.get("block_size", h.get("dflash_config", {}).get("block_size", 16)))) # dflash.target_layers[k] refers to the inputs going into the ith layer, which come from the (i-1)th layer's output. # The transformers configuration refers to the outputs being recorded. - self.gguf_writer.add_target_layers([int(x) + 1 for x in h["target_layer_ids"]]) + target_layer_ids = h.get("target_layer_ids") or h.get("dflash_config", {}).get("target_layer_ids", []) + self.gguf_writer.add_target_layers([int(x) + 1 for x in target_layer_ids]) + + dflash_config = h.get("dflash_config", {}) + if "conv_kernel_size" in dflash_config or "conv_kernel_size" in h: + self.gguf_writer.add_conv_kernel_size(int(dflash_config.get("conv_kernel_size", h["conv_kernel_size"]))) + self.gguf_writer.add_conv_group_size(int(dflash_config.get("conv_group_size", h["conv_group_size"]))) + self.gguf_writer.add_selector_rank(int(dflash_config.get("selector_rank", h["selector_rank"]))) + self.gguf_writer.add_selector_top_k(int(dflash_config.get("selector_top_k", h["selector_top_k"]))) if h.get("sliding_window") and h.get("layer_types"): self.gguf_writer.add_sliding_window(int(h["sliding_window"])) @@ -176,4 +184,9 @@ def set_gguf_parameters(self): def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: # DFlash defaults to NEOX (rotate_half) rope, matching transformers HF layout for Q/K, QK-norms # no permutation needed. + if name in ( + "model.candidate_selector.predecessor_codebook", + "model.candidate_selector.successor_codebook", + ): + name += ".weight" yield (self.map_tensor_name(name), data_torch) diff --git a/conversion/qwen.py b/conversion/qwen.py index b4ae528bf2d4..ca233538bbec 100644 --- a/conversion/qwen.py +++ b/conversion/qwen.py @@ -629,7 +629,7 @@ class Qwen3_5MoeTextModel(_Qwen35MRopeMixin, _LinearAttentionVReorderBase): model_arch = gguf.MODEL_ARCH.QWEN35MOE -@ModelBase.register("DFlashDraftModel") +@ModelBase.register("DFlashDraftModel", "DFlash2DraftModel") class DFlashModel(Qwen3Model): model_arch = gguf.MODEL_ARCH.DFLASH @@ -664,9 +664,31 @@ def set_vocab(self): def set_gguf_parameters(self): super().set_gguf_parameters() - block_size = self.hparams.get("block_size", 16) - self.gguf_writer.add_block_size(block_size) dflash_config = self.hparams.get("dflash_config", {}) + block_size = dflash_config.get("block_size", self.hparams.get("block_size", 16)) + self.gguf_writer.add_block_size(block_size) + + if "conv_kernel_size" in dflash_config: + self.gguf_writer.add_conv_kernel_size(int(dflash_config["conv_kernel_size"])) + self.gguf_writer.add_conv_group_size(int(dflash_config["conv_group_size"])) + self.gguf_writer.add_selector_rank(int(dflash_config["selector_rank"])) + self.gguf_writer.add_selector_top_k(int(dflash_config["selector_top_k"])) + + output_multiplier = dflash_config.get( + "output_multiplier", self.hparams.get("output_multiplier") + ) + if output_multiplier is not None: + self.gguf_writer.add_logit_scale(float(output_multiplier)) + softcap = dflash_config.get( + "final_logit_softcapping", self.hparams.get("final_logit_softcapping") + ) + if softcap is not None and float(softcap) > 0: + self.gguf_writer.add_final_logit_softcapping(float(softcap)) + embedding_scale = dflash_config.get( + "input_embedding_scale", self.hparams.get("input_embedding_scale") + ) + if embedding_scale is not None: + self.gguf_writer.add_embedding_scale(float(embedding_scale)) target_layer_ids = dflash_config.get("target_layer_ids", []) if target_layer_ids: @@ -681,6 +703,20 @@ def set_gguf_parameters(self): self.gguf_writer.add_sliding_window(sliding_window) self.gguf_writer.add_sliding_window_pattern(is_swa) + # M-RoPE target: the draft ropes on the temporal dim only + if self._target_uses_mrope(): + head_dim = self.hparams.get("head_dim") or self.hparams["hidden_size"] // self.hparams["num_attention_heads"] + self.gguf_writer.add_rope_dimension_sections([head_dim // 2, 0, 0, 0]) + + def _target_uses_mrope(self) -> bool: + if self.target_model_dir is None: + return False + with open(self.target_model_dir / "config.json", "r", encoding="utf-8") as f: + cfg = json.load(f) + cfg = cfg.get("text_config", cfg) + rope = cfg.get("rope_parameters") or cfg.get("rope_scaling") or {} + return "mrope_section" in rope + @classmethod def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None: name, gen = item @@ -688,6 +724,31 @@ def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Ca name = "model." + name return super().filter_tensors((name, gen)) + _ROPE_PERMUTE_SUFFIXES = ( + "self_attn.q_proj.weight", + "self_attn.k_proj.weight", + "self_attn.q_norm.weight", + "self_attn.k_norm.weight", + ) + + def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: + if name == "model.embed_tokens.weight" and not self.hparams.get("has_embed_tokens", True): + return + + # interleaved-rope checkpoints (rope_is_neox_style = false) -> NeoX layout + if not self.hparams.get("rope_is_neox_style", True) and name.endswith(self._ROPE_PERMUTE_SUFFIXES): + head_dim = self.hparams["head_dim"] + shape = data_torch.shape + data_torch = data_torch.reshape(-1, head_dim // 2, 2, *shape[1:]).transpose(1, 2).reshape(shape) + + if name in ( + "model.candidate_selector.predecessor_codebook", + "model.candidate_selector.successor_codebook", + ): + name += ".weight" + + yield from super().modify_tensors(data_torch, name, bid) + @ModelBase.register("Qwen3DSparkModel") class DSparkModel(DFlashModel): diff --git a/docs/speculative.md b/docs/speculative.md index 25abef1b602b..ff7495ddb3c4 100644 --- a/docs/speculative.md +++ b/docs/speculative.md @@ -74,9 +74,29 @@ llama-server -m Qwen3-4B.gguf -md Qwen3-4B-DFlash.gguf \ `--spec-draft-n-max` is clamped to the draft model's trained block size. +DFlash 2 drafts (`DFlash2DraftModel`, for example `z-lab/Qwen3.8-27B-DFlash2`) use the same +`--spec-type draft-dflash` flag. The runtime detects them from `selector_top_k` metadata and +runs the in-graph candidate selector plus local convolutions instead of independent per-position +argmax. Convert them the same way: + +```bash +python convert_hf_to_gguf.py z-lab/Qwen3.8-27B-DFlash2 \ + --target-model-dir Qwen/Qwen3.8-27B --outtype bf16 --outfile Qwen3.8-27B-DFlash2.gguf + +llama-server -m Qwen3.8-27B.gguf -md Qwen3.8-27B-DFlash2.gguf \ + --spec-type draft-dflash --spec-draft-n-max 7 -fa on --jinja +``` + +IQ1 target + higher-precision draft is the intended pairing: keep the target at IQ1_XS / IQ1_XXS / +IQ1_XXXS and quantize the DFlash draft to Q8_0 or Q4_K. If you IQ1-quantize a DFlash 2 draft, +selector, conv, and `fc.weight` stay Q8_0 so lattice scores stay usable. Do not IQ1 the draft +backbone; DFlash drafts a full block per step and IQ1 MMVQ is capped at batch 8. + See: - #22105 +- #27342 +- #27310 ### DSpark (`draft-dspark`) diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h index 8658562251cb..994ebd403de4 100644 --- a/ggml/src/ggml-cpu/arch-fallback.h +++ b/ggml/src/ggml-cpu/arch-fallback.h @@ -117,8 +117,6 @@ #define ggml_gemm_mxfp4_4x4_q8_0_generic ggml_gemm_mxfp4_4x4_q8_0 #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 -#define ggml_vec_dot_iq1_xs_q8_K_generic ggml_vec_dot_iq1_xs_q8_K -#define ggml_vec_dot_iq1_xxs_q8_K_generic ggml_vec_dot_iq1_xxs_q8_K #elif defined(__POWERPC__) || defined(__powerpc__) // ref: https://github.com/ggml-org/llama.cpp/pull/14146#issuecomment-2972561679 diff --git a/ggml/src/ggml-cpu/arch/x86/quants.c b/ggml/src/ggml-cpu/arch/x86/quants.c index a69391672cea..ca9c38ae8b9a 100644 --- a/ggml/src/ggml-cpu/arch/x86/quants.c +++ b/ggml/src/ggml-cpu/arch/x86/quants.c @@ -3710,6 +3710,103 @@ void ggml_vec_dot_iq1_s_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo #endif } +void ggml_vec_dot_iq1_xs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + assert(n % QK_K == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + +#if defined __AVX2__ + const block_iq1_xs * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + const int nb = n / QK_K; + float sumf = 0; + + for (int i = 0; i < nb; ++i) { + __m256i sumi = _mm256_setzero_si256(); + int correction = 0; + int sumi1 = 0; + for (int ib = 0; ib < QK_K/32; ++ib) { + const uint8_t * qs = x[i].qs + 4*ib; + const int qh = x[i].qh[ib]; + const int nib = (x[i].sc[ib/2] >> (4*(ib & 1))) & 0xf; + const int ls = 2*(nib & 7) + 1; + const __m256i q1 = _mm256_set_epi64x( + iq1_xs_grid[qs[3] | (((qh >> 6) & 3) << 8)], + iq1_xs_grid[qs[2] | (((qh >> 4) & 3) << 8)], + iq1_xs_grid[qs[1] | (((qh >> 2) & 3) << 8)], + iq1_xs_grid[qs[0] | ((qh & 3) << 8)]); + const __m256i q8 = _mm256_loadu_si256((const __m256i *)(y[i].qs + 32*ib)); +#if defined(__AVX512VNNI__) && defined(__AVX512VL__) + const __m256i q1u = _mm256_add_epi8(q1, _mm256_set1_epi8(1)); + const __m256i dot = _mm256_dpbusd_epi32(_mm256_setzero_si256(), q1u, q8); + sumi = _mm256_add_epi32(sumi, _mm256_mullo_epi32(dot, _mm256_set1_epi32(ls))); + correction += ls * (y[i].bsums[2*ib] + y[i].bsums[2*ib + 1]); +#else + const __m256i dot = mul_add_epi8(q1, q8); + sumi = _mm256_add_epi32(sumi, _mm256_madd_epi16(dot, _mm256_set1_epi16(ls))); +#endif + sumi1 += ls * (nib & 8 ? -1 : 1) * (y[i].bsums[2*ib] + y[i].bsums[2*ib + 1]); + } + const float d = GGML_CPU_FP16_TO_FP32(x[i].d) * y[i].d; + sumf += d * (hsum_i32_8(sumi) - correction + IQ1S_DELTA * sumi1); + } + *s = sumf; +#else + ggml_vec_dot_iq1_xs_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + +void ggml_vec_dot_iq1_xxs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + assert(n % QK_K == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + +#if defined __AVX2__ + const block_iq1_xxs * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + const int nb = n / QK_K; + float sumf = 0; + + for (int i = 0; i < nb; ++i) { + __m256i sumi = _mm256_setzero_si256(); + int correction = 0; + int sumi1 = 0; + for (int ib = 0; ib < QK_K/32; ++ib) { + const uint8_t * qs = x[i].qs + 4*ib; + const int qh = x[i].qh[ib]; + const int ls = 2*((qh >> 4) & 7) + 1; + const __m256i q1 = _mm256_set_epi64x( + iq1_xxs_grid[qs[3] | (((qh >> 3) & 1) << 8)], + iq1_xxs_grid[qs[2] | (((qh >> 2) & 1) << 8)], + iq1_xxs_grid[qs[1] | (((qh >> 1) & 1) << 8)], + iq1_xxs_grid[qs[0] | ((qh & 1) << 8)]); + const __m256i q8 = _mm256_loadu_si256((const __m256i *)(y[i].qs + 32*ib)); +#if defined(__AVX512VNNI__) && defined(__AVX512VL__) + const __m256i q1u = _mm256_add_epi8(q1, _mm256_set1_epi8(1)); + const __m256i dot = _mm256_dpbusd_epi32(_mm256_setzero_si256(), q1u, q8); + sumi = _mm256_add_epi32(sumi, _mm256_mullo_epi32(dot, _mm256_set1_epi32(ls))); + correction += ls * (y[i].bsums[2*ib] + y[i].bsums[2*ib + 1]); +#else + const __m256i dot = mul_add_epi8(q1, q8); + sumi = _mm256_add_epi32(sumi, _mm256_madd_epi16(dot, _mm256_set1_epi16(ls))); +#endif + sumi1 += ls * (qh & 0x80 ? -1 : 1) * (y[i].bsums[2*ib] + y[i].bsums[2*ib + 1]); + } + const float d = GGML_CPU_FP16_TO_FP32(x[i].d) * y[i].d; + sumf += d * (hsum_i32_8(sumi) - correction + IQ1S_DELTA * sumi1); + } + *s = sumf; +#else + ggml_vec_dot_iq1_xxs_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + void ggml_vec_dot_iq1_xxxs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { assert(n % QK_K == 0); assert(nrc == 1); diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index af6ccf7762ff..ec49a6213051 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -161,6 +161,10 @@ class LLM: TARGET_LAYERS = "{arch}.target_layers" TARGET_HIDDEN_SIZE = "{arch}.target_hidden_size" BLOCK_SIZE = "{arch}.block_size" + CONV_KERNEL_SIZE = "{arch}.conv_kernel_size" + CONV_GROUP_SIZE = "{arch}.conv_group_size" + SELECTOR_RANK = "{arch}.selector_rank" + SELECTOR_TOP_K = "{arch}.selector_top_k" NORM_BEFORE_RESIDUAL = "{arch}.norm_before_residual" NORM_BEFORE_FC = "{arch}.norm_before_fc" @@ -1077,6 +1081,13 @@ class MODEL_TENSOR(IntEnum): DSPARK_MARKOV_W1 = auto() # markov head: prev-token embed DSPARK_MARKOV_W2 = auto() # markov head: bias projection DSPARK_CONF_PROJ = auto() # confidence head + DFLASH_ATTN_CONV_BASE = auto() + DFLASH_ATTN_CONV_PROJ = auto() + DFLASH_FFN_CONV_BASE = auto() + DFLASH_FFN_CONV_PROJ = auto() + DFLASH_SELECTOR_PREV = auto() + DFLASH_SELECTOR_NEXT = auto() + DFLASH_SELECTOR_HIDDEN = auto() # lfm2 audio A_ENC_NORM_CONV = auto() A_ENC_LINEAR_POS = auto() @@ -1777,6 +1788,13 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.DSPARK_MARKOV_W1: "markov_w1", MODEL_TENSOR.DSPARK_MARKOV_W2: "markov_w2", MODEL_TENSOR.DSPARK_CONF_PROJ: "conf_proj", + MODEL_TENSOR.DFLASH_ATTN_CONV_BASE: "blk.{bid}.attn_conv_base", + MODEL_TENSOR.DFLASH_ATTN_CONV_PROJ: "blk.{bid}.attn_conv_proj", + MODEL_TENSOR.DFLASH_FFN_CONV_BASE: "blk.{bid}.ffn_conv_base", + MODEL_TENSOR.DFLASH_FFN_CONV_PROJ: "blk.{bid}.ffn_conv_proj", + MODEL_TENSOR.DFLASH_SELECTOR_PREV: "selector_predecessor", + MODEL_TENSOR.DFLASH_SELECTOR_NEXT: "selector_successor", + MODEL_TENSOR.DFLASH_SELECTOR_HIDDEN: "selector_hidden", MODEL_TENSOR.D2T: "d2t", } @@ -4671,6 +4689,13 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.DSPARK_MARKOV_W1, MODEL_TENSOR.DSPARK_MARKOV_W2, MODEL_TENSOR.DSPARK_CONF_PROJ, + MODEL_TENSOR.DFLASH_ATTN_CONV_BASE, + MODEL_TENSOR.DFLASH_ATTN_CONV_PROJ, + MODEL_TENSOR.DFLASH_FFN_CONV_BASE, + MODEL_TENSOR.DFLASH_FFN_CONV_PROJ, + MODEL_TENSOR.DFLASH_SELECTOR_PREV, + MODEL_TENSOR.DFLASH_SELECTOR_NEXT, + MODEL_TENSOR.DFLASH_SELECTOR_HIDDEN, ], MODEL_ARCH.MISTRAL4: [ MODEL_TENSOR.TOKEN_EMBD, diff --git a/gguf-py/gguf/gguf_writer.py b/gguf-py/gguf/gguf_writer.py index 81ae07c11a60..a6476af62fa0 100644 --- a/gguf-py/gguf/gguf_writer.py +++ b/gguf-py/gguf/gguf_writer.py @@ -981,6 +981,18 @@ def add_sliding_window(self, value: int) -> None: def add_block_size(self, value: int) -> None: self.add_uint32(Keys.LLM.BLOCK_SIZE.format(arch=self.arch), value) + def add_conv_kernel_size(self, value: int) -> None: + self.add_uint32(Keys.LLM.CONV_KERNEL_SIZE.format(arch=self.arch), value) + + def add_conv_group_size(self, value: int) -> None: + self.add_uint32(Keys.LLM.CONV_GROUP_SIZE.format(arch=self.arch), value) + + def add_selector_rank(self, value: int) -> None: + self.add_uint32(Keys.LLM.SELECTOR_RANK.format(arch=self.arch), value) + + def add_selector_top_k(self, value: int) -> None: + self.add_uint32(Keys.LLM.SELECTOR_TOP_K.format(arch=self.arch), value) + def add_target_layers(self, value: Sequence[int]) -> None: self.add_array(Keys.LLM.TARGET_LAYERS.format(arch=self.arch), value) diff --git a/gguf-py/gguf/tensor_mapping.py b/gguf-py/gguf/tensor_mapping.py index 79d270ab8fe9..ec2d1272aba5 100644 --- a/gguf-py/gguf/tensor_mapping.py +++ b/gguf-py/gguf/tensor_mapping.py @@ -1318,6 +1318,34 @@ class TensorNameMap: "model.confidence_head.proj", # dspark ), + MODEL_TENSOR.DFLASH_ATTN_CONV_BASE: ( + "model.layers.{bid}.attention_conv.base_kernel", + ), + + MODEL_TENSOR.DFLASH_ATTN_CONV_PROJ: ( + "model.layers.{bid}.attention_conv.kernel_projection", + ), + + MODEL_TENSOR.DFLASH_FFN_CONV_BASE: ( + "model.layers.{bid}.mlp_conv.base_kernel", + ), + + MODEL_TENSOR.DFLASH_FFN_CONV_PROJ: ( + "model.layers.{bid}.mlp_conv.kernel_projection", + ), + + MODEL_TENSOR.DFLASH_SELECTOR_PREV: ( + "model.candidate_selector.predecessor_codebook", + ), + + MODEL_TENSOR.DFLASH_SELECTOR_NEXT: ( + "model.candidate_selector.successor_codebook", + ), + + MODEL_TENSOR.DFLASH_SELECTOR_HIDDEN: ( + "model.candidate_selector.hidden_projection", + ), + MODEL_TENSOR.CLS: ( "classifier", # jina "classifier.dense", # roberta diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index 73fb8b981382..9a70ed27b732 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -323,6 +323,11 @@ static const std::map LLM_KV_NAMES = { { LLM_KV_CLASSIFIER_OUTPUT_LABELS, "%s.classifier.output_labels" }, { LLM_KV_TARGET_LAYERS, "%s.target_layers" }, + { LLM_KV_DFLASH_BLOCK_SIZE, "%s.block_size" }, + { LLM_KV_DFLASH_CONV_KERNEL_SIZE, "%s.conv_kernel_size" }, + { LLM_KV_DFLASH_CONV_GROUP_SIZE, "%s.conv_group_size" }, + { LLM_KV_DFLASH_SELECTOR_RANK, "%s.selector_rank" }, + { LLM_KV_DFLASH_SELECTOR_TOP_K, "%s.selector_top_k" }, { LLM_KV_TARGET_HIDDEN_SIZE, "%s.target_hidden_size" }, { LLM_KV_NORM_BEFORE_RESIDUAL, "%s.norm_before_residual" }, { LLM_KV_NORM_BEFORE_FC, "%s.norm_before_fc" }, @@ -627,6 +632,13 @@ static const std::map LLM_TENSOR_NAMES = { { LLM_TENSOR_DSPARK_MARKOV_W1, "markov_w1" }, { LLM_TENSOR_DSPARK_MARKOV_W2, "markov_w2" }, { LLM_TENSOR_DSPARK_CONF_PROJ, "conf_proj" }, + { LLM_TENSOR_DFLASH_ATTN_CONV_BASE, "blk.%d.attn_conv_base" }, + { LLM_TENSOR_DFLASH_ATTN_CONV_PROJ, "blk.%d.attn_conv_proj" }, + { LLM_TENSOR_DFLASH_FFN_CONV_BASE, "blk.%d.ffn_conv_base" }, + { LLM_TENSOR_DFLASH_FFN_CONV_PROJ, "blk.%d.ffn_conv_proj" }, + { LLM_TENSOR_DFLASH_SELECTOR_PREV, "selector_predecessor" }, + { LLM_TENSOR_DFLASH_SELECTOR_NEXT, "selector_successor" }, + { LLM_TENSOR_DFLASH_SELECTOR_HIDDEN, "selector_hidden" }, }; // declare information about the model weight tensors: @@ -885,6 +897,13 @@ static const std::map LLM_TENSOR_INFOS = { {LLM_TENSOR_DSPARK_MARKOV_W1, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_GET_ROWS}}, {LLM_TENSOR_DSPARK_MARKOV_W2, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL_MAT}}, {LLM_TENSOR_DSPARK_CONF_PROJ, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL_MAT}}, + {LLM_TENSOR_DFLASH_ATTN_CONV_BASE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}}, + {LLM_TENSOR_DFLASH_ATTN_CONV_PROJ, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, + {LLM_TENSOR_DFLASH_FFN_CONV_BASE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}}, + {LLM_TENSOR_DFLASH_FFN_CONV_PROJ, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, + {LLM_TENSOR_DFLASH_SELECTOR_PREV, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_GET_ROWS}}, + {LLM_TENSOR_DFLASH_SELECTOR_NEXT, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_GET_ROWS}}, + {LLM_TENSOR_DFLASH_SELECTOR_HIDDEN, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL_MAT}}, }; LLM_KV::LLM_KV(llm_arch arch, const char * suffix) : arch(arch), suffix(suffix) {} diff --git a/src/llama-arch.h b/src/llama-arch.h index 51dfd288ced2..a2fb6c3ac6ab 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -370,6 +370,11 @@ enum llm_kv { LLM_KV_TARGET_LAYERS, LLM_KV_TARGET_HIDDEN_SIZE, + LLM_KV_DFLASH_BLOCK_SIZE, + LLM_KV_DFLASH_CONV_KERNEL_SIZE, + LLM_KV_DFLASH_CONV_GROUP_SIZE, + LLM_KV_DFLASH_SELECTOR_RANK, + LLM_KV_DFLASH_SELECTOR_TOP_K, LLM_KV_NORM_BEFORE_RESIDUAL, LLM_KV_NORM_BEFORE_FC, @@ -635,6 +640,13 @@ enum llm_tensor { LLM_TENSOR_DSPARK_MARKOV_W1, LLM_TENSOR_DSPARK_MARKOV_W2, LLM_TENSOR_DSPARK_CONF_PROJ, + LLM_TENSOR_DFLASH_ATTN_CONV_BASE, + LLM_TENSOR_DFLASH_ATTN_CONV_PROJ, + LLM_TENSOR_DFLASH_FFN_CONV_BASE, + LLM_TENSOR_DFLASH_FFN_CONV_PROJ, + LLM_TENSOR_DFLASH_SELECTOR_PREV, + LLM_TENSOR_DFLASH_SELECTOR_NEXT, + LLM_TENSOR_DFLASH_SELECTOR_HIDDEN, }; diff --git a/src/llama-context.cpp b/src/llama-context.cpp index fcd3498dcbd2..0ec790392c91 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -1653,7 +1653,9 @@ int llama_context::decode(const llama_batch & batch_inp) { const int64_t n_vocab = vocab.n_tokens(); const bool mtp_embd = cparams.ctx_type == LLAMA_CONTEXT_TYPE_MTP && batch_inp.embd; - const int64_t n_embd = mtp_embd ? hparams.n_embd_out() : hparams.n_embd_inp(); + // DFlash embd batches carry fused target features at the encoder input width + const bool dflash_embd = model.arch == LLM_ARCH_DFLASH && batch_inp.embd; + const int64_t n_embd = mtp_embd ? hparams.n_embd_out() : dflash_embd ? hparams.n_embd_inp_enc() : hparams.n_embd_inp(); // when computing embeddings, all tokens are output const bool output_all = cparams.embeddings; @@ -2303,6 +2305,9 @@ uint32_t llama_context::graph_max_nodes(uint32_t n_tokens) const { model.arch == LLM_ARCH_NANBEIGE || model.arch == LLM_ARCH_MINIMAX_M3) { res = std::max(n_tokens * 40, 32u * model.n_tensors()); + } else if (model.arch == LLM_ARCH_DFLASH && model.hparams.dflash_selector_rank > 0) { + // DFlash2 convolutions and selector are shape work rather than matmuls + res = std::max(1024u, 12u*model.n_tensors()); } else { res = std::max(1024u, 8u*model.n_tensors()); for (const auto & lora : model.loras) { diff --git a/src/llama-ext.h b/src/llama-ext.h index 35d6e58adfa8..b001b62a5cb7 100644 --- a/src/llama-ext.h +++ b/src/llama-ext.h @@ -83,6 +83,8 @@ struct llama_device_memory_data { // TODO: convert to C-style data structure using llama_memory_breakdown = std::map; +LLAMA_API int32_t llama_model_dflash_selector_top_k(const struct llama_model * model); + LLAMA_API int32_t llama_model_n_expert (const struct llama_model * model); LLAMA_API int32_t llama_model_n_devices(const struct llama_model * model); diff --git a/src/llama-hparams.h b/src/llama-hparams.h index 57de808242bd..6a4711f0bdfc 100644 --- a/src/llama-hparams.h +++ b/src/llama-hparams.h @@ -202,6 +202,12 @@ struct llama_hparams { // output embedding dimension (0 = use n_embd) uint32_t n_embd_out_impl = 0; + uint32_t dflash_block_size = 0; + uint32_t dflash_conv_kernel_size = 0; + uint32_t dflash_conv_group_size = 0; + uint32_t dflash_selector_rank = 0; + uint32_t dflash_selector_top_k = 0; + // llama4 smallthinker uint32_t n_moe_layer_step = 0; uint32_t n_no_rope_layer_step = 4; diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 3bf3a22f26af..7e366ca58d17 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -2505,6 +2505,10 @@ int32_t llama_model_n_layer_nextn(const llama_model * model) { return model->hparams.n_layer_nextn; } +int32_t llama_model_dflash_selector_top_k(const llama_model * model) { + return model->hparams.dflash_selector_top_k; +} + int32_t llama_model_n_head(const llama_model * model) { return model->hparams.n_head(); } @@ -2697,6 +2701,10 @@ llama_rope_type llama_model_rope_type(const llama_model * model) { return LLAMA_ROPE_TYPE_NEOX; case LLM_ARCH_DFLASH: + // drafts for M-RoPE targets carry rope sections and follow the target's temporal dim + if (const auto & s = model->hparams.rope_sections; s[0] || s[1] || s[2] || s[3]) { + return LLAMA_ROPE_TYPE_MROPE; + } // DSV4 DSpark drafters use DeepSeek-V4's normal RoPE; legacy DFlash backbones are NeoX return model->hparams.dsv4_hc_mult > 0 ? LLAMA_ROPE_TYPE_NORM : LLAMA_ROPE_TYPE_NEOX; diff --git a/src/llama-model.h b/src/llama-model.h index 1dd0904387af..b1c4377591d1 100644 --- a/src/llama-model.h +++ b/src/llama-model.h @@ -357,6 +357,11 @@ struct llama_layer { struct ggml_tensor * ffn_exp_probs_b = nullptr; struct ggml_tensor * ffn_gate_tid2eid = nullptr; + struct ggml_tensor * dflash_attn_conv_base = nullptr; + struct ggml_tensor * dflash_attn_conv_proj = nullptr; + struct ggml_tensor * dflash_ffn_conv_base = nullptr; + struct ggml_tensor * dflash_ffn_conv_proj = nullptr; + // mamba proj struct ggml_tensor * ssm_in = nullptr; struct ggml_tensor * ssm_x = nullptr; @@ -633,6 +638,10 @@ struct llama_model { struct ggml_tensor * dspark_conf_proj = nullptr; struct ggml_tensor * dspark_conf_proj_b = nullptr; + struct ggml_tensor * dflash_selector_prev = nullptr; + struct ggml_tensor * dflash_selector_next = nullptr; + struct ggml_tensor * dflash_selector_hidden = nullptr; + // unified vector to store target-model extracted layer ids in eagle3, dflash, etc. std::vector target_layer_ids; diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index ca3e1cbd71ff..e8678b32d9d5 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -719,6 +719,18 @@ static ggml_type llama_tensor_get_type(quantize_state_impl & qs, const llama_mod new_type = llama_tensor_get_type_impl(qs, new_type, tensor, params->ftype, tm.category); } + if (!manual && qs.model.arch == LLM_ARCH_DFLASH && + (params->ftype == LLAMA_FTYPE_MOSTLY_IQ1_S || params->ftype == LLAMA_FTYPE_MOSTLY_IQ1_M || + ftype_is_iq1_narrow(params->ftype))) { + const std::string tensor_name(tensor->name); + if (tensor_name.rfind("selector_", 0) == 0 || + tensor_name.find(".attn_conv_") != std::string::npos || + tensor_name.find(".ffn_conv_") != std::string::npos || + tensor_name == "fc.weight") { + new_type = GGML_TYPE_Q8_0; + } + } + // incompatible tensor shapes are handled here - fallback to a compatible type new_type = tensor_type_fallback(qs, tensor, new_type); } diff --git a/src/models/dflash.cpp b/src/models/dflash.cpp index daff6e78f1cf..71ab34d85cc7 100644 --- a/src/models/dflash.cpp +++ b/src/models/dflash.cpp @@ -4,9 +4,24 @@ #include "llama-kv-cache.h" #include "llama-kv-cache-iswa.h" +#include +#include + void llama_model_dflash::load_arch_hparams(llama_model_loader & ml) { ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps); + ml.get_key(LLM_KV_LOGIT_SCALE, hparams.f_logit_scale, false); + hparams.f_final_logit_softcapping = 0.0f; + ml.get_key(LLM_KV_FINAL_LOGIT_SOFTCAPPING, hparams.f_final_logit_softcapping, false); + + // drafts for M-RoPE targets carry degenerate sections [n_rot/2, 0, 0, 0] + ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS, hparams.rope_sections, 4, false); + + ml.get_key(LLM_KV_DFLASH_BLOCK_SIZE, hparams.dflash_block_size, false); + ml.get_key(LLM_KV_DFLASH_CONV_KERNEL_SIZE, hparams.dflash_conv_kernel_size, false); + ml.get_key(LLM_KV_DFLASH_CONV_GROUP_SIZE, hparams.dflash_conv_group_size, false); + ml.get_key(LLM_KV_DFLASH_SELECTOR_RANK, hparams.dflash_selector_rank, false); + ml.get_key(LLM_KV_DFLASH_SELECTOR_TOP_K, hparams.dflash_selector_top_k, false); if (!ml.get_arr(LLM_KV_TARGET_LAYERS, target_layer_ids, false)) { throw std::runtime_error("DFlash model requires 'target_layers' in GGUF metadata"); @@ -96,6 +111,29 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) { LLAMA_LOG_INFO("%s: DFlash with DSpark markov head (rank = %lld)\n", __func__, (long long) dspark_markov_rank); } + const struct ggml_tensor * selector_meta = ml->get_tensor_meta("selector_hidden.weight"); + if (selector_meta) { + const int64_t rank = hparams.dflash_selector_rank; + if (rank <= 0 || hparams.dflash_block_size <= 0 || hparams.dflash_selector_top_k <= 0 || + hparams.dflash_conv_kernel_size <= 0 || hparams.dflash_conv_group_size <= 0) { + throw std::runtime_error("DFlash2 model is missing conv/selector metadata"); + } + if (n_embd % hparams.dflash_conv_group_size != 0) { + throw std::runtime_error("DFlash2 hidden size must be divisible by conv_group_size"); + } + if (n_embd < hparams.dflash_selector_top_k * (hparams.dflash_selector_top_k + 1)) { + throw std::runtime_error("DFlash2 hidden size is too small for the selector lattice"); + } + + dflash_selector_prev = create_tensor(tn(LLM_TENSOR_DFLASH_SELECTOR_PREV, "weight"), { rank, n_vocab }, 0); + dflash_selector_next = create_tensor(tn(LLM_TENSOR_DFLASH_SELECTOR_NEXT, "weight"), { rank, n_vocab }, 0); + dflash_selector_hidden = create_tensor(tn(LLM_TENSOR_DFLASH_SELECTOR_HIDDEN, "weight"), { n_embd, rank }, 0); + + LLAMA_LOG_INFO("%s: DFlash2 conv kernel = %u, group = %u, selector rank = %u, top-k = %u\n", __func__, + hparams.dflash_conv_kernel_size, hparams.dflash_conv_group_size, + hparams.dflash_selector_rank, hparams.dflash_selector_top_k); + } + fc = create_tensor(tn(LLM_TENSOR_FC, "weight"), { n_embd_inp, n_embd }, 0); output_norm_enc = create_tensor(tn(LLM_TENSOR_ENC_OUTPUT_NORM, "weight"), { n_embd }, 0); // encoder hidden_norm (after fc) output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), { n_embd }, 0); // decoder final norm @@ -167,6 +205,16 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) { layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), { n_embd, n_ff }, 0); layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd }, 0); layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), { n_embd, n_ff }, 0); + + if (selector_meta) { + const int64_t kernel = hparams.dflash_conv_kernel_size; + const int64_t groups = n_embd / hparams.dflash_conv_group_size; + const int64_t projected = 2 * kernel * groups; + layer.dflash_attn_conv_base = create_tensor(tn(LLM_TENSOR_DFLASH_ATTN_CONV_BASE, i), { n_embd, kernel, 2 }, 0); + layer.dflash_attn_conv_proj = create_tensor(tn(LLM_TENSOR_DFLASH_ATTN_CONV_PROJ, "weight", i), { n_embd, projected }, 0); + layer.dflash_ffn_conv_base = create_tensor(tn(LLM_TENSOR_DFLASH_FFN_CONV_BASE, i), { n_embd, kernel, 2 }, 0); + layer.dflash_ffn_conv_proj = create_tensor(tn(LLM_TENSOR_DFLASH_FFN_CONV_PROJ, "weight", i), { n_embd, projected }, 0); + } } } @@ -305,11 +353,166 @@ static void build_dspark_markov_head(llm_graph_context & g, const llama_model & ggml_build_forward_expand(g.gf, out); } +static ggml_tensor * build_dflash2_conv( + llm_graph_context & g, + ggml_tensor * hidden, + ggml_tensor * dynamic, + ggml_tensor * base, + int side) { + const auto & hparams = g.hparams; + const int64_t hidden_size = hidden->ne[0]; + const int64_t n_tokens = hidden->ne[1]; + const int64_t n_blocks = g.ubatch.n_seqs_unq; + const int64_t kernel_size = hparams.dflash_conv_kernel_size; + const int64_t group_size = hparams.dflash_conv_group_size; + const int64_t n_groups = hidden_size / group_size; + + GGML_ASSERT(n_blocks > 0 && n_tokens % n_blocks == 0); + GGML_ASSERT(dynamic && base && side >= 0 && side < 2); + + const int64_t block_size = n_tokens / n_blocks; + ggml_context * ctx0 = g.ctx0; + if (!ggml_is_contiguous(hidden) || hidden->ne[1] != n_tokens) { + hidden = ggml_cont_2d(ctx0, hidden, hidden_size, n_tokens); + } + if (!ggml_is_contiguous(dynamic) || dynamic->ne[1] != n_tokens) { + dynamic = ggml_cont_2d(ctx0, dynamic, dynamic->ne[0], n_tokens); + } + ggml_tensor * blocks = ggml_reshape_3d(ctx0, hidden, hidden_size, block_size, n_blocks); + ggml_tensor * coeffs = ggml_reshape_4d(ctx0, dynamic, n_groups, kernel_size, 2, n_tokens); + ggml_tensor * coeffs_side = ggml_view_3d(ctx0, coeffs, n_groups, kernel_size, n_tokens, + coeffs->nb[1], coeffs->nb[3], side * coeffs->nb[2]); + + ggml_tensor * coeff_all = ggml_cont(ctx0, coeffs_side); + coeff_all = ggml_reshape_4d(ctx0, coeff_all, 1, n_groups, kernel_size, n_tokens); + coeff_all = ggml_repeat_4d(ctx0, coeff_all, group_size, n_groups, kernel_size, n_tokens); + + ggml_tensor * base_side = ggml_reshape_4d(ctx0, + ggml_view_1d(ctx0, base, hidden_size * kernel_size, side * base->nb[2]), + group_size, n_groups, kernel_size, 1); + + ggml_tensor * weight_all = ggml_add(ctx0, coeff_all, base_side); + + ggml_tensor * result = nullptr; + for (int64_t tap = 0; tap < kernel_size; ++tap) { + ggml_tensor * values = blocks; + if (tap > 0) { + ggml_tensor * zeros = ggml_fill(ctx0, + ggml_new_tensor_3d(ctx0, hidden->type, hidden_size, std::min(tap, block_size), n_blocks), 0.0f); + if (tap < block_size) { + ggml_tensor * previous = ggml_view_3d(ctx0, blocks, hidden_size, block_size - tap, n_blocks, + blocks->nb[1], blocks->nb[2], 0); + values = ggml_concat(ctx0, zeros, previous, 1); + } else { + values = zeros; + } + } + values = ggml_reshape_2d(ctx0, values, hidden_size, n_tokens); + + ggml_tensor * weight = ggml_reshape_2d(ctx0, + ggml_cont(ctx0, ggml_view_4d(ctx0, weight_all, group_size, n_groups, 1, n_tokens, + weight_all->nb[1], weight_all->nb[2], weight_all->nb[3], tap * weight_all->nb[2])), + hidden_size, n_tokens); + + ggml_tensor * term = ggml_mul(ctx0, weight, values); + result = result ? ggml_add(ctx0, result, term) : term; + } + return result; +} + +// DFlash2 selector: top-k candidates per block position plus pairwise scores, +// packed into the nextn output slot for the CPU-side walk. +static void build_dflash2_selector(llm_graph_context & g, const llama_model & model, ggml_tensor * tokens) { + ggml_context * ctx0 = g.ctx0; + auto & res = g.res; + + const auto & hparams = g.hparams; + const int64_t n_tokens = g.n_tokens; + const int64_t n_embd = g.n_embd; + + const int64_t top_k = hparams.dflash_selector_top_k; + const int64_t rank = hparams.dflash_selector_rank; + const int64_t n_blocks = g.ubatch.n_seqs_unq; + GGML_ASSERT(n_blocks > 0 && n_tokens % n_blocks == 0); + GGML_ASSERT(res->t_logits->ne[1] == n_tokens); + if (!tokens) { + return; + } + + const int64_t tokens_per_block = n_tokens / n_blocks; + const int64_t block_size = std::min(tokens_per_block, hparams.dflash_block_size); + const int64_t row_used = top_k + top_k * top_k; + + ggml_tensor * candidates = ggml_top_k(ctx0, res->t_logits, top_k); + ggml_tensor * logits_rows = ggml_reshape_3d(ctx0, res->t_logits, 1, res->t_logits->ne[0], n_tokens); + ggml_tensor * unary = ggml_reshape_2d(ctx0, + ggml_get_rows(ctx0, logits_rows, candidates), top_k, n_tokens); + ggml_tensor * gate = g.build_lora_mm(model.dflash_selector_hidden, res->t_embd); + + ggml_tensor * cand_blk = ggml_reshape_3d(ctx0, candidates, top_k, tokens_per_block, n_blocks); + ggml_tensor * unary_blk = ggml_reshape_3d(ctx0, unary, top_k, tokens_per_block, n_blocks); + ggml_tensor * gate_blk = ggml_reshape_3d(ctx0, gate, rank, tokens_per_block, n_blocks); + + auto score_run = [&](int64_t beg_pos, int64_t n_pos, ggml_tensor * pred_ids) { + ggml_tensor * cand_run = ggml_cont(ctx0, ggml_view_3d(ctx0, cand_blk, top_k, n_pos, n_blocks, + cand_blk->nb[1], cand_blk->nb[2], beg_pos * cand_blk->nb[1])); + ggml_tensor * unary_run = ggml_cont(ctx0, ggml_view_3d(ctx0, unary_blk, top_k, n_pos, n_blocks, + unary_blk->nb[1], unary_blk->nb[2], beg_pos * unary_blk->nb[1])); + ggml_tensor * gate_run = ggml_cont(ctx0, ggml_view_3d(ctx0, gate_blk, rank, n_pos, n_blocks, + gate_blk->nb[1], gate_blk->nb[2], beg_pos * gate_blk->nb[1])); + + const int64_t n_pred = pred_ids->ne[0] / (n_pos * n_blocks); + + ggml_tensor * successor = ggml_reshape_4d(ctx0, + ggml_get_rows(ctx0, model.dflash_selector_next, ggml_reshape_1d(ctx0, cand_run, top_k * n_pos * n_blocks)), + rank, top_k, n_pos, n_blocks); + ggml_tensor * predecessor = ggml_reshape_4d(ctx0, + ggml_get_rows(ctx0, model.dflash_selector_prev, pred_ids), + rank, n_pred, n_pos, n_blocks); + + ggml_tensor * gate_bcast = ggml_reshape_4d(ctx0, gate_run, rank, 1, n_pos, n_blocks); + ggml_tensor * cond = ggml_mul(ctx0, predecessor, ggml_repeat(ctx0, gate_bcast, predecessor)); + ggml_tensor * score = ggml_mul_mat(ctx0, successor, cond); + if (n_pred == 1) { + score = ggml_repeat_4d(ctx0, score, top_k, top_k, n_pos, n_blocks); + } + ggml_tensor * unary_bcast = ggml_reshape_4d(ctx0, unary_run, top_k, 1, n_pos, n_blocks); + score = ggml_add(ctx0, score, ggml_repeat(ctx0, unary_bcast, score)); + + ggml_tensor * row = ggml_concat(ctx0, + ggml_cast(ctx0, cand_run, GGML_TYPE_F32), + ggml_reshape_3d(ctx0, score, top_k * top_k, n_pos, n_blocks), 0); + return ggml_pad(ctx0, row, n_embd - row_used, 0, 0, 0); + }; + + ggml_tensor * packed = ggml_fill(ctx0, + ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, n_embd, 1, n_blocks), 0.0f); + + if (block_size > 1) { + ggml_tensor * anchor_ids = ggml_cont_1d(ctx0, + ggml_view_2d(ctx0, tokens, 1, n_blocks, tokens_per_block * tokens->nb[0], 0), n_blocks); + packed = ggml_concat(ctx0, packed, score_run(1, 1, anchor_ids), 1); + } + if (block_size > 2) { + ggml_tensor * prev_ids = ggml_reshape_1d(ctx0, + ggml_cont(ctx0, ggml_view_3d(ctx0, cand_blk, top_k, block_size - 2, n_blocks, + cand_blk->nb[1], cand_blk->nb[2], cand_blk->nb[1])), + top_k * (block_size - 2) * n_blocks); + packed = ggml_concat(ctx0, packed, score_run(2, block_size - 2, prev_ids), 1); + } + + packed = ggml_reshape_2d(ctx0, packed, n_embd, block_size * n_blocks); + g.cb(packed, "dflash2_lattice", -1); + res->t_h_nextn = packed; + ggml_build_forward_expand(g.gf, packed); +} + // DFlash decoder, dual-mode by batch type: // * embd batch -> fused target features: project + inject K/V into the cache. // * token batch -> noise-block diffusion: attend over [committed, MASK...] to generate draft tokens template <> llama_model_dflash::graph::graph(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) { + const int64_t n_embd_inp = hparams.n_embd_inp_enc(); const int64_t n_embd_head = hparams.n_embd_head_v(); GGML_ASSERT(n_embd_head == hparams.n_embd_head_k()); @@ -329,18 +532,35 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra const float kq_scale = 1.0f/sqrtf(float(n_embd_head)); + int sections[4]; + std::copy(std::begin(hparams.rope_sections), std::begin(hparams.rope_sections) + 4, sections); + + auto build_rope = [&](ggml_tensor * cur, ggml_tensor * pos) { + return rope_type == GGML_ROPE_TYPE_MROPE + ? ggml_rope_multi(ctx0, cur, pos, nullptr, + n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow) + : ggml_rope_ext(ctx0, cur, pos, nullptr, + n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + }; + // KV cache injection if (ubatch.embd) { - auto inp = std::make_unique(n_embd); + auto inp = std::make_unique(n_embd_inp); - inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_embd, n_tokens); + inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_embd_inp, n_tokens); ggml_set_input(inp->embd); - ggml_tensor * inp_g = inp->embd; - cb(inp_g, "inp_g_embeddings", -1); + ggml_tensor * inp_target = inp->embd; + cb(inp_target, "inp_target_features", -1); res->add_input(std::move(inp)); + ggml_tensor * inp_g = build_lora_mm(model.fc, inp_target); + inp_g = build_norm(inp_g, model.output_norm_enc, NULL, LLM_NORM_RMS, -1); + cb(inp_g, "inp_g_embeddings", -1); + for (int il = 0; il < n_layer; ++il) { const auto & layer = model.layers[il]; @@ -351,11 +571,7 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra Vcur = ggml_reshape_3d(ctx0, Vcur, n_embd_head, n_head_kv, n_tokens); Kcur = build_norm(Kcur, layer.attn_k_norm, NULL, LLM_NORM_RMS, il); - Kcur = ggml_rope_ext( - ctx0, Kcur, inp_pos, nullptr, - n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, - ext_factor, attn_factor, beta_fast, beta_slow - ); + Kcur = build_rope(Kcur, inp_pos); cb(Kcur, "Kcur_injected", il); cb(Vcur, "Vcur_injected", il); @@ -409,6 +625,7 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra inp->tokens = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_tokens); ggml_set_input(inp->tokens); + res->t_inp_tokens = inp->tokens; ggml_tensor * inp_tokens = inp->tokens; @@ -423,6 +640,13 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra ggml_tensor * noise_norm = build_norm(inpL, layer.attn_norm, NULL, LLM_NORM_RMS, il); cb(noise_norm, "noise_norm", il); + ggml_tensor * attn_dynamic = nullptr; + if (layer.dflash_attn_conv_proj) { + attn_dynamic = build_lora_mm(layer.dflash_attn_conv_proj, noise_norm); + noise_norm = build_dflash2_conv(*this, noise_norm, attn_dynamic, layer.dflash_attn_conv_base, 0); + cb(noise_norm, "attn_conv_in", il); + } + ggml_tensor * Qcur = build_lora_mm(layer.wq, noise_norm); ggml_tensor * Kcur = build_lora_mm(layer.wk, noise_norm); ggml_tensor * Vcur = build_lora_mm(layer.wv, noise_norm); @@ -434,16 +658,8 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra Qcur = build_norm(Qcur, layer.attn_q_norm, NULL, LLM_NORM_RMS, il); Kcur = build_norm(Kcur, layer.attn_k_norm, NULL, LLM_NORM_RMS, il); - Qcur = ggml_rope_ext( - ctx0, Qcur, inp_pos, nullptr, - n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, - ext_factor, attn_factor, beta_fast, beta_slow - ); - Kcur = ggml_rope_ext( - ctx0, Kcur, inp_pos, nullptr, - n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, - ext_factor, attn_factor, beta_fast, beta_slow - ); + Qcur = build_rope(Qcur, inp_pos); + Kcur = build_rope(Kcur, inp_pos); cb(Qcur, "Qcur", il); cb(Kcur, "Kcur", il); cb(Vcur, "Vcur", il); @@ -453,12 +669,24 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra ? build_attn(inp_attn_iswa, layer.wo, NULL, NULL, Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il) : build_attn(inp_attn, layer.wo, NULL, NULL, Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il); + if (attn_dynamic) { + cur = build_dflash2_conv(*this, cur, attn_dynamic, layer.dflash_attn_conv_base, 1); + cb(cur, "attn_conv_out", il); + } + ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpL); cb(ffn_inp, "ffn_inp", il); cur = build_norm(ffn_inp, layer.ffn_norm, NULL, LLM_NORM_RMS, il); cb(cur, "ffn_norm", il); + ggml_tensor * ffn_dynamic = nullptr; + if (layer.dflash_ffn_conv_proj) { + ffn_dynamic = build_lora_mm(layer.dflash_ffn_conv_proj, cur); + cur = build_dflash2_conv(*this, cur, ffn_dynamic, layer.dflash_ffn_conv_base, 0); + cb(cur, "ffn_conv_in", il); + } + cur = build_ffn(cur, layer.ffn_up, NULL, NULL, layer.ffn_gate, NULL, NULL, @@ -467,6 +695,11 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra LLM_FFN_SILU, LLM_FFN_PAR, il); cb(cur, "ffn_out", il); + if (ffn_dynamic) { + cur = build_dflash2_conv(*this, cur, ffn_dynamic, layer.dflash_ffn_conv_base, 1); + cb(cur, "ffn_conv_out", il); + } + cur = ggml_add(ctx0, cur, ffn_inp); cb(cur, "l_out", il); @@ -488,6 +721,19 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra } cur = build_lora_mm(output, cur); + + // DFlash2 selector reads these logits in-graph, so apply the target's output transforms here + if (model.dflash_selector_hidden) { + if (hparams.f_logit_scale != 0.0f) { + cur = ggml_scale(ctx0, cur, hparams.f_logit_scale); + } + if (hparams.f_final_logit_softcapping > 0.0f) { + cur = ggml_scale(ctx0, cur, 1.0f / hparams.f_final_logit_softcapping); + cur = ggml_tanh(ctx0, cur); + cur = ggml_scale(ctx0, cur, hparams.f_final_logit_softcapping); + } + } + cb(cur, "result_output", -1); res->t_logits = cur; @@ -497,6 +743,10 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra if (model.dspark_markov_w1) { build_dspark_markov_head(*this, model, inp_tokens); } + + if (model.dflash_selector_hidden) { + build_dflash2_selector(*this, model, inp_tokens); + } } // DSV4 DSpark decoder, dual-mode by batch type (see the DFlash decoder above): @@ -504,6 +754,7 @@ llama_model_dflash::graph::graph(const llama_model & model, const llm_gra // * token batch -> noise block through 3 full DSV4 stages (hc + MLA + MoE), markov + confidence heads llama_model_dflash::graph_dsv4::graph_dsv4(const llama_model & model, const llm_graph_params & params) : llama_model_deepseek4::graph(params) { + const int64_t n_embd_inp = hparams.n_embd_inp_enc(); const int64_t n_embd_head = hparams.n_embd_head_k(); const int64_t n_embd_head_rope = hparams.n_rot(); const int64_t n_embd_head_nope = n_embd_head - n_embd_head_rope; @@ -514,16 +765,20 @@ llama_model_dflash::graph_dsv4::graph_dsv4(const llama_model & model, const llm_ // KV cache injection: fused target features from the encoder if (ubatch.embd) { - auto inp = std::make_unique(n_embd); + auto inp = std::make_unique(n_embd_inp); - inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_embd, n_tokens); + inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_embd_inp, n_tokens); ggml_set_input(inp->embd); - ggml_tensor * inp_g = inp->embd; - cb(inp_g, "inp_g_embeddings", -1); + ggml_tensor * inp_target = inp->embd; + cb(inp_target, "inp_target_features", -1); res->add_input(std::move(inp)); + ggml_tensor * inp_g = build_lora_mm(model.fc, inp_target); + inp_g = build_norm(inp_g, model.output_norm_enc, nullptr, LLM_NORM_RMS, -1); + cb(inp_g, "inp_g_embeddings", -1); + for (int il = 0; il < n_layer; ++il) { const auto & layer = model.layers[il]; diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp index d1c19a4b70af..fa65a8580b71 100644 --- a/tests/test-quantize-fns.cpp +++ b/tests/test-quantize-fns.cpp @@ -11,6 +11,10 @@ #include #if defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) +extern "C" void ggml_vec_dot_iq1_xs_q8_K_generic(int n, float * s, size_t bs, + const void * vx, size_t bx, const void * vy, size_t by, int nrc); +extern "C" void ggml_vec_dot_iq1_xxs_q8_K_generic(int n, float * s, size_t bs, + const void * vx, size_t bx, const void * vy, size_t by, int nrc); extern "C" void ggml_vec_dot_iq1_xxxs_q8_K_generic(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc); #endif @@ -206,39 +210,38 @@ static int test_vec_dot_q(bool verbose) { } #if defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) -static int test_iq1_xxxs_vec_dot_matches_generic(bool verbose) { +using iq1_generic_dot_fn = void (*)(int, float *, size_t, const void *, size_t, const void *, size_t, int); + +static int test_iq1_vec_dot_matches_generic(enum ggml_type type, iq1_generic_dot_fn generic, const char * name, bool verbose) { const size_t test_size = 32 * 128; - const auto * iq1 = ggml_get_type_traits_cpu(GGML_TYPE_IQ1_XXXS); + const auto * iq1 = ggml_get_type_traits_cpu(type); const auto * q8 = ggml_get_type_traits_cpu(iq1->vec_dot_type); + const size_t blk_size = ggml_type_size(type); std::vector activations(test_size); - std::vector quant_weights(ggml_row_size(GGML_TYPE_IQ1_XXXS, test_size)); + std::vector quant_weights(ggml_row_size(type, test_size)); std::vector quant_activations(2 * test_size); generate_data(1.75f, test_size, activations.data()); for (size_t block = 0; block < test_size / 256; ++block) { - uint8_t * data = quant_weights.data() + 38 * block; + uint8_t * data = quant_weights.data() + blk_size * block; data[0] = 0x00; data[1] = 0x3c; - for (size_t i = 0; i < 32; ++i) { - data[2 + i] = (uint8_t)(block * 37 + i * 11 + 3); - } - for (size_t i = 0; i < 4; ++i) { - data[34 + i] = (uint8_t)(((2 * i + 1) << 4) | (2 * i)); + for (size_t i = 2; i < blk_size; ++i) { + data[i] = (uint8_t)(block * 37 + i * 11 + 3); } } q8->from_float(activations.data(), quant_activations.data(), test_size); float optimized = 0.0f; - float generic = 0.0f; + float generic_s = 0.0f; iq1->vec_dot(test_size, &optimized, 0, quant_weights.data(), 0, quant_activations.data(), 0, 1); - ggml_vec_dot_iq1_xxxs_q8_K_generic(test_size, &generic, 0, - quant_weights.data(), 0, - quant_activations.data(), 0, 1); - const bool failed = fabsf(optimized - generic) > 1e-5f; + generic(test_size, &generic_s, 0, quant_weights.data(), 0, + quant_activations.data(), 0, 1); + const bool failed = fabsf(optimized - generic_s) > 1e-5f; if (failed || verbose) { - printf("iq1_xxxs optimized vs generic: %s (generic=%f optimized=%f)\n", - RESULT_STR[failed], generic, optimized); + printf("%s optimized vs generic: %s (generic=%f optimized=%f)\n", + name, RESULT_STR[failed], generic_s, optimized); } return failed; } @@ -266,7 +269,9 @@ int main(int argc, char * argv[]) { num_failed += test_vec_dot_f32(verbose); num_failed += test_vec_dot_q(verbose); #if defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) - num_failed += test_iq1_xxxs_vec_dot_matches_generic(verbose); + num_failed += test_iq1_vec_dot_matches_generic(GGML_TYPE_IQ1_XS, ggml_vec_dot_iq1_xs_q8_K_generic, "iq1_xs", verbose); + num_failed += test_iq1_vec_dot_matches_generic(GGML_TYPE_IQ1_XXS, ggml_vec_dot_iq1_xxs_q8_K_generic, "iq1_xxs", verbose); + num_failed += test_iq1_vec_dot_matches_generic(GGML_TYPE_IQ1_XXXS, ggml_vec_dot_iq1_xxxs_q8_K_generic, "iq1_xxxs", verbose); #endif if (num_failed || verbose) {