From 5a32f7b66ef6cfb3e60deea26e3454cc6ad3438c Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Fri, 21 Aug 2026 19:52:34 +0200 Subject: [PATCH 01/16] model: add dots3-note (#27060) * text: conversion * init impl * address review comments * fix rope * move to a new llama_kv_cache_dsa_iswa --- conversion/__init__.py | 3 + conversion/dots3.py | 195 +++++++++++++ gguf-py/gguf/constants.py | 43 +++ gguf-py/gguf/gguf_writer.py | 9 + gguf-py/gguf/tensor_mapping.py | 1 + src/CMakeLists.txt | 1 + src/llama-arch.cpp | 5 + src/llama-arch.h | 4 + src/llama-graph.cpp | 66 ++++- src/llama-graph.h | 35 +++ src/llama-hparams.h | 5 + src/llama-kv-cache-dsa-iswa.cpp | 341 +++++++++++++++++++++++ src/llama-kv-cache-dsa-iswa.h | 134 +++++++++ src/llama-kv-cache.cpp | 3 +- src/llama-model-saver.cpp | 1 + src/llama-model.cpp | 60 +++- src/llama-model.h | 1 + src/models/dots3note.cpp | 480 ++++++++++++++++++++++++++++++++ src/models/models.h | 12 + tests/test-llama-archs.cpp | 22 +- 20 files changed, 1412 insertions(+), 9 deletions(-) create mode 100644 conversion/dots3.py create mode 100644 src/llama-kv-cache-dsa-iswa.cpp create mode 100644 src/llama-kv-cache-dsa-iswa.h create mode 100644 src/models/dots3note.cpp diff --git a/conversion/__init__.py b/conversion/__init__.py index 4b8817ead..b4afdf0f8 100644 --- a/conversion/__init__.py +++ b/conversion/__init__.py @@ -64,6 +64,9 @@ TEXT_MODEL_MAP: dict[str, str] = { "DistilBertForSequenceClassification": "bert", "DistilBertModel": "bert", "Dots1ForCausalLM": "dots1", + "Dots3NoteForCausalLM": "dots3", + "Dots3NoteForConditionalGeneration": "dots3", + "Dots3NoteTextForCausalLM": "dots3", "DotsOCRForCausalLM": "qwen", "DreamModel": "dream", "Ernie4_5ForCausalLM": "ernie", diff --git a/conversion/dots3.py b/conversion/dots3.py new file mode 100644 index 000000000..7c36b8482 --- /dev/null +++ b/conversion/dots3.py @@ -0,0 +1,195 @@ +from __future__ import annotations + +import math +import re + +from typing import TYPE_CHECKING, Callable, Iterable + +if TYPE_CHECKING: + from torch import Tensor + +from .base import ModelBase, gguf + +from .deepseek import DeepseekV2Model + + +@ModelBase.register("Dots3NoteForCausalLM", "Dots3NoteForConditionalGeneration", "Dots3NoteTextForCausalLM") +class Dots3NoteModel(DeepseekV2Model): + model_arch = gguf.MODEL_ARCH.DOTS3NOTE + skip_mtp = False + supports_mtp_export = True + + # trunk layer count, stashed before indexing for filter_tensors (mirrors DeepseekV32Model) + _n_main_layers: int | None = None + + def index_tensors(self, remote_hf_model_id: str | None = None): + type(self)._n_main_layers = self.hparams["num_hidden_layers"] + return super().index_tensors(remote_hf_model_id=remote_hf_model_id) + + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + + hparams = self.hparams + + # config file doesn't specify MTP block, detect it from model weight + self.n_nextn = 1 if "model.mtp.embed_tokens.weight" in self.model_tensors else 0 + if self.n_nextn: + self.block_count += self.n_nextn + self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count) + + self.layer_types = hparams["layer_types"] + if len(self.layer_types) < hparams["num_hidden_layers"]: + raise ValueError("layer_types is shorter than num_hidden_layers") + + if hparams.get("use_dsa", True) is not True: + raise ValueError("dots3-note conversion requires use_dsa=true") + if hparams.get("normalization", "RMSNorm") != "RMSNorm" or hparams.get("final_norm", "RMSNorm") != "RMSNorm": + raise ValueError("dots3-note conversion only supports RMSNorm") + if hparams.get("k_rope_only_layernorm", True) is not True: + raise ValueError("dots3-note conversion requires k_rope_only_layernorm=true") + if hparams.get("topk_method", "noaux_tc") != "noaux_tc" or hparams.get("scoring_func") != "sigmoid": + raise ValueError("dots3-note conversion only supports noaux_tc/sigmoid expert gating") + if hparams.get("n_group", 1) != 1 or hparams.get("topk_group", 1) != 1: + raise ValueError("dots3-note conversion does not support grouped expert routing") + if hparams.get("use_dynamic_rsf", False) or hparams.get("moe_gating_fp32", False): + raise ValueError("dots3-note conversion does not support use_dynamic_rsf/moe_gating_fp32") + for key in ("attention_gate_type", "swa_attention_gate_type"): + if hparams.get(key, "headwise") != "headwise": + raise ValueError(f"dots3-note conversion only supports headwise attention gate, got {key}={hparams.get(key)!r}") + if hparams["swa_qk_nope_head_dim"] + hparams["swa_qk_rope_head_dim"] != hparams.get("swa_head_dim", 256): + raise ValueError("swa_head_dim must equal swa_qk_nope_head_dim + swa_qk_rope_head_dim") + if hparams["swa_qk_rope_head_dim"] != hparams["qk_rope_head_dim"]: + # both layer kinds share a single rope_dimension_count + raise ValueError("swa_qk_rope_head_dim must match qk_rope_head_dim") + + self.apply_lora_rescale = hparams.get("apply_mla_qkv_lora_rescale", False) + + def _is_swa_layer(self, bid: int) -> bool: + if bid >= self.hparams["num_hidden_layers"]: + # note: the NextN/MTP block uses the sliding-attention MLA + return True + return self.layer_types[bid] == "sliding_attention" + + def set_vocab(self): + from transformers import AutoTokenizer + tokenizer = AutoTokenizer.from_pretrained(self.dir_model) + special_vocab = gguf.SpecialVocab(self.dir_model, load_merges=True) + tokens, toktypes, tokpre = self.get_vocab_base() + self.gguf_writer.add_tokenizer_model("gpt2") + self.gguf_writer.add_tokenizer_pre(tokpre) + self.gguf_writer.add_token_list(tokens) + self.gguf_writer.add_token_types(toktypes) + special_vocab._set_special_token("eot", tokenizer.get_added_vocab()["<|endofassistant|>"]) # ty: ignore[unresolved-attribute] + special_vocab.add_to_gguf(self.gguf_writer) + + @classmethod + def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None: + if (titem := super().filter_tensors(item)) is None: + return None + name, gen = titem + if name.startswith(("vision_encoder.", "audio_encoder.")): + return None + + assert cls._n_main_layers is not None + is_mtp = name.startswith("model.mtp.") or \ + ((m := re.match(r"model\.layers\.(\d+)\.", name)) is not None and int(m.group(1)) >= cls._n_main_layers) + + # --no-mtp: drop the NextN/MTP block; --mtp: keep only that block plus the shared embeddings/norm/lm_head + if is_mtp and cls.no_mtp: + return None + if cls.mtp_only and not is_mtp and name not in ( + "model.embed_tokens.weight", "model.norm.weight", "lm_head.weight", + ): + return None + + return name, gen + + def set_gguf_parameters(self): + hparams = self.hparams + + # head_count is a per-layer array because the two layer kinds have different head counts + n_layer = hparams["num_hidden_layers"] + hparams["num_attention_heads"] = [ + hparams["swa_num_attention_heads"] if self._is_swa_layer(il) else hparams["num_attention_heads"] + for il in range(self.block_count) + ] + + # prevent the base class from emitting key/value_length from the unused head_dim + hparams.pop("head_dim", None) + + super().set_gguf_parameters() + + # MLA geometry of the sliding-window layers (rope.freq_base_swa is emitted by the base class) + swa_kv_lora_rank = hparams["swa_kv_lora_rank"] + self.gguf_writer.add_sliding_window(hparams["sliding_window_size"]) + self.gguf_writer.add_sliding_window_pattern([self._is_swa_layer(il) for il in range(n_layer)]) + self.gguf_writer.add_kv_lora_rank_swa(swa_kv_lora_rank) + self.gguf_writer.add_key_length_swa(swa_kv_lora_rank + hparams["swa_qk_rope_head_dim"]) + self.gguf_writer.add_value_length_swa(swa_kv_lora_rank) + self.gguf_writer.add_key_length_mla_swa(hparams["swa_qk_nope_head_dim"] + hparams["swa_qk_rope_head_dim"]) + self.gguf_writer.add_value_length_mla_swa(hparams["swa_v_head_dim"]) + if hparams["swa_q_lora_rank"] != hparams["q_lora_rank"]: + raise ValueError("dots3-note conversion assumes a shared q_lora_rank for both layer kinds") + + if self.n_nextn: + self.gguf_writer.add_nextn_predict_layers(self.n_nextn) + + # DSA indexer (full-attention layers only) + self.gguf_writer.add_indexer_head_count(hparams["index_n_heads"]) + self.gguf_writer.add_indexer_key_length(hparams["index_head_dim"]) + self.gguf_writer.add_indexer_top_k(hparams["index_topk"]) + self.gguf_writer.add_indexer_types([not self._is_swa_layer(il) for il in range(n_layer)]) + + def prepare_metadata(self, vocab_only: bool): + from_dir = self.fname_out.is_dir() + super().prepare_metadata(vocab_only=vocab_only) + + if not self.mtp_only or not from_dir: + return + + output_type: str = self.ftype.name.partition("_")[2] + fname_default: str = gguf.naming_convention( + self.metadata.name, self.metadata.basename, self.metadata.finetune, + self.metadata.version, size_label=None, output_type=output_type, model_type=None) + self.fname_out = self.fname_out.parent / f"mtp-{fname_default}.gguf" + + def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: + # move the MTP token embedding into the NextN block so the standard nextn mapping picks it up + if name == "model.mtp.embed_tokens.weight": + name = f"model.layers.{self.hparams['num_hidden_layers']}.embed_tokens.weight" + bid = self.hparams["num_hidden_layers"] + + # fold the activation rescale sqrt(n_embd/lora_rank) into the preceding RMSNorm weight + # this also covers the indexer wq_b, which reads the same rescaled q_lora activation + if self.apply_lora_rescale and bid is not None: + if name.endswith("q_a_layernorm.weight"): + data_torch = data_torch * math.sqrt(self.hparams["hidden_size"] / self.hparams["q_lora_rank"]) + elif name.endswith("kv_a_layernorm.weight"): + rank = self.hparams["swa_kv_lora_rank"] if self._is_swa_layer(bid) else self.hparams["kv_lora_rank"] + data_torch = data_torch * math.sqrt(self.hparams["hidden_size"] / rank) + + # MLA absorption: split kv_b_proj into k_b (transposed) and v_b, per-layer-kind geometry + if name.endswith("kv_b_proj.weight"): + assert bid is not None + if self._is_swa_layer(bid): + n_head = self.hparams["swa_num_attention_heads"] + qk_nope_head_dim = self.hparams["swa_qk_nope_head_dim"] + v_head_dim = self.hparams["swa_v_head_dim"] + else: + n_head = self.hparams["num_attention_heads"] + qk_nope_head_dim = self.hparams["qk_nope_head_dim"] + v_head_dim = self.hparams["v_head_dim"] + if isinstance(n_head, list): # set_gguf_parameters turns this into a per-layer array + n_head = n_head[bid] + + assert data_torch.shape[0] == n_head * (qk_nope_head_dim + v_head_dim) + + kv_b = data_torch.view(n_head, qk_nope_head_dim + v_head_dim, data_torch.shape[-1]) + k_b, v_b = kv_b.split([qk_nope_head_dim, v_head_dim], dim=1) + k_b = k_b.transpose(1, 2) + + yield from ModelBase.modify_tensors(self, k_b, name.replace("kv_b_proj", "k_b_proj"), bid) + yield from ModelBase.modify_tensors(self, v_b, name.replace("kv_b_proj", "v_b_proj"), bid) + return + + yield from super().modify_tensors(data_torch, name, bid) diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index fad8d1fd8..886253d88 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -205,6 +205,9 @@ class Keys: VALUE_LENGTH_MLA = "{arch}.attention.value_length_mla" KEY_LENGTH_SWA = "{arch}.attention.key_length_swa" VALUE_LENGTH_SWA = "{arch}.attention.value_length_swa" + KEY_LENGTH_MLA_SWA = "{arch}.attention.key_length_mla_swa" + VALUE_LENGTH_MLA_SWA = "{arch}.attention.value_length_mla_swa" + KV_LORA_RANK_SWA = "{arch}.attention.kv_lora_rank_swa" SHARED_KV_LAYERS = "{arch}.attention.shared_kv_layers" SLIDING_WINDOW_PATTERN = "{arch}.attention.sliding_window_pattern" TEMPERATURE_SCALE = "{arch}.attention.temperature_scale" @@ -558,6 +561,7 @@ class MODEL_ARCH(IntEnum): BAILINGMOE2 = auto() BAILINGMOE3 = auto() DOTS1 = auto() + DOTS3NOTE = auto() ARCEE = auto() AFMOE = auto() LAGUNA = auto() @@ -1275,6 +1279,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = { MODEL_ARCH.BAILINGMOE2: "bailingmoe2", MODEL_ARCH.BAILINGMOE3: "bailingmoe3", MODEL_ARCH.DOTS1: "dots1", + MODEL_ARCH.DOTS3NOTE: "dots3note", MODEL_ARCH.ARCEE: "arcee", MODEL_ARCH.AFMOE: "afmoe", MODEL_ARCH.LAGUNA: "laguna", @@ -4334,6 +4339,44 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = { MODEL_TENSOR.FFN_UP_EXP, MODEL_TENSOR.FFN_UP_SHEXP, ], + MODEL_ARCH.DOTS3NOTE: [ + MODEL_TENSOR.TOKEN_EMBD, + MODEL_TENSOR.OUTPUT_NORM, + MODEL_TENSOR.OUTPUT, + MODEL_TENSOR.ATTN_NORM, + MODEL_TENSOR.ATTN_Q_A, + MODEL_TENSOR.ATTN_Q_B, + MODEL_TENSOR.ATTN_KV_A_MQA, + MODEL_TENSOR.ATTN_K_B, + MODEL_TENSOR.ATTN_V_B, + MODEL_TENSOR.ATTN_Q_A_NORM, + MODEL_TENSOR.ATTN_KV_A_NORM, + MODEL_TENSOR.ATTN_K_NORM, + MODEL_TENSOR.ATTN_GATE, + MODEL_TENSOR.ATTN_OUT, + MODEL_TENSOR.FFN_NORM, + MODEL_TENSOR.FFN_GATE, + MODEL_TENSOR.FFN_DOWN, + MODEL_TENSOR.FFN_UP, + MODEL_TENSOR.FFN_GATE_INP, + MODEL_TENSOR.FFN_EXP_PROBS_B, + MODEL_TENSOR.FFN_GATE_EXP, + MODEL_TENSOR.FFN_DOWN_EXP, + MODEL_TENSOR.FFN_UP_EXP, + MODEL_TENSOR.FFN_GATE_SHEXP, + MODEL_TENSOR.FFN_DOWN_SHEXP, + MODEL_TENSOR.FFN_UP_SHEXP, + MODEL_TENSOR.INDEXER_K_NORM, + MODEL_TENSOR.INDEXER_PROJ, + MODEL_TENSOR.INDEXER_ATTN_K, + MODEL_TENSOR.INDEXER_ATTN_Q_B, + # NextN/MTP tensors - preserved but unused + MODEL_TENSOR.NEXTN_EH_PROJ, + MODEL_TENSOR.NEXTN_EMBED_TOKENS, + MODEL_TENSOR.NEXTN_ENORM, + MODEL_TENSOR.NEXTN_HNORM, + MODEL_TENSOR.NEXTN_SHARED_HEAD_NORM, + ], MODEL_ARCH.ARCEE: [ MODEL_TENSOR.TOKEN_EMBD, MODEL_TENSOR.OUTPUT_NORM, diff --git a/gguf-py/gguf/gguf_writer.py b/gguf-py/gguf/gguf_writer.py index 16ae9f999..496882025 100644 --- a/gguf-py/gguf/gguf_writer.py +++ b/gguf-py/gguf/gguf_writer.py @@ -785,6 +785,15 @@ class GGUFWriter: def add_key_length_swa(self, length: int) -> None: self.add_uint32(Keys.Attention.KEY_LENGTH_SWA.format(arch=self.arch), length) + def add_key_length_mla_swa(self, length: int) -> None: + self.add_uint32(Keys.Attention.KEY_LENGTH_MLA_SWA.format(arch=self.arch), length) + + def add_value_length_mla_swa(self, length: int) -> None: + self.add_uint32(Keys.Attention.VALUE_LENGTH_MLA_SWA.format(arch=self.arch), length) + + def add_kv_lora_rank_swa(self, length: int) -> None: + self.add_uint32(Keys.Attention.KV_LORA_RANK_SWA.format(arch=self.arch), length) + def add_value_length_swa(self, length: int) -> None: self.add_uint32(Keys.Attention.VALUE_LENGTH_SWA.format(arch=self.arch), length) diff --git a/gguf-py/gguf/tensor_mapping.py b/gguf-py/gguf/tensor_mapping.py index a0571ccd3..1ff4b61d9 100644 --- a/gguf-py/gguf/tensor_mapping.py +++ b/gguf-py/gguf/tensor_mapping.py @@ -723,6 +723,7 @@ class TensorNameMap: "model.layers.layers.{bid}.mixer.k", # plamo2 "model.layers.layers.{bid}.mixer.k_norm", # plamo3 "layers.{bid}.self_attn.k_norm", # qwen3-embedding + "model.layers.{bid}.self_attn.k_rope_only_layernorm", # dots3note "model.layers.{bid}.attention.key_layernorm", # apertus ), diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 39ba3061f..c6df19f2e 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -25,6 +25,7 @@ add_library(llama llama-kv-cache.cpp llama-kv-cache-iswa.cpp llama-kv-cache-dsa.cpp + llama-kv-cache-dsa-iswa.cpp llama-kv-cache-msa.cpp llama-kv-cache-dsv4.cpp llama-memory.cpp diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index c9b504c33..60b8e461c 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -110,6 +110,7 @@ static const std::map LLM_ARCH_NAMES = { { LLM_ARCH_BAILINGMOE2, "bailingmoe2" }, { LLM_ARCH_BAILINGMOE3, "bailingmoe3" }, { LLM_ARCH_DOTS1, "dots1" }, + { LLM_ARCH_DOTS3NOTE, "dots3note" }, { LLM_ARCH_ARCEE, "arcee" }, { LLM_ARCH_AFMOE, "afmoe" }, { LLM_ARCH_LAGUNA, "laguna" }, @@ -273,6 +274,9 @@ static const std::map LLM_KV_NAMES = { { LLM_KV_ATTENTION_VALUE_LENGTH_MLA, "%s.attention.value_length_mla" }, { LLM_KV_ATTENTION_KEY_LENGTH_SWA, "%s.attention.key_length_swa" }, { LLM_KV_ATTENTION_VALUE_LENGTH_SWA, "%s.attention.value_length_swa" }, + { LLM_KV_ATTENTION_KEY_LENGTH_MLA_SWA, "%s.attention.key_length_mla_swa" }, + { LLM_KV_ATTENTION_VALUE_LENGTH_MLA_SWA, "%s.attention.value_length_mla_swa" }, + { LLM_KV_ATTENTION_KV_LORA_RANK_SWA, "%s.attention.kv_lora_rank_swa" }, { LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, "%s.attention.indexer.head_count" }, { LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, "%s.attention.indexer.key_length" }, { LLM_KV_ATTENTION_INDEXER_TOP_K, "%s.attention.indexer.top_k" }, @@ -1056,6 +1060,7 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) { case LLM_ARCH_DEEPSEEK2: case LLM_ARCH_DEEPSEEK32: case LLM_ARCH_DEEPSEEK4: + case LLM_ARCH_DOTS3NOTE: case LLM_ARCH_GLM_DSA: case LLM_ARCH_BITNET: case LLM_ARCH_T5: diff --git a/src/llama-arch.h b/src/llama-arch.h index 48fe051a9..7159e23bf 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -115,6 +115,7 @@ enum llm_arch { LLM_ARCH_BAILINGMOE2, LLM_ARCH_BAILINGMOE3, LLM_ARCH_DOTS1, + LLM_ARCH_DOTS3NOTE, LLM_ARCH_ARCEE, LLM_ARCH_AFMOE, LLM_ARCH_LAGUNA, @@ -278,6 +279,9 @@ enum llm_kv { LLM_KV_ATTENTION_VALUE_LENGTH_MLA, LLM_KV_ATTENTION_KEY_LENGTH_SWA, LLM_KV_ATTENTION_VALUE_LENGTH_SWA, + LLM_KV_ATTENTION_KEY_LENGTH_MLA_SWA, + LLM_KV_ATTENTION_VALUE_LENGTH_MLA_SWA, + LLM_KV_ATTENTION_KV_LORA_RANK_SWA, LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, LLM_KV_ATTENTION_INDEXER_TOP_K, diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index 5212e19a2..8fca8e1bc 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -9,6 +9,7 @@ #include "llama-kv-cache.h" #include "llama-kv-cache-iswa.h" #include "llama-kv-cache-dsa.h" +#include "llama-kv-cache-dsa-iswa.h" #include "llama-kv-cache-msa.h" #include "llama-kv-cache-dsv4.h" #include "llama-memory-hybrid.h" @@ -507,10 +508,12 @@ void llm_graph_input_attn_k::set_input(const llama_ubatch * ubatch) { } bool llm_graph_input_attn_k::can_reuse(const llm_graph_params & params) { - const auto * mctx = static_cast(params.mctx); + mctx = static_cast(params.mctx); - this->mctx = mctx; + return can_reuse_impl(params); +} +bool llm_graph_input_attn_k::can_reuse_impl(const llm_graph_params & params) { bool res = true; res &= self_k_idxs->ne[0] == params.ubatch.n_tokens; @@ -567,10 +570,12 @@ void llm_graph_input_attn_k_dsa::set_input(const llama_ubatch * ubatch) { } bool llm_graph_input_attn_k_dsa::can_reuse(const llm_graph_params & params) { - const auto * mctx = static_cast(params.mctx); + mctx = static_cast(params.mctx); - this->mctx = mctx; + return can_reuse_impl(params); +} +bool llm_graph_input_attn_k_dsa::can_reuse_impl(const llm_graph_params & params) { bool res = true; res &= self_k_idxs_mla->ne[0] == params.ubatch.n_tokens; @@ -582,6 +587,25 @@ bool llm_graph_input_attn_k_dsa::can_reuse(const llm_graph_params & params) { return res; } +void llm_graph_input_attn_k_dsa_iswa::set_input(const llama_ubatch * ubatch) { + inp_dsa->set_input(ubatch); + inp_swa->set_input(ubatch); +} + +bool llm_graph_input_attn_k_dsa_iswa::can_reuse(const llm_graph_params & params) { + mctx = static_cast(params.mctx); + + inp_dsa->mctx = mctx->get_dsa(); + inp_swa->mctx = mctx->get_swa(); + + bool res = true; + + res &= inp_dsa->can_reuse_impl(params); + res &= inp_swa->can_reuse_impl(params); + + return res; +} + void llm_graph_input_attn_kv_iswa::set_input(const llama_ubatch * ubatch) { // base tensors may not be allocated if there are no non-SWA attention layers if (self_k_idxs && self_k_idxs->buffer) { @@ -3210,8 +3234,12 @@ ggml_tensor * llm_graph_context::build_attn( return cur; } -llm_graph_input_attn_k_dsa * llm_graph_context::build_attn_inp_k_dsa() const { - const auto * mctx_cur = static_cast(mctx); +static std::unique_ptr build_attn_inp_k_dsa_impl( + ggml_context * ctx0, + const llama_ubatch & ubatch, + const llama_hparams & hparams, + const llama_cparams & cparams, + const llama_kv_cache_dsa_context * mctx_cur) { auto inp = std::make_unique(hparams, cparams, mctx_cur); @@ -3235,9 +3263,35 @@ llm_graph_input_attn_k_dsa * llm_graph_context::build_attn_inp_k_dsa() const { inp->self_k_rot_lid = mctx_cur->get_lid()->build_input_k_rot(ctx0); } + return inp; +} + +llm_graph_input_attn_k_dsa * llm_graph_context::build_attn_inp_k_dsa() const { + const auto * mctx_cur = static_cast(mctx); + + auto inp = build_attn_inp_k_dsa_impl(ctx0, ubatch, hparams, cparams, mctx_cur); + return (llm_graph_input_attn_k_dsa *) res->add_input(std::move(inp)); } +llm_graph_input_attn_k_dsa_iswa * llm_graph_context::build_attn_inp_k_dsa_iswa() const { + const auto * mctx_cur = static_cast(mctx); + + auto inp_dsa = build_attn_inp_k_dsa_impl(ctx0, ubatch, hparams, cparams, mctx_cur->get_dsa()); + + // build_attn_inp_k_impl rejects SWA caches, so construct the input directly + auto inp_swa = std::make_unique(hparams, cparams, mctx_cur->get_swa()); + + inp_swa->self_k_idxs = mctx_cur->get_swa()->build_input_k_idxs(ctx0, ubatch); + + inp_swa->self_kq_mask = build_attn_inp_kq_mask(ctx0, mctx_cur->get_swa(), ubatch, cparams); + inp_swa->self_kq_mask_cnv = inp_swa->self_kq_mask; + + auto inp = std::make_unique(std::move(inp_dsa), std::move(inp_swa), mctx_cur); + + return (llm_graph_input_attn_k_dsa_iswa *) res->add_input(std::move(inp)); +} + llm_graph_input_attn_kv_msa * llm_graph_context::build_attn_inp_kv_msa(bool msa_enabled) const { const auto * mctx_cur = static_cast(mctx); diff --git a/src/llama-graph.h b/src/llama-graph.h index 94324c745..b388e028c 100644 --- a/src/llama-graph.h +++ b/src/llama-graph.h @@ -23,6 +23,7 @@ struct llama_memory_context_i; class llama_kv_cache_context; class llama_kv_cache_dsa_context; +class llama_kv_cache_dsa_iswa_context; class llama_kv_cache_msa_context; class llama_kv_cache_dsv4_raw_context; class llama_kv_cache_dsv4_context; @@ -374,6 +375,9 @@ public: bool can_reuse(const llm_graph_params & params) override; + // like can_reuse, but does not re-bind mctx + bool can_reuse_impl(const llm_graph_params & params); + ggml_tensor * get_k_idxs() const { return self_k_idxs; } ggml_tensor * get_kq_mask() const { return self_kq_mask_cnv; } @@ -405,6 +409,9 @@ public: bool can_reuse(const llm_graph_params & params) override; + // like can_reuse, but does not re-bind mctx + bool can_reuse_impl(const llm_graph_params & params); + ggml_tensor * get_k_idxs_mla() const { return self_k_idxs_mla; } ggml_tensor * get_k_idxs_lid() const { return self_k_idxs_lid; } @@ -427,6 +434,32 @@ public: const llama_kv_cache_dsa_context * mctx; }; +// DSA input (full-attention layers + indexer) with K-only input for the SWA layers +class llm_graph_input_attn_k_dsa_iswa : public llm_graph_input_i { +public: + llm_graph_input_attn_k_dsa_iswa( + std::unique_ptr inp_dsa, + std::unique_ptr inp_swa, + const llama_kv_cache_dsa_iswa_context * mctx) : + inp_dsa(std::move(inp_dsa)), + inp_swa(std::move(inp_swa)), + mctx(mctx) { + } + ~llm_graph_input_attn_k_dsa_iswa() = default; + + void set_input(const llama_ubatch * ubatch) override; + + bool can_reuse(const llm_graph_params & params) override; + + llm_graph_input_attn_k_dsa * get_dsa() const { return inp_dsa.get(); } + llm_graph_input_attn_k * get_swa() const { return inp_swa.get(); } + + std::unique_ptr inp_dsa; + std::unique_ptr inp_swa; + + const llama_kv_cache_dsa_iswa_context * mctx; +}; + // standard K/V attention input against the base cache, plus destination indices for the indexer key cache class llm_graph_input_attn_kv_msa : public llm_graph_input_attn_kv { public: @@ -1191,6 +1224,8 @@ struct llm_graph_context { llm_graph_input_attn_k_dsa * build_attn_inp_k_dsa() const; + llm_graph_input_attn_k_dsa_iswa * build_attn_inp_k_dsa_iswa() const; + llm_graph_input_attn_kv_msa * build_attn_inp_kv_msa(bool msa_enabled) const; ggml_tensor * build_attn( diff --git a/src/llama-hparams.h b/src/llama-hparams.h index f6af36436..c3c14292c 100644 --- a/src/llama-hparams.h +++ b/src/llama-hparams.h @@ -101,6 +101,11 @@ struct llama_hparams { uint32_t n_group_used = 0; uint32_t n_group_experts = 0; + // MLA + SWA (i.e. dots3note) + uint32_t n_lora_kv_swa = 0; + uint32_t n_embd_head_k_mla_swa = 0; + uint32_t n_embd_head_v_mla_swa = 0; + float expert_group_scale = 0.05f; float expert_weights_scale = 0.0f; bool expert_weights_norm = false; diff --git a/src/llama-kv-cache-dsa-iswa.cpp b/src/llama-kv-cache-dsa-iswa.cpp new file mode 100644 index 000000000..dc10342a1 --- /dev/null +++ b/src/llama-kv-cache-dsa-iswa.cpp @@ -0,0 +1,341 @@ +#include "llama-kv-cache-dsa-iswa.h" + +#include "llama-impl.h" +#include "llama-batch.h" +#include "llama-model.h" + +#include +#include + +// +// llama_kv_cache_dsa_iswa +// + +llama_kv_cache_dsa_iswa::llama_kv_cache_dsa_iswa( + const llama_model & model, + ggml_type type_k, + ggml_type type_v, + bool v_trans, + bool offload, + bool swa_full, + bool unified, + uint32_t kv_size, + uint32_t n_seq_max, + uint32_t n_ubatch, + uint32_t n_pad, + const layer_filter_cb & filter_mla, + const layer_filter_cb & filter_lid, + const layer_reuse_cb & reuse) : unified(unified) { + + const auto & hparams = model.hparams; + + // chain filters + const layer_filter_cb filter_dsa = [&](int32_t il) { + if (filter_mla && !filter_mla(il)) { + return false; + } + + return !hparams.is_swa(il); + }; + + const layer_filter_cb filter_swa = [&](int32_t il) { + if (filter_mla && !filter_mla(il)) { + return false; + } + + return hparams.is_swa(il); + }; + + const uint32_t size_dsa = kv_size; + + // note: the SWA cache is always padded to 256 for performance + // https://github.com/ggml-org/llama.cpp/issues/17037 + uint32_t size_swa = GGML_PAD(std::min(size_dsa, hparams.n_swa*(unified ? n_seq_max : 1) + n_ubatch), 256); + + // when using full-size SWA cache, we set the SWA cache size to be equal to the base cache size + if (swa_full) { + LLAMA_LOG_WARN("%s: using full-size SWA cache (ref: %s)\n", + __func__, "https://github.com/ggml-org/llama.cpp/pull/13194#issuecomment-2868343055"); + + size_swa = size_dsa; + } + + LLAMA_LOG_INFO("%s: creating DSA KV cache, size = %u cells\n", __func__, size_dsa); + + kv_dsa = std::make_unique( + model, type_k, type_v, + v_trans, offload, unified, size_dsa, n_seq_max, n_pad, + 0, LLAMA_SWA_TYPE_NONE, filter_dsa, filter_lid, reuse); + + LLAMA_LOG_INFO("%s: creating SWA KV cache, size = %u cells\n", __func__, size_swa); + + kv_swa = std::make_unique( + model, hparams, type_k, type_v, + v_trans, offload, unified, size_swa, n_seq_max, n_pad, + hparams.n_swa, hparams.swa_type, nullptr, filter_swa, reuse, nullptr); +} + +void llama_kv_cache_dsa_iswa::clear(bool data) { + kv_dsa->clear(data); + kv_swa->clear(data); +} + +bool llama_kv_cache_dsa_iswa::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1) { + bool res = true; + + res = res & kv_dsa->seq_rm(seq_id, p0, p1); + res = res & kv_swa->seq_rm(seq_id, p0, p1); + + return res; +} + +void llama_kv_cache_dsa_iswa::seq_cp(llama_seq_id seq_id_src, llama_seq_id seq_id_dst, llama_pos p0, llama_pos p1) { + kv_dsa->seq_cp(seq_id_src, seq_id_dst, p0, p1); + kv_swa->seq_cp(seq_id_src, seq_id_dst, p0, p1); +} + +void llama_kv_cache_dsa_iswa::seq_keep(llama_seq_id seq_id) { + kv_dsa->seq_keep(seq_id); + kv_swa->seq_keep(seq_id); +} + +void llama_kv_cache_dsa_iswa::seq_add(llama_seq_id seq_id, llama_pos p0, llama_pos p1, llama_pos shift) { + kv_dsa->seq_add(seq_id, p0, p1, shift); + kv_swa->seq_add(seq_id, p0, p1, shift); +} + +void llama_kv_cache_dsa_iswa::seq_div(llama_seq_id seq_id, llama_pos p0, llama_pos p1, int d) { + kv_dsa->seq_div(seq_id, p0, p1, d); + kv_swa->seq_div(seq_id, p0, p1, d); +} + +llama_pos llama_kv_cache_dsa_iswa::seq_pos_min(llama_seq_id seq_id) const { + // the DSA cache is a superset of the SWA cache, so we can just check the SWA cache + return kv_swa->seq_pos_min(seq_id); +} + +llama_pos llama_kv_cache_dsa_iswa::seq_pos_max(llama_seq_id seq_id) const { + return kv_swa->seq_pos_max(seq_id); +} + +std::map llama_kv_cache_dsa_iswa::memory_breakdown() const { + std::map mb = kv_dsa->memory_breakdown(); + for (const auto & buft_size : kv_swa->memory_breakdown()) { + mb[buft_size.first] += buft_size.second; + } + return mb; +} + +llama_memory_context_ptr llama_kv_cache_dsa_iswa::init_batch(llama_batch_allocr & balloc, uint32_t n_ubatch, bool embd_all) { + GGML_UNUSED(embd_all); + + // first try simple split + do { + if (!unified) { + // requires equal splits, so we skip the simple split + break; + } + + balloc.split_reset(); + + std::vector ubatches; + while (true) { + auto ubatch = balloc.split_simple(n_ubatch); + + if (ubatch.n_tokens == 0) { + break; + } + + ubatches.push_back(std::move(ubatch)); // NOLINT + } + + if (balloc.get_n_used() < balloc.get_n_tokens()) { + // failed to find a suitable split + break; + } + + auto sinfos_mla = kv_dsa->get_mla()->prepare(ubatches); + if (sinfos_mla.empty()) { + break; + } + + auto sinfos_lid = kv_dsa->get_lid()->prepare(ubatches); + if (sinfos_lid.empty()) { + break; + } + + auto sinfos_swa = kv_swa->prepare(ubatches); + if (sinfos_swa.empty()) { + break; + } + + assert(sinfos_mla.size() == sinfos_swa.size()); + + return std::make_unique( + this, std::move(sinfos_mla), std::move(sinfos_lid), std::move(sinfos_swa), std::move(ubatches)); + } while (false); + + // if it fails, try equal split + do { + balloc.split_reset(); + + std::vector ubatches; + while (true) { + auto ubatch = balloc.split_equal(n_ubatch, !unified, 0); + + if (ubatch.n_tokens == 0) { + break; + } + + ubatches.push_back(std::move(ubatch)); // NOLINT + } + + if (balloc.get_n_used() < balloc.get_n_tokens()) { + // failed to find a suitable split + break; + } + + auto sinfos_mla = kv_dsa->get_mla()->prepare(ubatches); + if (sinfos_mla.empty()) { + break; + } + + auto sinfos_lid = kv_dsa->get_lid()->prepare(ubatches); + if (sinfos_lid.empty()) { + break; + } + + auto sinfos_swa = kv_swa->prepare(ubatches); + if (sinfos_swa.empty()) { + break; + } + + assert(sinfos_mla.size() == sinfos_swa.size()); + + return std::make_unique( + this, std::move(sinfos_mla), std::move(sinfos_lid), std::move(sinfos_swa), std::move(ubatches)); + } while (false); + + return std::make_unique(LLAMA_MEMORY_STATUS_FAILED_PREPARE); +} + +llama_memory_context_ptr llama_kv_cache_dsa_iswa::init_full() { + return std::make_unique(this); +} + +llama_memory_context_ptr llama_kv_cache_dsa_iswa::init_update(llama_context * lctx, bool optimize) { + return std::make_unique(this, lctx, optimize); +} + +bool llama_kv_cache_dsa_iswa::get_can_shift() const { + return kv_dsa->get_can_shift() && + kv_swa->get_can_shift() && + kv_dsa->get_mla()->get_size() == kv_swa->get_size(); +} + +void llama_kv_cache_dsa_iswa::state_write(llama_io_write_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) const { + if ((flags & LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY) == 0) { + kv_dsa->state_write(io, seq_id, flags); + } + + kv_swa->state_write(io, seq_id, flags); +} + +void llama_kv_cache_dsa_iswa::state_read(llama_io_read_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) { + if ((flags & LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY) == 0) { + kv_dsa->state_read(io, seq_id, flags); + } + + kv_swa->state_read(io, seq_id, flags); +} + +llama_kv_cache_dsa * llama_kv_cache_dsa_iswa::get_dsa() const { + return kv_dsa.get(); +} + +llama_kv_cache * llama_kv_cache_dsa_iswa::get_swa() const { + return kv_swa.get(); +} + +// +// llama_kv_cache_dsa_iswa_context +// + +llama_kv_cache_dsa_iswa_context::llama_kv_cache_dsa_iswa_context(llama_memory_status status) : status(status) {} + +llama_kv_cache_dsa_iswa_context::llama_kv_cache_dsa_iswa_context( + llama_kv_cache_dsa_iswa * kv) : + ctx_dsa(kv->get_dsa()->init_full()), + ctx_swa(kv->get_swa()->init_full()), + status(llama_memory_status_combine(ctx_dsa->get_status(), ctx_swa->get_status())) { +} + +llama_kv_cache_dsa_iswa_context::llama_kv_cache_dsa_iswa_context( + llama_kv_cache_dsa_iswa * kv, + llama_context * lctx, + bool optimize) : + ctx_dsa(kv->get_dsa()->init_update(lctx, optimize)), + ctx_swa(kv->get_swa()->init_update(lctx, optimize)), + status(llama_memory_status_combine(ctx_dsa->get_status(), ctx_swa->get_status())) { +} + +llama_kv_cache_dsa_iswa_context::llama_kv_cache_dsa_iswa_context( + llama_kv_cache_dsa_iswa * kv, + slot_info_vec_t sinfos_mla, + slot_info_vec_t sinfos_lid, + slot_info_vec_t sinfos_swa, + std::vector ubatches) : + ubatches(std::move(ubatches)), + // note: here we copy the ubatches. not sure if this is ideal + ctx_dsa(new llama_kv_cache_dsa_context(kv->get_dsa(), std::move(sinfos_mla), std::move(sinfos_lid), this->ubatches)), + ctx_swa(new llama_kv_cache_context(kv->get_swa(), std::move(sinfos_swa), this->ubatches)), + status(llama_memory_status_combine(ctx_dsa->get_status(), ctx_swa->get_status())) { +} + +llama_kv_cache_dsa_iswa_context:: ~llama_kv_cache_dsa_iswa_context() = default; + +bool llama_kv_cache_dsa_iswa_context::next() { + assert(status == LLAMA_MEMORY_STATUS_SUCCESS); + + ctx_dsa->next(); + ctx_swa->next(); + + if (++i_next >= ubatches.size()) { + return false; + } + + return true; +} + +bool llama_kv_cache_dsa_iswa_context::apply() { + assert(!llama_memory_status_is_fail(status)); + + bool res = true; + + res = res & ctx_dsa->apply(); + res = res & ctx_swa->apply(); + + return res; +} + +llama_memory_status llama_kv_cache_dsa_iswa_context::get_status() const { + return status; +} + +const llama_ubatch & llama_kv_cache_dsa_iswa_context::get_ubatch() const { + assert(status == LLAMA_MEMORY_STATUS_SUCCESS); + + return ubatches[i_next]; +} + +const llama_kv_cache_dsa_context * llama_kv_cache_dsa_iswa_context::get_dsa() const { + assert(status == LLAMA_MEMORY_STATUS_SUCCESS); + + return static_cast(ctx_dsa.get()); +} + +const llama_kv_cache_context * llama_kv_cache_dsa_iswa_context::get_swa() const { + assert(status == LLAMA_MEMORY_STATUS_SUCCESS); + + return static_cast(ctx_swa.get()); +} diff --git a/src/llama-kv-cache-dsa-iswa.h b/src/llama-kv-cache-dsa-iswa.h new file mode 100644 index 000000000..28cf95bf0 --- /dev/null +++ b/src/llama-kv-cache-dsa-iswa.h @@ -0,0 +1,134 @@ +#pragma once + +#include "llama-kv-cache-dsa.h" + +#include + +// +// llama_kv_cache_dsa_iswa +// + +// utilizes two child memories: llama_kv_cache_dsa for the full-attention (DSA) layers and llama_kv_cache for the SWA layers + +class llama_kv_cache_dsa_iswa : public llama_memory_i { +public: + llama_kv_cache_dsa_iswa( + const llama_model & model, + ggml_type type_k, + ggml_type type_v, + bool v_trans, + bool offload, + bool swa_full, + bool unified, + uint32_t kv_size, + uint32_t n_seq_max, + uint32_t n_ubatch, + uint32_t n_pad, + const layer_filter_cb & filter_mla, + const layer_filter_cb & filter_lid, + const layer_reuse_cb & reuse); + + ~llama_kv_cache_dsa_iswa() = default; + + // + // llama_memory_i + // + + llama_memory_context_ptr init_batch( + llama_batch_allocr & balloc, + uint32_t n_ubatch, + bool embd_all) override; + + llama_memory_context_ptr init_full() override; + + llama_memory_context_ptr init_update(llama_context * lctx, bool optimize) override; + + bool get_can_shift() const override; + + void clear(bool data) override; + + bool seq_rm (llama_seq_id seq_id, llama_pos p0, llama_pos p1) override; + void seq_cp (llama_seq_id seq_id_src, llama_seq_id seq_id_dst, llama_pos p0, llama_pos p1) override; + void seq_keep(llama_seq_id seq_id) override; + void seq_add (llama_seq_id seq_id, llama_pos p0, llama_pos p1, llama_pos shift) override; + void seq_div (llama_seq_id seq_id, llama_pos p0, llama_pos p1, int d) override; + + llama_pos seq_pos_min(llama_seq_id seq_id) const override; + llama_pos seq_pos_max(llama_seq_id seq_id) const override; + + std::map memory_breakdown() const override; + + // state write/load + + void state_write(llama_io_write_i & io, llama_seq_id seq_id = -1, llama_state_seq_flags flags = 0) const override; + void state_read (llama_io_read_i & io, llama_seq_id seq_id = -1, llama_state_seq_flags flags = 0) override; + + // + // llama_kv_cache_dsa_iswa specific API + // + + llama_kv_cache_dsa * get_dsa() const; + llama_kv_cache * get_swa() const; + +private: + const bool unified; + + std::unique_ptr kv_dsa; + std::unique_ptr kv_swa; +}; + +class llama_kv_cache_dsa_iswa_context : public llama_memory_context_i { +public: + using slot_info_vec_t = llama_kv_cache::slot_info_vec_t; + + // used for errors + llama_kv_cache_dsa_iswa_context(llama_memory_status status); + + // used to create a full-cache context + llama_kv_cache_dsa_iswa_context( + llama_kv_cache_dsa_iswa * kv); + + // used to create an update context + llama_kv_cache_dsa_iswa_context( + llama_kv_cache_dsa_iswa * kv, + llama_context * lctx, + bool optimize); + + // used to create a batch processing context from a batch + llama_kv_cache_dsa_iswa_context( + llama_kv_cache_dsa_iswa * kv, + slot_info_vec_t sinfos_mla, + slot_info_vec_t sinfos_lid, + slot_info_vec_t sinfos_swa, + std::vector ubatches); + + virtual ~llama_kv_cache_dsa_iswa_context(); + + // + // llama_memory_context_i + // + + bool next() override; + bool apply() override; + + llama_memory_status get_status() const override; + const llama_ubatch & get_ubatch() const override; + + // + // llama_kv_cache_dsa_iswa_context specific API + // + + const llama_kv_cache_dsa_context * get_dsa() const; + const llama_kv_cache_context * get_swa() const; + +private: + // the index of the next ubatch to process + size_t i_next = 0; + + std::vector ubatches; + + const llama_memory_context_ptr ctx_dsa; + const llama_memory_context_ptr ctx_swa; + + const llama_memory_status status; +}; diff --git a/src/llama-kv-cache.cpp b/src/llama-kv-cache.cpp index 5382cd726..2e2bd7dc6 100644 --- a/src/llama-kv-cache.cpp +++ b/src/llama-kv-cache.cpp @@ -323,7 +323,8 @@ llama_kv_cache::llama_kv_cache( hparams.n_embd_head_k() % 64 == 0; // always create Hadamard rotation tensors for DeepSeek lightning indexers - if ((model.arch == LLM_ARCH_DEEPSEEK32 || model.arch == LLM_ARCH_DEEPSEEK4 || model.arch == LLM_ARCH_GLM_DSA) && + if ((model.arch == LLM_ARCH_DEEPSEEK32 || model.arch == LLM_ARCH_DEEPSEEK4 || + model.arch == LLM_ARCH_GLM_DSA || model.arch == LLM_ARCH_DOTS3NOTE) && hparams.n_embd_head_k_full == hparams.indexer_head_size) { attn_rot_k = true; } diff --git a/src/llama-model-saver.cpp b/src/llama-model-saver.cpp index b9e0a6009..0d39e6de8 100644 --- a/src/llama-model-saver.cpp +++ b/src/llama-model-saver.cpp @@ -31,6 +31,7 @@ bool llama_model_saver_supports_arch(llm_arch arch) { case LLM_ARCH_MELLUM: case LLM_ARCH_LAGUNA: case LLM_ARCH_GRANITE_SWA: + case LLM_ARCH_DOTS3NOTE: // TODO: need to handle SWA pattern and MLA+SWA config return false; default: return true; diff --git a/src/llama-model.cpp b/src/llama-model.cpp index d7874e0a9..de0d3c1a6 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -11,6 +11,7 @@ #include "llama-kv-cache.h" #include "llama-kv-cache-iswa.h" #include "llama-kv-cache-dsa.h" +#include "llama-kv-cache-dsa-iswa.h" #include "llama-kv-cache-msa.h" #include "llama-kv-cache-dsv4.h" #include "llama-memory-hybrid.h" @@ -194,6 +195,8 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params return new llama_model_deepseek2ocr(params); case LLM_ARCH_DEEPSEEK32: return new llama_model_deepseek32(params); + case LLM_ARCH_DOTS3NOTE: + return new llama_model_dots3note(params); case LLM_ARCH_DEEPSEEK4: return new llama_model_deepseek4(params); case LLM_ARCH_GLM_DSA: @@ -851,6 +854,7 @@ const char * llm_type_name(llm_type type) { case LLM_TYPE_230B_A10B: return "230B.A10B"; case LLM_TYPE_428B_A23B: return "428B.A23B"; case LLM_TYPE_235B_A22B: return "235B.A22B"; + case LLM_TYPE_288B_A19B: return "288B.A19B"; case LLM_TYPE_300B_A47B: return "300B.A47B"; case LLM_TYPE_310B_A15B: return "310B.A15B"; case LLM_TYPE_355B_A32B: return "355B.A32B"; @@ -1924,7 +1928,9 @@ void llama_model::print_info() const { LLAMA_LOG_INFO("%s: expert_weights_scale = %.1f\n", __func__, hparams.expert_weights_scale); } - if (arch == LLM_ARCH_DEEPSEEK2 || arch == LLM_ARCH_DEEPSEEK2OCR || arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_MISTRAL4) { + if (arch == LLM_ARCH_DEEPSEEK2 || arch == LLM_ARCH_DEEPSEEK2OCR || + arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || + arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_MISTRAL4) { LLAMA_LOG_INFO("%s: n_layer_dense_lead = %d\n", __func__, hparams.n_layer_dense_lead); LLAMA_LOG_INFO("%s: n_lora_q = %d\n", __func__, hparams.n_lora_q); LLAMA_LOG_INFO("%s: n_lora_kv = %d\n", __func__, hparams.n_lora_kv); @@ -2193,6 +2199,57 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, nullptr); } } break; + case LLM_ARCH_DOTS3NOTE: + { + GGML_ASSERT(hparams.swa_type != LLAMA_SWA_TYPE_NONE); + + if (params.ctx_type == LLAMA_CONTEXT_TYPE_MTP && hparams.n_layer_nextn > 0) { + // MTP draft context: plain attention KV cache holding only the nextn layer + llama_kv_cache::layer_filter_cb filter = + [&](uint32_t il) { return il >= hparams.n_layer(); }; + + res = new llama_kv_cache( + *this, + hparams, + params.type_k, + params.type_v, + !cparams.flash_attn, + cparams.offload_kqv, + cparams.kv_unified, + cparams.n_ctx_seq, + cparams.n_seq_max, + 1, + hparams.n_swa, + hparams.swa_type, + nullptr, + filter, + nullptr, + nullptr); + } else { + // main context: DSA cache for the trunk full-attention layers plus a window-sized SWA cache + llama_kv_cache::layer_filter_cb filter_mla = nullptr; + if (hparams.n_layer_nextn > 0) { + filter_mla = [&](uint32_t il) { return il < hparams.n_layer(); }; + } + llama_kv_cache::layer_filter_cb filter_lid = [&](uint32_t il) { return il < hparams.n_layer() && hparams.is_indexer_full(il); }; + + res = new llama_kv_cache_dsa_iswa( + *this, + params.type_k, + params.type_v, + !cparams.flash_attn, + cparams.offload_kqv, + params.swa_full, + cparams.kv_unified, + cparams.n_ctx_seq, + cparams.n_seq_max, + cparams.n_ubatch, + 1, + filter_mla, + filter_lid, + nullptr); + } + } break; case LLM_ARCH_DEEPSEEK4: { GGML_ASSERT(hparams.swa_type != LLAMA_SWA_TYPE_NONE); @@ -2661,6 +2718,7 @@ llama_rope_type llama_model_rope_type(const llama_model * model) { case LLM_ARCH_LLAMA_EMBED: case LLM_ARCH_MAINCODER: case LLM_ARCH_GLM_DSA: + case LLM_ARCH_DOTS3NOTE: case LLM_ARCH_NANBEIGE: case LLM_ARCH_POCKETTTS: return LLAMA_ROPE_TYPE_NORM; diff --git a/src/llama-model.h b/src/llama-model.h index 4412ef08e..44bd96757 100644 --- a/src/llama-model.h +++ b/src/llama-model.h @@ -140,6 +140,7 @@ enum llm_type { LLM_TYPE_230B_A10B, // Minimax M2 LLM_TYPE_428B_A23B, // Minimax M3 LLM_TYPE_235B_A22B, + LLM_TYPE_288B_A19B, // dots3-note LLM_TYPE_300B_A47B, // Ernie MoE big LLM_TYPE_310B_A15B, // /MiMo-V2-Flash LLM_TYPE_355B_A32B, // GLM-4.5 diff --git a/src/models/dots3note.cpp b/src/models/dots3note.cpp new file mode 100644 index 000000000..00a008c2c --- /dev/null +++ b/src/models/dots3note.cpp @@ -0,0 +1,480 @@ +#include "models.h" + +#include "llama-kv-cache.h" +#include "llama-kv-cache-dsa.h" + +// note: code adapted from deepseek32.cpp (DSA indexer + absorbed MLA) and step35.cpp (head-wise output gate) + +void llama_model_dots3note::load_arch_hparams(llama_model_loader & ml) { + ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps); + hparams.f_norm_eps = 1e-6; // eps for the indexer k_norm layer norm + + // TODO: use MTP layer + ml.get_key(LLM_KV_NEXTN_PREDICT_LAYERS, hparams.n_layer_nextn, false); + GGML_ASSERT(hparams.n_layer_nextn < hparams.n_layer_all && "n_layer_nextn must be < n_layer_all"); + + // MoE parameters + ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared); + ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp); + ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead); + ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false); + ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false); + ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func); + + // MLA parameters of the full-attention layers + ml.get_key(LLM_KV_ATTENTION_Q_LORA_RANK, hparams.n_lora_q); + ml.get_key(LLM_KV_ATTENTION_KV_LORA_RANK, hparams.n_lora_kv); + ml.get_key(LLM_KV_ATTENTION_KEY_LENGTH_MLA, hparams.n_embd_head_k_mla_impl); + ml.get_key(LLM_KV_ATTENTION_VALUE_LENGTH_MLA, hparams.n_embd_head_v_mla_impl); + + // MLA parameters of the sliding-window layers + ml.get_key(LLM_KV_ATTENTION_KV_LORA_RANK_SWA, hparams.n_lora_kv_swa); + ml.get_key(LLM_KV_ATTENTION_KEY_LENGTH_MLA_SWA, hparams.n_embd_head_k_mla_swa); + ml.get_key(LLM_KV_ATTENTION_VALUE_LENGTH_MLA_SWA, hparams.n_embd_head_v_mla_swa); + + hparams.swa_type = LLAMA_SWA_TYPE_STANDARD; + ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa); + ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa); + ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl); + + // DSA parameters + ml.get_key(LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, hparams.indexer_n_head); + ml.get_key(LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, hparams.indexer_head_size); + ml.get_key(LLM_KV_ATTENTION_INDEXER_TOP_K, hparams.indexer_top_k); + ml.get_arr(LLM_KV_ATTENTION_INDEXER_TYPES, hparams.is_indexer_full_impl); + + switch (hparams.n_layer()) { + case 46: type = LLM_TYPE_288B_A19B; break; + default: type = LLM_TYPE_UNKNOWN; + } +} + +void llama_model_dots3note::load_arch_tensors(llama_model_loader & ml) { + LLAMA_LOAD_LOCALS; + GGML_UNUSED(ml); + + if (!hparams.is_mla()) { + throw std::runtime_error("DOTS3NOTE architecture requires MLA"); + } + + const int64_t n_embd_head_qk_rope = hparams.n_rot(); + + const int64_t q_lora_rank = hparams.n_lora_q; + const int64_t n_ff_exp = hparams.n_ff_exp; + const int64_t n_expert_shared = hparams.n_expert_shared; + + tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0); + + output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0); + output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED); + if (!output) { + output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED); + } + + for (int i = 0; i < n_layer_all; ++i) { + auto & layer = layers[i]; + + const bool is_mtp = i >= n_layer; + // the NextN/MTP block uses the sliding-attention geometry + const bool is_swa = is_mtp || hparams.is_swa(i); + + // MTP tensors are preserved in the GGUF but there is no MTP graph yet + const int flags = is_mtp ? TENSOR_SKIP | TENSOR_NOT_REQUIRED : 0; + + const int64_t n_head_l = hparams.n_head(i); + + const int64_t kv_lora_rank = is_swa ? hparams.n_lora_kv_swa : hparams.n_lora_kv; + const int64_t n_embd_head_k_mla = is_swa ? hparams.n_embd_head_k_mla_swa : hparams.n_embd_head_k_mla(); + const int64_t n_embd_head_v_mla = is_swa ? hparams.n_embd_head_v_mla_swa : hparams.n_embd_head_v_mla(); + const int64_t n_embd_head_qk_nope = n_embd_head_k_mla - n_embd_head_qk_rope; + + layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, flags); + layer.attn_q_a_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_A_NORM, "weight", i), {q_lora_rank}, flags); + layer.attn_kv_a_norm = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_NORM, "weight", i), {kv_lora_rank}, flags); + // norm applied on the shared rope key before rope + layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {n_embd_head_qk_rope}, flags); + + layer.wq_a = create_tensor(tn(LLM_TENSOR_ATTN_Q_A, "weight", i), {n_embd, q_lora_rank}, flags); + layer.wq_b = create_tensor(tn(LLM_TENSOR_ATTN_Q_B, "weight", i), {q_lora_rank, n_head_l * n_embd_head_k_mla}, flags); + + layer.wkv_a_mqa = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_MQA, "weight", i), {n_embd, kv_lora_rank + n_embd_head_qk_rope}, flags); + + layer.wk_b = create_tensor(tn(LLM_TENSOR_ATTN_K_B, "weight", i), {n_embd_head_qk_nope, kv_lora_rank, n_head_l}, flags); + layer.wv_b = create_tensor(tn(LLM_TENSOR_ATTN_V_B, "weight", i), {kv_lora_rank, n_embd_head_v_mla, n_head_l}, flags); + + layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_head_l * n_embd_head_v_mla, n_embd}, flags); + + // head-wise sigmoid output gate + layer.wqkv_gate = create_tensor(tn(LLM_TENSOR_ATTN_GATE, "weight", i), {n_embd, n_head_l}, flags); + + layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, flags); + + // DSA indexer + if (!is_mtp && hparams.is_indexer_full(i)) { + layer.indexer_k_norm = create_tensor(tn(LLM_TENSOR_INDEXER_K_NORM, "weight", i), {hparams.indexer_head_size}, flags); + layer.indexer_k_norm_b = create_tensor(tn(LLM_TENSOR_INDEXER_K_NORM, "bias", i), {hparams.indexer_head_size}, flags); + layer.indexer_proj = create_tensor(tn(LLM_TENSOR_INDEXER_PROJ, "weight", i), {n_embd, hparams.indexer_n_head}, flags); + layer.indexer_attn_k = create_tensor(tn(LLM_TENSOR_INDEXER_ATTN_K, "weight", i), {n_embd, hparams.indexer_head_size}, flags); + layer.indexer_attn_q_b = create_tensor(tn(LLM_TENSOR_INDEXER_ATTN_Q_B, "weight", i), {q_lora_rank, hparams.indexer_n_head * hparams.indexer_head_size}, flags); + } + + if (is_mtp || i < (int) hparams.n_layer_dense_lead) { + layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, flags); + layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd}, flags); + layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, flags); + } else { + if (n_expert == 0 || n_expert_used == 0) { + throw std::runtime_error("n_expert and n_expert_used must be > 0"); + } + + layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, flags); + layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, flags); + + layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, flags); + layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, flags); + layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, flags); + + layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", i), {n_embd, n_ff_exp * n_expert_shared}, flags); + layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", i), { n_ff_exp * n_expert_shared, n_embd}, flags); + layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", i), {n_embd, n_ff_exp * n_expert_shared}, flags); + } + + if (is_mtp) { + layer.nextn.eh_proj = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", i), { 2 * n_embd, n_embd }, flags); + layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), { n_embd }, flags); + layer.nextn.hnorm = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM, "weight", i), { n_embd }, flags); + layer.nextn.embed_tokens = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS, "weight", i), { n_embd, n_vocab }, flags); + layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", i), { n_embd }, flags); + } + } +} + +std::unique_ptr llama_model_dots3note::build_arch_graph(const llm_graph_params & params) const { + return std::make_unique(*this, params); +} + +llama_model_dots3note::graph::graph(const llama_model & model, const llm_graph_params & params) : + llm_graph_context(params) { + GGML_ASSERT(hparams.is_mla()); + + const int64_t n_embd_head_qk_rope = hparams.n_rot(); + + const int64_t n_indexer_head = hparams.indexer_n_head; + const int64_t n_embd_indexer_head = hparams.indexer_head_size; + const uint32_t n_indexer_top_k = hparams.indexer_top_k; + + // the indexer head layout is [rope | nope] + GGML_ASSERT(hparams.n_rot() <= n_embd_indexer_head); + + ggml_tensor * cur; + ggml_tensor * inpL; + + inpL = build_inp_embd(model.tok_embd); + + ggml_tensor * inp_pos = build_inp_pos(); + + llm_graph_input_attn_k_dsa_iswa * inp_attn = build_attn_inp_k_dsa_iswa(); + + ggml_tensor * inp_out_ids = build_inp_out_ids(); + + for (int il = 0; il < n_layer; ++il) { + ggml_tensor * inpSA = inpL; + + const bool is_swa = hparams.is_swa(il); + + const int64_t n_head_l = hparams.n_head(il); + + const int64_t kv_lora_rank = is_swa ? hparams.n_lora_kv_swa : hparams.n_lora_kv; + const int64_t n_embd_head_k_mla = is_swa ? hparams.n_embd_head_k_mla_swa : hparams.n_embd_head_k_mla(); + const int64_t n_embd_head_v_mla = is_swa ? hparams.n_embd_head_v_mla_swa : hparams.n_embd_head_v_mla(); + const int64_t n_embd_head_qk_nope = n_embd_head_k_mla - n_embd_head_qk_rope; + + const float kq_scale = 1.0f/sqrtf(float(n_embd_head_k_mla)); + const float freq_base_l = model.get_rope_freq_base(cparams, il); + + // norm + cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "attn_norm", il); + + // self_attention + { + ggml_tensor * attn_inp = cur; + + ggml_tensor * qr = ggml_mul_mat(ctx0, model.layers[il].wq_a, cur); + cb(qr, "qr", il); + + qr = build_norm(qr, model.layers[il].attn_q_a_norm, nullptr, LLM_NORM_RMS, il); + cb(qr, "qr", il); + + ggml_tensor * top_k = nullptr; + + // lightning indexer (full-attention layers only) + if (!is_swa) { + ggml_tensor * indexer_q = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_q_b, qr); + cb(indexer_q, "indexer_q", il); + + // {n_embd_indexer_head, n_indexer_head, n_tokens} + indexer_q = ggml_reshape_3d(ctx0, indexer_q, n_embd_indexer_head, n_indexer_head, n_tokens); + indexer_q = ggml_rope_ext(ctx0, indexer_q, inp_pos, nullptr, n_rot, + LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + cb(indexer_q, "indexer_q", il); + + ggml_tensor * indexer_k = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_k, cur); + cb(indexer_k, "indexer_k", il); + + indexer_k = build_norm(indexer_k, model.layers[il].indexer_k_norm, model.layers[il].indexer_k_norm_b, LLM_NORM, il); + cb(indexer_k, "indexer_k", il); + + // {n_embd_indexer_head, 1, n_tokens} + indexer_k = ggml_reshape_3d(ctx0, indexer_k, n_embd_indexer_head, 1, n_tokens); + indexer_k = ggml_rope_ext(ctx0, indexer_k, inp_pos, nullptr, n_rot, + LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + cb(indexer_k, "indexer_k", il); + + // perform Hadamard transform on indexer q and k + indexer_q = ggml_mul_mat(ctx0, inp_attn->get_dsa()->self_k_rot_lid, indexer_q); + cb(indexer_q, "indexer_q", il); + indexer_k = ggml_mul_mat(ctx0, inp_attn->get_dsa()->self_k_rot_lid, indexer_k); + cb(indexer_k, "indexer_k", il); + + // store indexer keys to KV cache + const auto * mctx_lid = inp_attn->get_dsa()->mctx->get_lid(); + const auto & k_idxs_lid = inp_attn->get_dsa()->get_k_idxs_lid(); + ggml_build_forward_expand(gf, mctx_lid->cpy_k(ctx0, indexer_k, k_idxs_lid, il)); + + ggml_tensor * indexer_weights = ggml_mul_mat(ctx0, model.layers[il].indexer_proj, cur); + cb(indexer_weights, "indexer_weights", il); + + indexer_k = mctx_lid->get_k(ctx0, il); + + // split the batch into streams if needed + const auto n_stream = indexer_k->ne[3]; + indexer_q = ggml_view_4d(ctx0, indexer_q, indexer_q->ne[0], indexer_q->ne[1], indexer_q->ne[2]/n_stream, n_stream, indexer_q->nb[1], indexer_q->nb[2], indexer_q->nb[3]/n_stream, 0); + indexer_weights = ggml_view_4d(ctx0, indexer_weights, indexer_weights->ne[0], indexer_weights->ne[1]/n_stream, indexer_weights->ne[2], n_stream, indexer_weights->nb[1], indexer_weights->nb[2]/n_stream, indexer_weights->nb[3]/n_stream, 0); + + // pre-scale weights to avoid scaling operations on huge indexer_score tensor + indexer_weights = ggml_scale(ctx0, indexer_weights, 1.0f / sqrtf(float(n_embd_indexer_head * n_indexer_head))); + cb(indexer_weights, "indexer_weights", il); + + ggml_tensor * indexer_score = nullptr; + if (cparams.fused_lid) { + indexer_score = ggml_lightning_indexer(ctx0, indexer_q, indexer_k, indexer_weights, inp_attn->get_dsa()->get_kq_mask_lid()); + cb(indexer_score, "indexer_score", il); + res->add_fused_node({LLM_FUSED_OP_LIGHTNING_INDEXER, indexer_score, il}); + } else { + indexer_q = ggml_permute(ctx0, indexer_q, 0, 2, 1, 3); + cb(indexer_q, "indexer_q", il); + indexer_k = ggml_permute(ctx0, indexer_k, 0, 2, 1, 3); + cb(indexer_k, "indexer_k", il); + + ggml_tensor * indexer_kq = ggml_mul_mat(ctx0, indexer_k, indexer_q); + cb(indexer_kq, "indexer_kq", il); + + // ReLU requires contiguous tensors + indexer_kq = ggml_cont(ctx0, ggml_permute(ctx0, indexer_kq, 2, 1, 0, 3)); + cb(indexer_kq, "indexer_kq", il); + + indexer_score = ggml_relu(ctx0, indexer_kq); + cb(indexer_score, "indexer_score", il); + + indexer_score = ggml_mul(ctx0, indexer_score, indexer_weights); + cb(indexer_score, "indexer_score", il); + + // sum by q n_indexer_head dimension + indexer_score = ggml_sum_rows(ctx0, indexer_score); + cb(indexer_score, "indexer_score", il); + + // permute result to match KQ mask + indexer_score = ggml_cont(ctx0, ggml_permute(ctx0, indexer_score, 2, 1, 0, 3)); + cb(indexer_score, "indexer_score", il); + + ggml_tensor * indexer_kq_mask = inp_attn->get_dsa()->get_kq_mask_lid(); + indexer_score = ggml_add(ctx0, indexer_score, indexer_kq_mask); + cb(indexer_score, "indexer_score", il); + } + + // get indices of top k indexer scores + uint32_t n_top_k = indexer_score->ne[0] < n_indexer_top_k ? indexer_score->ne[0] : n_indexer_top_k; + top_k = ggml_cont(ctx0, ggml_top_k(ctx0, indexer_score, n_top_k)); + cb(top_k, "top_k", il); + } + + ggml_tensor * q = ggml_mul_mat(ctx0, model.layers[il].wq_b, qr); + cb(q, "q", il); + + // split into {n_embd_head_qk_nope, n_head_l, n_tokens} + ggml_tensor * q_nope = + ggml_view_3d(ctx0, q, n_embd_head_qk_nope, n_head_l, n_tokens, ggml_row_size(q->type, n_embd_head_k_mla), + ggml_row_size(q->type, n_embd_head_k_mla) * n_head_l, 0); + cb(q_nope, "q_nope", il); + + // and {n_embd_head_qk_rope, n_head_l, n_tokens} + ggml_tensor * q_pe = ggml_view_3d( + ctx0, q, n_embd_head_qk_rope, n_head_l, n_tokens, ggml_row_size(q->type, n_embd_head_k_mla), + ggml_row_size(q->type, n_embd_head_k_mla) * n_head_l, ggml_row_size(q->type, n_embd_head_qk_nope)); + cb(q_pe, "q_pe", il); + + ggml_tensor * kv_cmpr_pe = ggml_mul_mat(ctx0, model.layers[il].wkv_a_mqa, cur); + cb(kv_cmpr_pe, "kv_cmpr_pe", il); + + // split into {kv_lora_rank, n_tokens} + ggml_tensor * kv_cmpr = + ggml_view_2d(ctx0, kv_cmpr_pe, kv_lora_rank, n_tokens, + ggml_row_size(kv_cmpr_pe->type, kv_lora_rank + n_embd_head_qk_rope), 0); + cb(kv_cmpr, "kv_cmpr", il); + + // and {n_embd_head_qk_rope, 1, n_tokens} + ggml_tensor * k_pe = ggml_view_3d(ctx0, kv_cmpr_pe, n_embd_head_qk_rope, 1, n_tokens, + ggml_row_size(kv_cmpr_pe->type, kv_lora_rank + n_embd_head_qk_rope), + ggml_row_size(kv_cmpr_pe->type, kv_lora_rank + n_embd_head_qk_rope), + ggml_row_size(kv_cmpr_pe->type, kv_lora_rank)); + cb(k_pe, "k_pe", il); + + // norm on the shared rope key, applied before rope + k_pe = build_norm(k_pe, model.layers[il].attn_k_norm, nullptr, LLM_NORM_RMS, il); + cb(k_pe, "k_pe", il); + + q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base_l, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + cb(q_pe, "q_pe", il); + + k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base_l, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + cb(k_pe, "k_pe", il); + + kv_cmpr = build_norm(kv_cmpr, model.layers[il].attn_kv_a_norm, nullptr, LLM_NORM_RMS, il); + cb(kv_cmpr, "kv_cmpr", il); + + // MLA attention with the absorption optimization + { + // {n_embd_head_qk_nope, n_tokens, n_head_l} + q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3); + cb(q_nope, "q_nope_perm", il); + + // {n_embd_head_qk_nope, kv_lora_rank, n_head_l} x {n_embd_head_qk_nope, n_tokens, n_head_l} + ggml_tensor * q_nope_absorbed = ggml_mul_mat(ctx0, model.layers[il].wk_b, q_nope); + cb(q_nope_absorbed, "q_nope_absorbed", il); + + // {kv_lora_rank, n_head_l, n_tokens} + q_nope_absorbed = ggml_permute(ctx0, q_nope_absorbed, 0, 2, 1, 3); + cb(q_nope_absorbed, "q_nope_absorbed_perm", il); + + // {n_embd_head_qk_rope + kv_lora_rank, n_head_l, n_tokens} + ggml_tensor * Qcur = ggml_concat(ctx0, q_nope_absorbed, q_pe, 0); + cb(Qcur, "Qcur", il); + + kv_cmpr = ggml_reshape_3d(ctx0, kv_cmpr, kv_lora_rank, 1, n_tokens); + cb(kv_cmpr, "kv_cmpr_reshape", il); + + // {n_embd_head_qk_rope + kv_lora_rank, 1, n_tokens} + ggml_tensor * Kcur = ggml_concat(ctx0, kv_cmpr, k_pe, 0); + cb(Kcur, "Kcur", il); + + // {kv_lora_rank, 1, n_tokens} + ggml_tensor * Vcur = kv_cmpr; + cb(Vcur, "Vcur", il); + + // apply the head-wise output gate before o_proj, so wo stays out of build_attn + if (is_swa) { + cur = build_attn(inp_attn->get_swa(), + nullptr, nullptr, nullptr, + Qcur, Kcur, Vcur, nullptr, nullptr, model.layers[il].wv_b, kq_scale, il); + } else { + cur = build_attn(inp_attn->get_dsa(), + nullptr, nullptr, nullptr, + Qcur, Kcur, Vcur, nullptr, nullptr, model.layers[il].wv_b, top_k, kq_scale, il); + } + cb(cur, "attn_out", il); + + ggml_tensor * gate = build_lora_mm(model.layers[il].wqkv_gate, attn_inp); + cb(gate, "attn_gate", il); + + gate = ggml_sigmoid(ctx0, gate); + cb(gate, "attn_gate_sigmoid", il); + + // broadcast the per-head gate over the head dimension + ggml_tensor * attn_3d = ggml_reshape_3d(ctx0, cur, n_embd_head_v_mla, n_head_l, n_tokens); + ggml_tensor * gate_3d = ggml_reshape_3d(ctx0, gate, 1, n_head_l, n_tokens); + attn_3d = ggml_mul(ctx0, attn_3d, gate_3d); + cb(attn_3d, "attn_gated", il); + + cur = ggml_reshape_2d(ctx0, attn_3d, n_embd_head_v_mla * n_head_l, n_tokens); + + cur = build_lora_mm(model.layers[il].wo, cur, model.layers[il].wo_s); + cb(cur, "attn_output", il); + } + } + + if (il == n_layer - 1 && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); + inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids); + } + + ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA); + cb(ffn_inp, "ffn_inp", il); + + cur = build_norm(ffn_inp, model.layers[il].ffn_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "ffn_norm", il); + + if ((uint32_t) il < hparams.n_layer_dense_lead) { + cur = build_ffn(cur, + model.layers[il].ffn_up, NULL, model.layers[il].ffn_up_s, + model.layers[il].ffn_gate, NULL, model.layers[il].ffn_gate_s, + model.layers[il].ffn_down, NULL, model.layers[il].ffn_down_s, + NULL, LLM_FFN_SILU, LLM_FFN_PAR, il); + cb(cur, "ffn_out", il); + } else { + ggml_tensor * moe_out = build_moe_ffn(cur, + model.layers[il].ffn_gate_inp, + model.layers[il].ffn_up_exps, + model.layers[il].ffn_gate_exps, + model.layers[il].ffn_down_exps, + model.layers[il].ffn_exp_probs_b, + n_expert, n_expert_used, + LLM_FFN_SILU, hparams.expert_weights_norm, + hparams.expert_weights_scale, + (llama_expert_gating_func_type) hparams.expert_gating_func, + il, + nullptr, + model.layers[il].ffn_gate_up_exps, + model.layers[il].ffn_up_exps_s, + model.layers[il].ffn_gate_exps_s, + model.layers[il].ffn_down_exps_s); + cb(moe_out, "ffn_moe_out", il); + + ggml_tensor * ffn_shexp = + build_ffn(cur, + model.layers[il].ffn_up_shexp, NULL, model.layers[il].ffn_up_shexp_s, + model.layers[il].ffn_gate_shexp, NULL, model.layers[il].ffn_gate_shexp_s, + model.layers[il].ffn_down_shexp, NULL, model.layers[il].ffn_down_shexp_s, + NULL, LLM_FFN_SILU, LLM_FFN_PAR, il); + cb(ffn_shexp, "ffn_shexp", il); + + cur = ggml_add(ctx0, moe_out, ffn_shexp); + cb(cur, "ffn_out", il); + } + + cur = ggml_add(ctx0, cur, ffn_inp); + + cur = build_cvec(cur, il); + cb(cur, "l_out", il); + + inpL = cur; + } + + cur = inpL; + + cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); + + cb(cur, "result_norm", -1); + res->t_embd = cur; + + cur = ggml_mul_mat(ctx0, model.output, cur); + + cb(cur, "result_output", -1); + res->t_logits = cur; + + ggml_build_forward_expand(gf, cur); +} diff --git a/src/models/models.h b/src/models/models.h index 1dd30dfd1..157b05dc0 100644 --- a/src/models/models.h +++ b/src/models/models.h @@ -1156,6 +1156,18 @@ struct llama_model_deepseek32 : public llama_model_base { }; +struct llama_model_dots3note : public llama_model_base { + llama_model_dots3note(const struct llama_model_params & params) : llama_model_base(params) {} + void load_arch_hparams(llama_model_loader & ml) override; + void load_arch_tensors(llama_model_loader & ml) override; + + struct graph : public llm_graph_context { + graph(const llama_model & model, const llm_graph_params & params); + }; + + std::unique_ptr build_arch_graph(const llm_graph_params & params) const override; +}; + struct llama_model_deepseek4 : public llama_model_base { llama_model_deepseek4(const struct llama_model_params & params) : llama_model_base(params) {} void load_arch_hparams(llama_model_loader & ml) override; diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp index 4eb3763cb..07e3a7a11 100644 --- a/tests/test-llama-archs.cpp +++ b/tests/test-llama-archs.cpp @@ -104,6 +104,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) { } else if (arch == LLM_ARCH_DEEPSEEK2 || arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA + || arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_KIMI_LINEAR || arch == LLM_ARCH_BAILINGMOE3 || arch == LLM_ARCH_KIMI_K3 @@ -166,6 +167,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) { if (arch == LLM_ARCH_DEEPSEEK2 || arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA + || arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_KIMI_LINEAR || arch == LLM_ARCH_BAILINGMOE3 || arch == LLM_ARCH_KIMI_K3 @@ -175,6 +177,22 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) { ms.add_kv(LLM_KV_ROPE_DIMENSION_COUNT, uint32_t(64)); ms.add_kv(LLM_KV_ATTENTION_KEY_LENGTH_MLA, uint32_t(192)); ms.add_kv(LLM_KV_ATTENTION_VALUE_LENGTH_MLA, uint32_t(128)); + if (arch == LLM_ARCH_DOTS3NOTE) { + // SWA layers reuse the same MLA geometry as the full layers in this fixture + ms.add_kv(LLM_KV_ATTENTION_KV_LORA_RANK_SWA, uint32_t(512)); + ms.add_kv(LLM_KV_ATTENTION_KEY_LENGTH_SWA, uint32_t(576)); + ms.add_kv(LLM_KV_ATTENTION_VALUE_LENGTH_SWA, uint32_t(512)); + ms.add_kv(LLM_KV_ATTENTION_KEY_LENGTH_MLA_SWA, uint32_t(192)); + ms.add_kv(LLM_KV_ATTENTION_VALUE_LENGTH_MLA_SWA, uint32_t(128)); + ms.add_kv(LLM_KV_ROPE_FREQ_BASE_SWA, 10000.0f); + // indexer on the full-attention layers (inverse of the swa pattern) + std::vector indexer_types; + indexer_types.reserve(n_layer); + for (uint32_t il = 0; il < n_layer; il++) { + indexer_types.push_back(il % 2 ? 0 : 1); + } + ms.add_kv(LLM_KV_ATTENTION_INDEXER_TYPES, indexer_types); + } } else if (arch == LLM_ARCH_MINIMAX_M3) { // partial rotary: n_rot must not exceed the indexer key length (64) ms.add_kv(LLM_KV_ROPE_DIMENSION_COUNT, uint32_t(64)); @@ -197,7 +215,8 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) { ms.add_kv(LLM_KV_ROPE_FREQ_BASE_SWA, 10000.0f); // SWA pattern: every 5th layer is full attention (matches E2B layer_types) ms.add_kv(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, uint32_t(5)); - } else if (arch == LLM_ARCH_COHERE2MOE || arch == LLM_ARCH_MIMO2 || arch == LLM_ARCH_STEP35 || arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_GRANITE_SWA) { + } else if (arch == LLM_ARCH_COHERE2MOE || arch == LLM_ARCH_MIMO2 || arch == LLM_ARCH_STEP35 || + arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_GRANITE_SWA || arch == LLM_ARCH_DOTS3NOTE) { std::vector pattern; pattern.reserve(n_layer); for (uint32_t il = 0; il < n_layer; il++) { @@ -365,6 +384,7 @@ static bool moe_mandatory(const llm_arch arch) { case LLM_ARCH_DEEPSEEK: case LLM_ARCH_DEEPSEEK2: case LLM_ARCH_DEEPSEEK32: + case LLM_ARCH_DOTS3NOTE: case LLM_ARCH_GLM4_MOE: case LLM_ARCH_GLM_DSA: case LLM_ARCH_EXAONE_MOE: From a3b9c23ead5054832c5ed4b43f652fa1123869a3 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 21 Aug 2026 21:41:25 +0300 Subject: [PATCH 02/16] ci : fix empty release_id in make-release upload step (#27516) The 'Create release' step had no id, so steps.create_release.outputs.id resolved to an empty string in the 'Upload nightly-tag.txt' step. The uploadReleaseAsset call then hit /releases//assets and failed with HTTP 404 (Unhandled error: HttpError), e.g. run 32513839499. Add id: create_release to the step; the action already exposes the id output. Assisted-by: pi:llama.cpp/Qwen3.8-27B --- .github/workflows/make-release.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/make-release.yml b/.github/workflows/make-release.yml index 451ec261f..7d7c9d1e5 100644 --- a/.github/workflows/make-release.yml +++ b/.github/workflows/make-release.yml @@ -70,6 +70,7 @@ jobs: cat nightly-tag.txt - name: Create release + id: create_release if: ${{ github.event.inputs.dry_run == 'false' }} uses: ggml-org/action-create-release@v1 env: From 9a286ac98d2cab74231bd3f1fc3f2b8bdf05422e Mon Sep 17 00:00:00 2001 From: Niklas Wenzel Date: Fri, 21 Aug 2026 20:49:27 +0200 Subject: [PATCH 03/16] docs: improve Windows build instructions (#27381) --- docs/build.md | 25 +++++++++++++++---------- 1 file changed, 15 insertions(+), 10 deletions(-) diff --git a/docs/build.md b/docs/build.md index 45fe7f17a..ed48e7a05 100644 --- a/docs/build.md +++ b/docs/build.md @@ -70,18 +70,23 @@ cmake --build build --config Release - Tab Workload: Desktop-development with C++ - Tab Components (select quickly via search): C++-_CMake_ Tools for Windows, _Git_ for Windows, C++-_Clang_ Compiler for Windows, MS-Build Support for LLVM-Toolset (clang) - Please remember to always use a Developer Command Prompt / PowerShell for VS2022 for git, build, test - - For Windows on ARM (arm64, WoA) build with: - ```bash - cmake --preset arm64-windows-llvm-release -D GGML_OPENMP_FETCH=ON - cmake --build build-arm64-windows-llvm-release - ``` - `GGML_OPENMP_FETCH` downloads the official LLVM OpenMP runtime and requires Clang, 7-Zip and network access during configuration. CMake selects the runtime from the target architecture, so this also works when cross-compiling for WoA from x64. The extracted header, import library, DLL and OpenMP license are placed under `build/_deps`. The build copies `libomp.dll` and `LICENSE-LLVM-OpenMP` to the runtime output directory and installs them together. Omit the option to use CMake's normal OpenMP detection, or pass `-D GGML_OPENMP=OFF` to disable OpenMP. - For building with ninja generator and clang compiler as default: - -set path:set LIB=C:\Program Files (x86)\Windows Kits\10\Lib\10.0.22621.0\um\x64;C:\Program Files\Microsoft Visual Studio\2022\Community\VC\Tools\MSVC\14.41.34120\lib\x64\uwp;C:\Program Files (x86)\Windows Kits\10\Lib\10.0.22621.0\ucrt\x64 + - For Windows on ARM (arm64, WoA), build with: ```bash - cmake --preset x64-windows-llvm-release - cmake --build build-x64-windows-llvm-release + cmake --preset arm64-windows-llvm-release -D GGML_OPENMP_FETCH=ON + cmake --build build-arm64-windows-llvm-release ``` + - Use `ARM64 Native Tools Command Prompt for VS 2022` if you are building on an ARM64 machine. + - `GGML_OPENMP_FETCH` downloads the official LLVM OpenMP runtime and requires Clang, 7-Zip and network access during configuration. CMake selects the runtime from the target architecture, so this also works when cross-compiling for WoA from x64. The extracted header, import library, DLL and OpenMP license are placed under `build/_deps`. The build copies `libomp.dll` and `LICENSE-LLVM-OpenMP` to the runtime output directory and installs them together. Omit the option to use CMake's normal OpenMP detection, or pass `-D GGML_OPENMP=OFF` to disable OpenMP. + - For building with ninja generator and clang compiler as default: + - Set path: + ``` + set LIB=C:\Program Files (x86)\Windows Kits\10\Lib\10.0.22621.0\um\x64;C:\Program Files\Microsoft Visual Studio\2022\Community\VC\Tools\MSVC\14.41.34120\lib\x64\uwp;C:\Program Files (x86)\Windows Kits\10\Lib\10.0.22621.0\ucrt\x64 + ``` + - Run: + ```bash + cmake --preset x64-windows-llvm-release + cmake --build build-x64-windows-llvm-release + ``` - If you want HTTPS/TLS features, you may install OpenSSL development libraries. If not installed, the project will build and run without SSL support. - **Debian / Ubuntu:** `sudo apt-get install libssl-dev` - **Fedora / RHEL / Rocky / Alma:** `sudo dnf install openssl-devel` From 3af988fabcf79fd81f8720505e684d2aa5bfc786 Mon Sep 17 00:00:00 2001 From: Hongqiang Wang Date: Fri, 21 Aug 2026 14:24:33 -0700 Subject: [PATCH 04/16] opencl: fold the gpt-oss MoE per-expert bias adds into the epilogue (op/kernel fusion) (#26431) * opencl: fold the gpt-oss MoE bias adds into swiglu_oai Default on, opt out with GGML_OPENCL_FUSE_MOE_BIAS_GLU=0. * opencl: fold the MoE down-projection bias into the combine Default on, opt out with GGML_OPENCL_FUSE_MOE_BIAS_COMBINE=0. --- ggml/src/ggml-opencl/CMakeLists.txt | 1 + ggml/src/ggml-opencl/ggml-opencl.cpp | 353 ++++++++++++++++++ .../src/ggml-opencl/kernels/moe_add_id_glu.cl | 76 ++++ ggml/src/ggml-opencl/kernels/moe_combine.cl | 43 +++ 4 files changed, 473 insertions(+) create mode 100644 ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index 72334d5ce..1f62ce1c6 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -63,6 +63,7 @@ endfunction() set(GGML_OPENCL_KERNELS add add_id + moe_add_id_glu argsort tri fill diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 84f854cc2..64f3325b2 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -577,6 +577,12 @@ struct ggml_backend_opencl_context { // whether fuse moe combine cl_uint fuse_moe_combine; + // whether to fold the MoE bias adds into swiglu_oai + cl_uint fuse_moe_bias_glu; + + // whether to fold the MoE down-projection bias add into the combine + cl_uint fuse_moe_bias_combine; + bool adreno_has_large_buffer; bool adreno_use_large_buffer; bool adreno_use_bin_kernels; @@ -658,6 +664,7 @@ struct ggml_backend_opencl_context { cl_program program_add; cl_program program_add_id; + cl_program program_moe_add_id_glu; cl_program program_clamp; cl_program program_cvt; cl_program program_diag_mask_inf; @@ -723,6 +730,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_div, kernel_div_row, kernel_div_f16, kernel_div_row_f16; cl_kernel kernel_sub, kernel_sub_row, kernel_sub_f16, kernel_sub_row_f16; cl_kernel kernel_add_id; + cl_kernel kernel_add_id_add_id_swiglu_oai; cl_kernel kernel_scale_f32, kernel_scale_f32_4; cl_kernel kernel_sqr_cont_f32, kernel_sqr_cont_f32_4, kernel_sqr_cont_f16, kernel_sqr_cont_f16_4; cl_kernel kernel_sqrt_cont_f32, kernel_sqrt_cont_f32_4, kernel_sqrt_cont_f16, kernel_sqrt_cont_f16_4; @@ -899,6 +907,7 @@ struct ggml_backend_opencl_context { cl_kernel kernel_moe_histogram, kernel_moe_scan, kernel_moe_fill, kernel_moe_scatter; cl_kernel kernel_moe_scatter_stable = nullptr; // deterministic slot assignment cl_kernel kernel_moe_combine_f32 = nullptr; // fused router-weight mul + cross-expert sum + cl_kernel kernel_moe_combine_bias_f32 = nullptr; // same, with the down-projection bias add folded in cl_kernel kernel_mul_mv_id_q4_0_f32_8x_flat; cl_kernel kernel_mul_mv_id_q8_0_f32, kernel_mul_mv_id_q8_0_f32_flat; cl_kernel kernel_mul_mv_id_mxfp4_f32; @@ -1346,6 +1355,23 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { GGML_LOG_CONT("."); } + // moe_add_id_glu + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "moe_add_id_glu.cl.h" + }; +#else + const std::string kernel_src = read_file("moe_add_id_glu.cl"); +#endif + backend_ctx->program_moe_add_id_glu = + build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_add_id_add_id_swiglu_oai = + clCreateKernel(backend_ctx->program_moe_add_id_glu, "kernel_add_id_add_id_swiglu_oai", &err), err)); + GGML_LOG_CONT("."); + } + // tri { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -3276,6 +3302,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { backend_ctx, kernel_src.c_str(), compile_opts); CL_CHECK((backend_ctx->kernel_moe_combine_f32 = clCreateKernel(prog, "kernel_moe_combine_f32", &err), err)); + CL_CHECK((backend_ctx->kernel_moe_combine_bias_f32 = + clCreateKernel(prog, "kernel_moe_combine_bias_f32", &err), err)); CL_CHECK(clReleaseProgram(prog)); GGML_LOG_CONT("."); } @@ -6012,6 +6040,12 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) { backend_ctx->adreno_moe_ragged_skip_gran = (ragged_gran_env != NULL) ? atoi(ragged_gran_env) : 8; // whether fuse moe combine + static const char * fuse_moe_bias_glu_env = getenv("GGML_OPENCL_FUSE_MOE_BIAS_GLU"); + backend_ctx->fuse_moe_bias_glu = fuse_moe_bias_glu_env == NULL ? 1 : (atoi(fuse_moe_bias_glu_env) != 0); + + static const char * fuse_moe_bias_combine_env = getenv("GGML_OPENCL_FUSE_MOE_BIAS_COMBINE"); + backend_ctx->fuse_moe_bias_combine = fuse_moe_bias_combine_env == NULL ? 1 : (atoi(fuse_moe_bias_combine_env) != 0); + static const char * fuse_moe_combine_env = getenv("GGML_OPENCL_FUSE_MOE_COMBINE"); backend_ctx->fuse_moe_combine = fuse_moe_combine_env == NULL ? 1 : (atoi(fuse_moe_combine_env) != 0); @@ -6880,6 +6914,300 @@ static bool ggml_opencl_can_fuse_moe_combine(const struct ggml_cgraph * cgraph, return true; } +// Detect the gpt-oss MoE bias+activation epilogue on the PREFILL path: +// {MUL_MAT_ID(gate), ADD_ID(gate_bias), MUL_MAT_ID(up), ADD_ID(up_bias), GLU(swiglu_oai)}. +// The two matmuls still run as their own dispatches (the prefill GEMM is the vendor's); +// what collapses is the epilogue — both add_id passes are in-place read-modify-writes of a +// tensor the GLU immediately reads again, so they are three full passes over the same +// [n_ff, n_expert_used, n_tokens] f32 tensor where one suffices. +// +// The decode counterpart is handled by the mxfp4 fused GEMV arm in ggml_opencl_can_fuse, +// which folds the matmul too; this one deliberately fires only when that cannot (ne[2] > 1). +static bool ggml_opencl_can_fuse_moe_bias_glu(const struct ggml_cgraph * cgraph, int node_idx) { + if (node_idx + 4 >= cgraph->n_nodes) { + return false; + } + + const enum ggml_op mg_ops[] = { GGML_OP_MUL_MAT_ID, GGML_OP_ADD_ID, GGML_OP_MUL_MAT_ID, GGML_OP_ADD_ID, GGML_OP_GLU }; + const int mg_out[] = { node_idx + 4 }; + if (!ggml_can_fuse_subgraph(cgraph, node_idx, 5, mg_ops, mg_out, 1)) { + return false; + } + + const ggml_tensor * gmm = cgraph->nodes[node_idx]; + const ggml_tensor * gad = cgraph->nodes[node_idx+1]; + const ggml_tensor * umm = cgraph->nodes[node_idx+2]; + const ggml_tensor * uad = cgraph->nodes[node_idx+3]; + const ggml_tensor * glu = cgraph->nodes[node_idx+4]; + + if (ggml_get_glu_op(glu) != GGML_GLU_OP_SWIGLU_OAI) { + return false; + } + // Prefill only — at one token the mxfp4 arm above folds the matmul as well. + if (gmm->src[1]->ne[2] == 1) { + return false; + } + // Wiring: both matmuls share the activation and the expert selection, each add_id + // biases its own matmul, and the GLU consumes the two biased results as separate + // operands (so the same-buffer ne00_off/ne10_off split path is not in play). + if (gad->src[0] != gmm || uad->src[0] != umm || + glu->src[0] != gad || glu->src[1] != uad || + umm->src[1] != gmm->src[1] || umm->src[2] != gmm->src[2]) { + return false; + } + // A swapped GLU would exchange the gate/up roles the fused kernel hard-codes. + if (ggml_get_op_params_i32(glu, 1)) { + return false; + } + if (gad->type != GGML_TYPE_F32 || uad->type != GGML_TYPE_F32 || glu->type != GGML_TYPE_F32) { + return false; + } + if (!gad->src[1] || gad->src[1]->type != GGML_TYPE_F32 || + !uad->src[1] || uad->src[1]->type != GGML_TYPE_F32) { + return false; + } + if (!gad->src[2] || gad->src[2]->type != GGML_TYPE_I32 || uad->src[2] != gad->src[2]) { + return false; + } + // Full width on both operands: the kernel writes one output element per input pair. + if (!ggml_are_same_shape(gad, uad) || glu->ne[0] != gad->ne[0] || + glu->ne[1] != gad->ne[1] || glu->ne[2] != gad->ne[2] || glu->ne[3] != gad->ne[3]) { + return false; + } + if (gad->ne[3] != 1) { + return false; + } + // The destination is addressed by (expert slot, token) rather than the GLU's flat row + // walk; those agree only for a contiguous destination. + if (!ggml_is_contiguous(glu) || !ggml_is_contiguous(gmm) || !ggml_is_contiguous(umm)) { + return false; + } + return true; +} + +static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst); + +// Runs the gate and up matmuls unchanged, then one kernel in place of +// add_id(gate) + add_id(up) + swiglu_oai. See ggml_opencl_can_fuse_moe_bias_glu. +static void ggml_cl_moe_bias_glu_fused(ggml_backend_t backend, ggml_tensor * gate_mm, const ggml_tensor * gate_add, + ggml_tensor * up_mm, const ggml_tensor * up_add, const ggml_tensor * glu) { + ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; + + ggml_cl_mul_mat_id(backend, gate_mm->src[0], gate_mm->src[1], gate_mm); + ggml_cl_mul_mat_id(backend, up_mm->src[0], up_mm->src[1], up_mm); + + const ggml_tensor * gbias = gate_add->src[1]; + const ggml_tensor * ubias = up_add->src[1]; + const ggml_tensor * ids = gate_add->src[2]; + + ggml_tensor_extra_cl * eg = (ggml_tensor_extra_cl *)gate_mm->extra; + ggml_tensor_extra_cl * egb = (ggml_tensor_extra_cl *)gbias->extra; + ggml_tensor_extra_cl * eu = (ggml_tensor_extra_cl *)up_mm->extra; + ggml_tensor_extra_cl * eub = (ggml_tensor_extra_cl *)ubias->extra; + ggml_tensor_extra_cl * ei = (ggml_tensor_extra_cl *)ids->extra; + ggml_tensor_extra_cl * ed = (ggml_tensor_extra_cl *)glu->extra; + + cl_ulong off_g = eg->offset + gate_mm->view_offs; + cl_ulong off_gb = egb->offset + gbias->view_offs; + cl_ulong off_u = eu->offset + up_mm->view_offs; + cl_ulong off_ub = eub->offset + ubias->view_offs; + cl_ulong off_i = ei->offset + ids->view_offs; + cl_ulong off_d = ed->offset + glu->view_offs; + + const cl_ulong nb01_g = gate_mm->nb[1]; + const cl_ulong nb02_g = gate_mm->nb[2]; + const cl_ulong nb01_u = up_mm->nb[1]; + const cl_ulong nb02_u = up_mm->nb[2]; + const cl_ulong nb11_g = gbias->nb[1]; + const cl_ulong nb11_u = ubias->nb[1]; + const cl_ulong nb21 = ids->nb[1]; + const cl_ulong nbd1 = glu->nb[1]; + const cl_ulong nbd2 = glu->nb[2]; + + const int ne0 = (int)glu->ne[0]; + const float alpha = ggml_get_op_params_f32(glu, 2); + const float limit = ggml_get_op_params_f32(glu, 3); + + cl_kernel kernel = backend_ctx->kernel_add_id_add_id_swiglu_oai; + + int i = 0; + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_mem), &eg->data_device)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &off_g)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_mem), &egb->data_device)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &off_gb)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_mem), &eu->data_device)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &off_u)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_mem), &eub->data_device)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &off_ub)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_mem), &ei->data_device)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &off_i)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_mem), &ed->data_device)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &off_d)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nb01_g)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nb02_g)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nb01_u)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nb02_u)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nb11_g)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nb11_u)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nb21)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nbd1)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(cl_ulong), &nbd2)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(float), &limit)); + CL_CHECK(clSetKernelArg(kernel, i++, sizeof(float), &alpha)); + + const int nth = MIN(ne0, (int) backend_ctx->get_kernel_workgroup_size(kernel)); + size_t global_work_size[] = { (size_t)glu->ne[1]*nth, (size_t)glu->ne[2], 1 }; + size_t local_work_size[] = { (size_t)nth, 1, 1 }; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, (ggml_tensor *)glu); +} + +// Fusion B: the MoE down-projection bias add feeding the combine. +// +// The graph runs ADD_ID(down_bias) and then immediately the combine subgraph +// {MUL(router weights), k VIEWs, k-1 ADDs}, and the ADD_ID's only consumer is that +// MUL. Since the ADD_ID is an in-place read-modify-write of a tensor the combine +// reads once more, the bias can be added inside the combine instead, dropping a +// full pass over [n_embd, k, n_tokens]. +// +// Shape checks for the combine tail are delegated to ggml_opencl_can_fuse_moe_combine +// (which also owns the n_nodes >= 32 bail and the experts/dst aliasing bail); what is +// added here is the ADD_ID wiring plus a subgraph check over the WHOLE run, so that +// the intermediate bias result is confirmed not to escape. +static bool ggml_opencl_can_fuse_moe_bias_combine(const struct ggml_cgraph * cgraph, int node_idx, + const ggml_tensor ** out_final_add) { + if (node_idx + 1 >= cgraph->n_nodes) { + return false; + } + const ggml_tensor * add = cgraph->nodes[node_idx]; + if (add->op != GGML_OP_ADD_ID) { + return false; + } + const ggml_tensor * mul = cgraph->nodes[node_idx+1]; + if (mul->op != GGML_OP_MUL || mul->src[0] != add) { + return false; + } + + const ggml_tensor * final_add = NULL; + if (!ggml_opencl_can_fuse_moe_combine(cgraph, node_idx+1, &final_add)) { + return false; + } + + const ggml_tensor * raw = add->src[0]; + const ggml_tensor * bias = add->src[1]; + const ggml_tensor * ids = add->src[2]; + if (!raw || !bias || !ids) { + return false; + } + if (raw->type != GGML_TYPE_F32 || bias->type != GGML_TYPE_F32 || + ids->type != GGML_TYPE_I32 || add->type != GGML_TYPE_F32) { + return false; + } + // The combine reads the raw matmul output with the strides it computed from the + // add_id result, so the two must have the same layout. + if (!ggml_are_same_shape(raw, add) || !ggml_is_contiguous(raw)) { + return false; + } + if (raw->nb[1] != add->nb[1] || raw->nb[2] != add->nb[2]) { + return false; + } + // ids is indexed as [expert slot, token]; the combine walks the same two axes. + if (ids->ne[0] < add->ne[1] || ids->ne[1] < add->ne[2]) { + return false; + } + + // Whole-run escape check: ADD_ID + MUL + k VIEWs + (k-1) ADDs, only the last node escapes. + const int k = (int)add->ne[1]; + const int n_nodes = 2 + k + (k - 1); + if (n_nodes >= 32 || node_idx + n_nodes > cgraph->n_nodes) { + return false; + } + enum ggml_op ops[32]; + int n = 0; + ops[n++] = GGML_OP_ADD_ID; + ops[n++] = GGML_OP_MUL; + for (int j = 0; j < k; ++j) ops[n++] = GGML_OP_VIEW; + for (int j = 0; j < k - 1; ++j) ops[n++] = GGML_OP_ADD; + const int outs[] = { node_idx + n_nodes - 1 }; + if (!ggml_can_fuse_subgraph(cgraph, node_idx, n_nodes, ops, outs, 1)) { + return false; + } + + *out_final_add = final_add; + return true; +} + + +// Fusion B dispatch: the combine, reading the RAW matmul output and adding the +// per-expert bias row inline. See ggml_opencl_can_fuse_moe_bias_combine. +static void ggml_cl_moe_bias_combine_fused(ggml_backend_t backend, const ggml_tensor * add, + const ggml_tensor * mul, const ggml_tensor * dst) { + ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; + + const ggml_tensor * experts = add->src[0]; // raw matmul output, bias not yet applied + const ggml_tensor * bias = add->src[1]; + const ggml_tensor * ids = add->src[2]; + const ggml_tensor * weights = mul->src[1]; + + ggml_tensor_extra_cl * ee = (ggml_tensor_extra_cl *)experts->extra; + ggml_tensor_extra_cl * eb = (ggml_tensor_extra_cl *)bias->extra; + ggml_tensor_extra_cl * ei = (ggml_tensor_extra_cl *)ids->extra; + ggml_tensor_extra_cl * ew = (ggml_tensor_extra_cl *)weights->extra; + ggml_tensor_extra_cl * ed = (ggml_tensor_extra_cl *)dst->extra; + cl_ulong off_e = ee->offset + experts->view_offs; + cl_ulong off_b = eb->offset + bias->view_offs; + cl_ulong off_i = ei->offset + ids->view_offs; + cl_ulong off_w = ew->offset + weights->view_offs; + cl_ulong off_d = ed->offset + dst->view_offs; + + const int n_embd4 = (int)(experts->ne[0] / 4); + const int k = (int)experts->ne[1]; + const int nt = (int)experts->ne[2]; + const cl_uint e1 = (cl_uint)(experts->nb[1] / sizeof(float)); + const cl_uint e2 = (cl_uint)(experts->nb[2] / sizeof(float)); + const cl_uint w1 = (cl_uint)(weights->nb[1] / sizeof(float)); + const cl_uint w2 = (cl_uint)(weights->nb[2] / sizeof(float)); + const cl_uint d1 = (cl_uint)(dst->nb[1] / sizeof(float)); + const cl_ulong nb_b1 = bias->nb[1]; + const cl_ulong nb_i1 = ids->nb[1]; + + const size_t w_bytes = ggml_nbytes(weights); + backend_ctx->prealloc_moe_combine_w.allocate(backend_ctx->context, w_bytes); + CL_CHECK(clEnqueueCopyBuffer(backend_ctx->queue, ew->data_device, backend_ctx->prealloc_moe_combine_w.buffer, + off_w, 0, w_bytes, 0, NULL, NULL)); + cl_mem w_dev = backend_ctx->prealloc_moe_combine_w.buffer; + cl_ulong w_off = 0; + + cl_kernel kernel = backend_ctx->kernel_moe_combine_bias_f32; + int a = 0; + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_mem), &ee->data_device)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_ulong), &off_e)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_mem), &w_dev)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_ulong), &w_off)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_mem), &eb->data_device)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_ulong), &off_b)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_mem), &ei->data_device)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_ulong), &off_i)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_mem), &ed->data_device)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_ulong), &off_d)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(int), &n_embd4)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(int), &k)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(int), &nt)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_uint), &e1)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_uint), &e2)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_uint), &w1)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_uint), &w2)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_uint), &d1)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_ulong), &nb_b1)); + CL_CHECK(clSetKernelArg(kernel, a++, sizeof(cl_ulong), &nb_i1)); + + size_t lws[2] = { 64, 1 }; + size_t gws[2] = { (size_t)(((n_embd4 + 63) / 64) * 64), (size_t)nt }; + backend_ctx->enqueue_ndrange_kernel(kernel, 2, gws, lws, (ggml_tensor *)dst); +} + + static void ggml_cl_moe_combine_fused(ggml_backend_t backend, const ggml_tensor * mul, const ggml_tensor * dst) { ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context; const ggml_tensor * experts = mul->src[0]; @@ -7034,6 +7362,31 @@ static ggml_status ggml_backend_opencl_graph_compute(ggml_backend_t backend, ggm } // Fuse the MoE combine: router-weight mul + cross-expert add chain -> // one weighted-sum-across-experts kernel. + // Fold the gpt-oss MoE bias epilogue: add_id(gate_bias) + add_id(up_bias) + + // glu(swiglu_oai) -> one kernel, leaving the two matmuls as their own dispatches. + // Both add_ids are in-place passes over a tensor the GLU reads again, so this + // drops two full read+write passes per layer. Opt out GGML_OPENCL_FUSE_MOE_BIAS_GLU=0. + if (backend_ctx->fuse_moe_bias_glu && !backend_ctx->disable_fusion && + ggml_opencl_can_fuse_moe_bias_glu(cgraph, i)) { + ggml_cl_moe_bias_glu_fused(backend, node, cgraph->nodes[i+1], cgraph->nodes[i+2], + cgraph->nodes[i+3], cgraph->nodes[i+4]); + i += 4; + continue; + } + + // Fold the MoE down-projection bias into the combine: add_id(down_bias) + the whole + // combine subgraph -> one kernel. Checked before the plain combine arm so the longer + // pattern wins. Opt out GGML_OPENCL_FUSE_MOE_BIAS_COMBINE=0. + if (backend_ctx->fuse_moe_bias_combine && backend_ctx->fuse_moe_combine && + !backend_ctx->disable_fusion) { + const ggml_tensor * bias_combine_out = nullptr; + if (ggml_opencl_can_fuse_moe_bias_combine(cgraph, i, &bias_combine_out)) { + ggml_cl_moe_bias_combine_fused(backend, node, cgraph->nodes[i+1], bias_combine_out); + i += 2 * (int)node->ne[1]; // ADD_ID + MUL + k VIEWs + (k-1) ADDs + continue; + } + } + if (backend_ctx->fuse_moe_combine && !backend_ctx->disable_fusion) { const ggml_tensor * combine_out = nullptr; if (ggml_opencl_can_fuse_moe_combine(cgraph, i, &combine_out)) { diff --git a/ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl b/ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl new file mode 100644 index 000000000..6a8e4fb17 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl @@ -0,0 +1,76 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +//------------------------------------------------------------------------------ +// add_id(gate) + add_id(up) + swiglu_oai, fused +// +// gpt-oss-class MoE FFNs run three full passes over the same +// [n_ff, n_expert_used, n_tokens] f32 tensor: a per-expert bias add on the gate +// matmul output, the same on the up matmul output, then swiglu_oai over the +// two. Both bias adds are in-place, so each costs a full read plus a full write +// of a tensor that is only read once more. Folding them into the swiglu pass +// leaves two reads and one write instead of six passes. +// +// Grouping matches kernel_add_id: group 0 = expert slot (i1), group 1 = token +// (i2). For a contiguous destination that addressing is identical to the flat +// row walk kernel_swiglu_oai uses, since row i1 + i2*ne1 sits at +// i1*nb1 + i2*ne1*nb1. +//------------------------------------------------------------------------------ +kernel void kernel_add_id_add_id_swiglu_oai( + global char * src_g, + ulong offset_g, + global char * src_gb, + ulong offset_gb, + global char * src_u, + ulong offset_u, + global char * src_ub, + ulong offset_ub, + global char * src_ids, + ulong offset_ids, + global char * dst, + ulong offsetd, + ulong nb01_g, + ulong nb02_g, + ulong nb01_u, + ulong nb02_u, + ulong nb11_g, + ulong nb11_u, + ulong nb21, + ulong nbd1, + ulong nbd2, + int ne0, + float limit, + float alpha +) { + src_g = (global char *)(src_g + offset_g); + src_gb = (global char *)(src_gb + offset_gb); + src_u = (global char *)(src_u + offset_u); + src_ub = (global char *)(src_ub + offset_ub); + src_ids = (global char *)(src_ids + offset_ids); + dst = (global char *)(dst + offsetd); + + const int i1 = get_group_id(0); + const int i2 = get_group_id(1); + + // The ids tensor is a view into a [n_expert, n_tokens] buffer, so its row + // stride is nb21 and the k selected ids are NOT contiguous per token. + const int i11 = *((global const int *) (src_ids + i1*sizeof(int) + i2*nb21)); + + global const float * g_row = (global const float *)(src_g + i1*nb01_g + i2*nb02_g); + global const float * u_row = (global const float *)(src_u + i1*nb01_u + i2*nb02_u); + global const float * gb_row = (global const float *)(src_gb + i11*nb11_g); + global const float * ub_row = (global const float *)(src_ub + i11*nb11_u); + global float * d_row = (global float *)(dst + i1*nbd1 + i2*nbd2); + + for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) { + float x0 = g_row[i0] + gb_row[i0]; + float x1 = u_row[i0] + ub_row[i0]; + + x0 = min(x0, limit); + x1 = max(min(x1, limit), -limit); + + float out_glu = x0 / (1.0f + exp(-x0 * alpha)); + out_glu = out_glu * (1.0f + x1); + + d_row[i0] = out_glu; + } +} diff --git a/ggml/src/ggml-opencl/kernels/moe_combine.cl b/ggml/src/ggml-opencl/kernels/moe_combine.cl index c195f1472..acd08dbe6 100644 --- a/ggml/src/ggml-opencl/kernels/moe_combine.cl +++ b/ggml/src/ggml-opencl/kernels/moe_combine.cl @@ -8,6 +8,49 @@ // buffer and the k-1 elementwise add round-trips). Vectorized float4 over rows. // strides e1/e2/w1/w2/d1 are in ELEMENTS (floats). +// Same weighted sum, with the per-expert bias add folded in. +// +// The MoE down projection's bias is applied by an in-place add_id whose only +// consumer is this combine, so it costs a full read plus a full write of a +// tensor that is read once more immediately afterwards. Reading the raw matmul +// output here and adding the bias row while it is already in registers removes +// that pass. Kept as a separate kernel so the unfused path is untouched. +__kernel void kernel_moe_combine_bias_f32( + __global const char * e_buf, ulong off_e, + __global const char * w_buf, ulong off_w, + __global const char * b_buf, ulong off_b, // per-expert bias rows + __global const char * i_buf, ulong off_i, // expert ids + __global char * d_buf, ulong off_d, + int n_embd4, // n_embd / 4 + int k, // n_expert_used + int n_tokens, + uint e1, uint e2, // experts strides (elements): per-expert, per-token + uint w1, uint w2, // weights strides (elements) + uint d1, // dst per-token stride (elements) + ulong nb_b1, // bias row stride (bytes) + ulong nb_i1) // ids row stride (bytes) - ids is a view, not packed +{ + const uint r4 = get_global_id(0); + const uint tok = get_global_id(1); + if (r4 >= (uint)n_embd4 || tok >= (uint)n_tokens) return; + + __global const float * E = (__global const float *)(e_buf + off_e) + tok*e2 + r4*4u; + __global const float * W = (__global const float *)(w_buf + off_w) + tok*w2; + __global const char * B = b_buf + off_b; + __global const char * I = i_buf + off_i + (ulong)tok*nb_i1; + + float4 acc = (float4)(0.0f); + for (int e = 0; e < k; ++e) { + const int i11 = *((__global const int *)(I + (ulong)e*sizeof(int))); + __global const float * Brow = (__global const float *)(B + (ulong)i11*nb_b1) + r4*4u; + const float4 v = vload4(0, E + (uint)e*e1) + vload4(0, Brow); + acc = mad(v, (float4)(W[(uint)e*w1]), acc); + } + + __global float * D = (__global float *)(d_buf + off_d) + tok*d1 + r4*4u; + vstore4(acc, 0, D); +} + __kernel void kernel_moe_combine_f32( __global const char * e_buf, ulong off_e, __global const char * w_buf, ulong off_w, From d775b8967a46d8beb110d444aa3b8938179e0dd8 Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Sat, 22 Aug 2026 01:38:05 +0200 Subject: [PATCH 05/16] mtmd: support webp via ffmpeg (#27520) --- tools/mtmd/mtmd-helper.cpp | 49 ++++++++++++++++++++++++++++++++++++++ tools/mtmd/mtmd-helper.h | 1 + 2 files changed, 50 insertions(+) diff --git a/tools/mtmd/mtmd-helper.cpp b/tools/mtmd/mtmd-helper.cpp index 8b8e5f5f6..f719323c0 100644 --- a/tools/mtmd/mtmd-helper.cpp +++ b/tools/mtmd/mtmd-helper.cpp @@ -358,6 +358,15 @@ static bool decode_audio_from_buf(const unsigned char * buf_in, size_t len, int } // namespace audio_helpers +static bool is_webp_file(const unsigned char * buf, size_t len) { + // WEBP ref: https://developers.google.com/speed/webp/docs/riff_container + return len >= 12 && memcmp(buf, "RIFF", 4) == 0 && memcmp(buf + 8, "WEBP", 4) == 0; +} + +#ifdef MTMD_VIDEO +static mtmd_bitmap * decode_webp_with_ffmpeg(mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder); +#endif + mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(mtmd_context * ctx, const unsigned char * buf, size_t len, bool placeholder) { // calculate the hash if needed std::string id; @@ -397,6 +406,19 @@ mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(mtmd_context * ctx, // otherwise, fallthrough to video decoding (if supported) } +#ifdef MTMD_VIDEO + // stb_image does not support webp; decode it with ffmpeg as a single frame + if (!result && is_webp_file(buf, len)) { + result = decode_webp_with_ffmpeg(ctx, buf, len, placeholder); + if (!result) { + LOG_ERR("%s: failed to decode webp buffer\n", __func__); + return {nullptr, nullptr}; + } + mtmd_bitmap_set_id(result, id.empty() ? nullptr : id.c_str()); + return {result, nullptr}; + } +#endif + // last try: load as video #ifdef MTMD_VIDEO if (!result) { @@ -820,6 +842,33 @@ static std::string video_resolve_bin(const char * bin_dir, const char * name) { return result; } +#ifdef MTMD_VIDEO +static mtmd_bitmap * decode_webp_with_ffmpeg(mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder) { + auto params = mtmd_helper_video_init_params_default(); + mtmd_helper_video vctx; + vctx.mctx = mctx; + vctx.input_buf.assign(buf, buf + len); + vctx.ffmpeg_bin = video_resolve_bin(params.ffmpeg_bin_dir, "ffmpeg"); + vctx.ffprobe_bin = video_resolve_bin(params.ffmpeg_bin_dir, "ffprobe"); + if (!vctx.probe(0.0f)) { + return nullptr; + } + if (placeholder) { + return mtmd_bitmap_init(vctx.info.width, vctx.info.height, nullptr); + } + // still image: the fps filter would output no frame, so disable it + vctx.fps_target = 0.0f; + if (!vctx.start_ffmpeg(0.0f)) { + return nullptr; + } + mtmd_bitmap * frame = vctx.read_next_frame(); + if (frame) { + mtmd_bitmap_set_mergeable(frame, false); + } + return frame; +} +#endif + mtmd_helper_video * mtmd_helper_video_init( mtmd_context * mctx, const char * path, diff --git a/tools/mtmd/mtmd-helper.h b/tools/mtmd/mtmd-helper.h index 5c6b92419..58dfb1525 100644 --- a/tools/mtmd/mtmd-helper.h +++ b/tools/mtmd/mtmd-helper.h @@ -45,6 +45,7 @@ MTMD_API struct mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_file(mtm // helper function to construct a mtmd_bitmap from a buffer containing a file // supported formats: // image: formats supported by stb_image: jpg, png, bmp, gif, etc. +// webp is decoded via ffmpeg, requires MTMD_VIDEO build with ffmpeg in PATH // audio: formats supported by miniaudio: wav, mp3, flac // note: // - for now, video input is only supported via C++ helper functions From 2100e592600e4538496fd5201cbe6a7f8fbeb1e0 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Sat, 22 Aug 2026 08:25:00 +0300 Subject: [PATCH 06/16] readme : update badges (#27531) --- README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 85726c8c0..33a7b9816 100644 --- a/README.md +++ b/README.md @@ -7,9 +7,9 @@ LLM inference in C/C++ [![License: MIT](https://img.shields.io/badge/license-MIT-blue.svg)](https://opensource.org/licenses/MIT) -[![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp?filter=v*)](https://github.com/ggml-org/llama.cpp/releases?q=tag:v0) -[![Nightly](https://img.shields.io/github/v/release/ggml-org/llama.cpp?label=nightly)](https://github.com/ggml-org/llama.cpp/releases) -[![Server](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml) +[![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp?filter=v*&color=brightgreen)](https://github.com/ggml-org/llama.cpp/releases?q=tag:v0) +[![Nightly](https://img.shields.io/github/v/release/ggml-org/llama.cpp?label=nightly&filter=b*&color=orange)](https://github.com/ggml-org/llama.cpp/releases?q=b) +[![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/server.yml?label=Server)](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml) [![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/docker.yml?label=Docker)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml) [![Winget](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/winget.yml?label=Winget)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml) From 3aeb924628c7113b9764b2ec2f3aebc6c51b27a5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?= Date: Sat, 22 Aug 2026 09:08:07 +0200 Subject: [PATCH 07/16] readme : fix server badge alt (#27533) --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 33a7b9816..1f2076038 100644 --- a/README.md +++ b/README.md @@ -9,7 +9,7 @@ [![License: MIT](https://img.shields.io/badge/license-MIT-blue.svg)](https://opensource.org/licenses/MIT) [![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp?filter=v*&color=brightgreen)](https://github.com/ggml-org/llama.cpp/releases?q=tag:v0) [![Nightly](https://img.shields.io/github/v/release/ggml-org/llama.cpp?label=nightly&filter=b*&color=orange)](https://github.com/ggml-org/llama.cpp/releases?q=b) -[![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/server.yml?label=Server)](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml) +[![Server](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/server.yml?label=Server)](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml) [![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/docker.yml?label=Docker)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml) [![Winget](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/winget.yml?label=Winget)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml) From 86722900390abf479fd9719eda12c299a2b25bbf Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?= Date: Sat, 22 Aug 2026 09:09:26 +0200 Subject: [PATCH 08/16] sycl : add Q2_K reordered MMVQ and ESIMD kernels (again) (#27490) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Revert "Revert "sycl : add Q2_K reordered MMVQ and ESIMD kernels (#26336)" (#…" This reverts commit 7a0e42fd01fb0acda644e4f04b1f1acbbb9e23ba. * add gate params --- ggml/src/ggml-sycl/convert.cpp | 25 ++++++++- ggml/src/ggml-sycl/dequantize.hpp | 41 +++++++++++++++ ggml/src/ggml-sycl/dmmv.cpp | 27 +++++++++- ggml/src/ggml-sycl/esimd.hpp | 87 +++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/ggml-sycl.cpp | 2 + ggml/src/ggml-sycl/mmvq.cpp | 76 ++++++++++++++++++++++++++- ggml/src/ggml-sycl/quants.hpp | 23 ++++++++ ggml/src/ggml-sycl/vecdotq.hpp | 33 ++++++++++++ 8 files changed, 310 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index 9ec927695..b660b56ab 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -76,6 +76,19 @@ static void dequantize_row_q2_K_sycl(const void *vx, dst_t *y, const int64_t k, #endif } +template +static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int64_t k, + dpct::queue_ptr stream) { + const int64_t nb = k / QK_K; + + dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 }); + stream->parallel_for( + sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), + [=](sycl::nd_item<3> item_ct1) { + dequantize_block_q2_K_reorder(vx, y, item_ct1, nb); + }); +} + template static void dequantize_row_q3_K_sycl(const void *vx, dst_t *y, const int64_t k, dpct::queue_ptr stream) { @@ -667,7 +680,11 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) { return dequantize_block_sycl; } case GGML_TYPE_Q2_K: - return dequantize_row_q2_K_sycl; + if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { + return dequantize_row_q2_K_sycl_reorder; + } else { + return dequantize_row_q2_K_sycl; + } case GGML_TYPE_Q3_K: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { return dequantize_row_q3_K_sycl_reorder; @@ -753,7 +770,11 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) { return dequantize_block_sycl; } case GGML_TYPE_Q2_K: - return dequantize_row_q2_K_sycl; + if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { + return dequantize_row_q2_K_sycl_reorder; + } else { + return dequantize_row_q2_K_sycl; + } case GGML_TYPE_Q3_K: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { return dequantize_row_q3_K_sycl_reorder; diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index 876ba1b44..1b13e0f1a 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -943,6 +943,47 @@ static void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restri } +template +static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, + const sycl::nd_item<3> & item_ct1, int64_t n_blocks) { +#if QK_K == 256 + const int64_t i = item_ct1.get_group(2); + if (i >= n_blocks) { + return; + } + + const uint8_t * base = static_cast(vx); + const size_t qs_offset = i * (QK_K / 4); + const size_t scales_offset = n_blocks * (QK_K / 4) + i * (QK_K / 16); + const size_t dm_offset = n_blocks * (QK_K / 4) + n_blocks * (QK_K / 16) + i * sizeof(ggml_half2); + + const uint8_t * qs = base + qs_offset; + const uint8_t * scales = base + scales_offset; + const ggml_half2 * dm = reinterpret_cast(base + dm_offset); + + const int64_t tid = item_ct1.get_local_id(2); + const int64_t n = tid / 32; + const int64_t l = tid - 32 * n; + const int64_t is = 8 * n + l / 16; + + const uint8_t q = qs[32 * n + l]; + dst_t * y = yy + i * QK_K + 128 * n; + + const float dall = (*dm)[0]; + const float dmin = (*dm)[1]; + y[l+ 0] = dall * (scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (scales[is+0] >> 4); + y[l+32] = dall * (scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (scales[is+2] >> 4); + y[l+64] = dall * (scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (scales[is+4] >> 4); + y[l+96] = dall * (scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (scales[is+6] >> 4); +#else + GGML_UNUSED(vx); + GGML_UNUSED(yy); + GGML_UNUSED(item_ct1); + GGML_UNUSED(n_blocks); + GGML_ABORT("Q2_K reorder dequantize not supported for QK_K != 256"); +#endif +} + template static void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restrict__ yy, const sycl::nd_item<3> &item_ct1) { diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp index fdcadbf91..d47d6831a 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -1921,6 +1921,23 @@ ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd( } } +static void dequantize_mul_mat_vec_q2_K_sycl_reorder_esimd(const void *vx, const float *y, + float *dst, const int ncols, + const int nrows, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + const int workgroups = (nrows + 1) / 2; + stream->submit([&](sycl::handler &h) { + sycl::local_accessor lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h); + h.parallel_for( + sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)), + [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] { + dequantize_mul_mat_vec_reorder_esimd( + vx, y, dst, ncols, nrows, lmem, it); + }); + }); +} + static void dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(const void *vx, const float *y, float *dst, const int ncols, const int nrows, @@ -2111,7 +2128,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec( case GGML_TYPE_Q2_K: if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { - dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); +#ifdef GGML_SYCL_DMMV_HAS_ESIMD + if (g_ggml_sycl_enable_esimd) { + dequantize_mul_mat_vec_q2_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); + } + else +#endif + { + dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); + } } else { dequantize_mul_mat_vec_q2_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); } diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index 04e596ed3..0485ff0ce 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -61,6 +61,93 @@ static ESIMD_INLINE void unpack_scale_min_k4( min_f = convert(m) * (-dmin); } +// --------------------------------------------------------------------------- +// Q2_K, SOA reorder layout produced by reorder_qw_q2_k: +// [qs: nb*(QK_K/4)] [scales: nb*(QK_K/16)] [dm: nb*sizeof(half2)] +// with nb = nrows*num_blocks_per_row. +// +// 2 bits per weight. The 8 output chunks of 32 (matching dequantize_row_q2_K) +// map to super-chunk s (0..7): byte base 32*(s/4) into the 64-byte qs array, +// bit shift 2*(s%4); the low 16 lanes use scales[2s], the high 16 use +// scales[2s+1], with dl = d*(sc & 0xF), ml = dmin*(sc >> 4), deq = dl*q - ml. +// --------------------------------------------------------------------------- +template <> struct esimd_reorder_q_traits { + struct ptrs { + const uint8_t * qs; + const uint8_t * scales; + const sycl::half * dm; + }; + + static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) { + const uint8_t * qs = (const uint8_t *) vx; + const uint8_t * scales = qs + nb * (QK_K / 4); + const sycl::half * dm = (const sycl::half *) (scales + nb * (QK_K / 16)); + return { qs, scales, dm }; + } + + static ESIMD_INLINE void mac_pair( + const ptrs & pa, size_t bia, + const ptrs & pb, size_t bib, bool has_b, + sycl::ext::intel::esimd::simd & y_vec, + sycl::ext::intel::esimd::simd & acc_a, + sycl::ext::intel::esimd::simd & acc_b) { + using namespace sycl::ext::intel::esimd; + + simd qs_a = block_load(pa.qs + bia * (QK_K / 4)); + simd qs_b = 0; + simd scales_a = block_load(pa.scales + bia * (QK_K / 16)); + simd scales_b = 0; + + const float dall_a = (float) pa.dm[bia * 2 + 0]; + const float dmin_a = (float) pa.dm[bia * 2 + 1]; + float dall_b = 0.0f; + float dmin_b = 0.0f; + if (has_b) { + qs_b = block_load(pb.qs + bib * (QK_K / 4)); + scales_b = block_load(pb.scales + bib * (QK_K / 16)); + dall_b = (float) pb.dm[bib * 2 + 0]; + dmin_b = (float) pb.dm[bib * 2 + 1]; + } + + // per-chunk scale (d * (sc & 0xF)) and min (-dmin * (sc >> 4)), all 16 codes; + // min carries the negation so the dequant epilogue adds (matches Q4_K/Q5_K) + simd scale_f_a = convert(scales_a & simd(0x0F)) * dall_a; + simd min_f_a = convert(scales_a >> simd(4)) * (-dmin_a); + simd scale_f_b = convert(scales_b & simd(0x0F)) * dall_b; + simd min_f_b = convert(scales_b >> simd(4)) * (-dmin_b); + +#pragma unroll + for (int s = 0; s < 8; ++s) { + const int byte_base = 32 * (s / 4); + const uint8_t shift = (uint8_t) (2 * (s % 4)); + simd y_s = y_vec.select<32, 1>(s * 32); + + simd qa = (qs_a.select<32, 1>(byte_base) >> shift) & simd(3); + simd qb = (qs_b.select<32, 1>(byte_base) >> shift) & simd(3); + + const float scale_a_lo = scale_f_a[2 * s + 0]; + const float scale_a_hi = scale_f_a[2 * s + 1]; + const float min_a_lo = min_f_a[2 * s + 0]; + const float min_a_hi = min_f_a[2 * s + 1]; + const float scale_b_lo = scale_f_b[2 * s + 0]; + const float scale_b_hi = scale_f_b[2 * s + 1]; + const float min_b_lo = min_f_b[2 * s + 0]; + const float min_b_hi = min_f_b[2 * s + 1]; + + simd scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi); + simd min_vec_a = splat_lo_hi(min_a_lo, min_a_hi); + simd scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi); + simd min_vec_b = splat_lo_hi(min_b_lo, min_b_hi); + + simd deq_a = convert(qa) * scale_vec_a + min_vec_a; + simd deq_b = convert(qb) * scale_vec_b + min_vec_b; + + acc_a += y_s * deq_a; + acc_b += y_s * deq_b; + } + } +}; + // --------------------------------------------------------------------------- // Q3_K, SOA reorder layout produced by reorder_qw_q3_k: // [qs: nb*(QK_K/4)] [hmask: nb*(QK_K/8)] [scales: nb*12] [d: nb*sizeof(half)] diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index de56ea5b9..3f82020f4 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3796,6 +3796,7 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) { case GGML_TYPE_Q1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: @@ -3809,6 +3810,7 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) { static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) { #ifdef GGML_SYCL_DMMV_HAS_ESIMD switch (type) { + case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index 123b2a2f0..f6a8591e3 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1401,6 +1401,65 @@ static void mul_mat_vec_q2_K_q8_1_sycl_switch_ncols( } } +static void reorder_mul_mat_vec_q2_k_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, + const int nrows, dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + + // Round up to a whole number of subgroup-sized workgroups; out-of-range rows are skipped inside the kernel. + constexpr size_t num_subgroups = WARP_SIZE; + const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups); + const sycl::range<3> block_nums(1, 1, block_num_y); + const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + mul_mat_vec_q_reorder>(vx, vy, dst, ncols, nrows, + nd_item); + }); + }); +} + +template +static void reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols( + const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, + const int stride_col_y_bytes, const int stride_col_dst, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + constexpr size_t num_subgroups = WARP_SIZE; + const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups); + const sycl::range<3> block_nums(1, 1, block_num_y); + const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); + + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + mul_mat_vec_q_reorder_ncols, ncols_dst>( + vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, + /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item); + }); + }); +} + +static void reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols( + const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, const int ncols_dst, + const int stride_col_y_bytes, const int stride_col_dst, + dpct::queue_ptr stream) { + switch (ncols_dst) { + case 1: reorder_mul_mat_vec_q2_k_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); break; + case 2: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 3: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 4: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 5: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 6: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<6>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 7: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<7>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + case 8: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<8>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; + default: GGML_ABORT("unsupported ncols_dst=%d for Q2_K reorder multi-col MMVQ", ncols_dst); + } +} + static void mul_mat_vec_q3_K_q8_1_sycl(const void *vx, const void *vy, float *dst, const int ncols, const int nrows, @@ -2297,7 +2356,21 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens } break; case GGML_TYPE_Q2_K: - if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { + if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && + ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { + if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { + const int stride_col_y_bytes = src1_padded_col_size * q8_1_ts / q8_1_bs; + const int stride_col_dst = dst->ne[0]; + GGML_SYCL_DEBUG("Calling reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); + reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols( + src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, + src1_ncols, stride_col_y_bytes, stride_col_dst, stream); + return; + } else { + GGML_SYCL_DEBUG("Calling reorder_mul_mat_vec_q2_k_q8_1_sycl\n"); + reorder_mul_mat_vec_q2_k_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); + } + } else if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { const int stride_col_y = src1_padded_col_size / QK8_1; const int stride_col_dst = dst->ne[0]; GGML_SYCL_DEBUG("Calling mul_mat_vec_q2_K_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); @@ -2306,6 +2379,7 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens src1_ncols, stride_col_y, stride_col_dst, stream); return; } else if (i == 0 || src1_ncols == 1) { + GGML_SYCL_DEBUG("Calling mul_mat_vec_q2_K_q8_1_sycl\n"); mul_mat_vec_q2_K_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); } break; diff --git a/ggml/src/ggml-sycl/quants.hpp b/ggml/src/ggml-sycl/quants.hpp index 95287f175..a26a6ce6e 100644 --- a/ggml/src/ggml-sycl/quants.hpp +++ b/ggml/src/ggml-sycl/quants.hpp @@ -58,6 +58,29 @@ template <> struct block_q_t { static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; } }; +template <> struct block_q_t { + struct traits { + static constexpr uint32_t qk = QK_K; + static constexpr uint32_t qi = QI2_K; + static constexpr uint32_t qr = QR2_K; + static constexpr uint32_t vdr_mmvq = 1; + }; + + // Reordered layout: [qs (QK_K/4 per block)] [scales (QK_K/16 per block)] [dm] + static constexpr std::pair get_block_offset(const int block_index, const int /* n_blocks */) { + return { block_index * (QK_K / 4), 0 }; + } + + static constexpr std::pair get_d_offset(int nrows, int ncols, const int block_index) { + auto nblocks = (nrows * (ncols / QK_K)); + auto total_qs_bytes = nblocks * (QK_K / 4); + return { total_qs_bytes + block_index * (QK_K / 16), + total_qs_bytes + nblocks * (QK_K / 16) + block_index * sizeof(ggml_half2) }; + } + + static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; } +}; + template <> struct block_q_t { struct traits { static constexpr uint32_t qk = QK_K; diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index c11a6e8f9..3ad4cee93 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -429,6 +429,39 @@ template <> struct reorder_vec_dot_q_sycl { } }; +template <> struct reorder_vec_dot_q_sycl { + static constexpr ggml_type gtype = GGML_TYPE_Q2_K; + + using q2_k_block = ggml_sycl_reordered::block_q_t; + using q2_k_traits = typename q2_k_block::traits; + + __dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair ibx_offset, + const std::pair d_offset, const int8_t * q8_1_quant_ptr, + const sycl::half2 * q8_1_ds, const int & iqs) { + const uint8_t * base = static_cast(vbq); + const uint8_t * qs = base + ibx_offset.first; + const uint8_t * scales = base + d_offset.first; + const ggml_half2 * dm = reinterpret_cast(base + d_offset.second); + + const int bq8_offset = QR2_K * (iqs / QI8_1); + const int scale_offset = iqs - iqs % QI8_1 + (iqs % QI8_1) / (QI8_1 / 2); + + const int v = get_int_from_uint8_aligned(qs, iqs); + + int u[QR2_K]; + float d8[QR2_K]; + +#pragma unroll + for (int i = 0; i < QR2_K; ++i) { + const int8_t * quant_base_ptr = q8_1_quant_ptr + (bq8_offset + i) * QK8_1; + u[i] = get_int_from_int8_aligned(quant_base_ptr, iqs % QI8_1); + d8[i] = (*(q8_1_ds + bq8_offset + i))[0]; + } + + return vec_dot_q2_K_q8_1_impl_mmvq(v, u, scales + scale_offset, *dm, d8); + } +}; + template <> struct reorder_vec_dot_q_sycl { static constexpr ggml_type gtype = GGML_TYPE_Q3_K; From 2c6b141efb3b0868fd39d3cae73f69606e1d654c Mon Sep 17 00:00:00 2001 From: Shahir BIn Zulfiker <119410932+aorko01@users.noreply.github.com> Date: Sat, 22 Aug 2026 13:44:22 +0600 Subject: [PATCH 09/16] common : fix draft-mtp with embeddings (#26352, #27299) (#27400) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * common: fix draft-mtp with embeddings (#26352) * --whitespace --------- Co-authored-by: Sigbjørn Skjæret --- common/speculative.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/common/speculative.cpp b/common/speculative.cpp index 89e9b2782..8461e4107 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -2322,6 +2322,9 @@ common_params common_base_params_to_speculative(const common_params & params) { const auto & params_spec = params.speculative.draft; common_params result = params; + result.embedding = false; + result.pooling_type = LLAMA_POOLING_TYPE_UNSPECIFIED; + if (has_draft) { result.devices = params_spec.devices; result.model = params_spec.mparams; From 369e1cd6140b7cbdb552cc3f87613aeb9e122422 Mon Sep 17 00:00:00 2001 From: Kartik Sirohi <99896785+sirohikartik@users.noreply.github.com> Date: Sat, 22 Aug 2026 14:00:31 +0530 Subject: [PATCH 10/16] ggml: optimize concat op by replacing per-element memcpy with row-level memcpy (#24575) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * ggml: optimize concat op by replacing per-element memcpy with row-level memcpy * ggml: fix concat offsets for row-level copies * ggml: add concat row contiguity asserts * ggml: move concat block size asserts * ggml: remove redundant concat asserts * Update ggml/src/ggml-cpu/ops.cpp Co-authored-by: Sigbjørn Skjæret --------- Co-authored-by: Sigbjørn Skjæret --- ggml/src/ggml-cpu/ops.cpp | 50 ++++++++++++++++++--------------------- 1 file changed, 23 insertions(+), 27 deletions(-) diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 2b5f68444..b869f4bdd 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -1896,7 +1896,6 @@ void ggml_compute_forward_repeat_back( } // ggml_compute_forward_concat - static void ggml_compute_forward_concat_any( const ggml_compute_params * params, ggml_tensor * dst) { @@ -1904,8 +1903,6 @@ static void ggml_compute_forward_concat_any( const ggml_tensor * src0 = dst->src[0]; const ggml_tensor * src1 = dst->src[1]; - const size_t len = ggml_type_size(src0->type); - const int ith = params->ith; const int nth = params->nth; @@ -1914,31 +1911,38 @@ static void ggml_compute_forward_concat_any( const int32_t dim = ggml_get_op_params_i32(dst, 0); GGML_ASSERT(dim >= 0 && dim < 4); + GGML_ASSERT(ggml_is_contiguous_rows(src0)); + GGML_ASSERT(ggml_is_contiguous_rows(src1)); int64_t o[4] = {0, 0, 0, 0}; + if (dim == 0) { + GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); + GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); + o[dim] = src0->ne[dim]/ggml_blck_size(src0->type); } else { o[dim] = src0->ne[dim]; } - const char * x; + // Region 1: copy rows from src0 + for (int i3 = 0; i3 < ne03; i3++) { + for (int i2 = ith; i2 < ne02; i2 += nth) { + for (int i1 = 0; i1 < ne01; i1++) { + const char * x = (const char *) src0->data + i1*nb01 + i2*nb02 + i3*nb03; + char * y = ( char *) dst->data + i1*nb1 + i2*nb2 + i3*nb3; + memcpy(y, x, ggml_row_size(src0->type, ne00)); + } + } + } - // TODO: smarter multi-theading - for (int i3 = 0; i3 < ne3; i3++) { - for (int i2 = ith; i2 < ne2; i2 += nth) { - for (int i1 = 0; i1 < ne1; i1++) { - for (int i0 = 0; i0 < ne0/ggml_blck_size(dst->type); i0++) { - if (i0 < ne00/ggml_blck_size(src0->type) && i1 < ne01 && i2 < ne02 && i3 < ne03) { - x = (const char *)src0->data + (i0 )*nb00 + (i1 )*nb01 + (i2 )*nb02 + (i3 )*nb03; - } else { - x = (const char *)src1->data + (i0 - o[0])*nb10 + (i1 - o[1])*nb11 + (i2 - o[2])*nb12 + (i3 - o[3])*nb13; - } - - char * y = (char *)dst->data + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3; - - memcpy(y, x, len); - } + // Region 2: copy rows from src1, offset into dst by o[] + for (int i3 = 0; i3 < ne13; i3++) { + for (int i2 = ith; i2 < ne12; i2 += nth) { + for (int i1 = 0; i1 < ne11; i1++) { + const char * x = (const char *) src1->data + i1*nb11 + i2*nb12 + i3*nb13; + char * y = ( char *) dst->data + (i1 + o[1])*nb1 + (i2 + o[2])*nb2 + (i3 + o[3])*nb3 + o[0]*nb0; + memcpy(y, x, ggml_row_size(src1->type, ne10)); } } } @@ -2078,14 +2082,6 @@ void ggml_compute_forward_concat( ggml_tensor * dst) { const ggml_tensor * src0 = dst->src[0]; - const ggml_tensor * src1 = dst->src[1]; - - if (ggml_is_quantized(src0->type)) { - GGML_ASSERT(ggml_is_contiguous_rows(src0)); - GGML_ASSERT(ggml_is_contiguous_rows(src1)); - GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); - GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); - } switch (src0->type) { case GGML_TYPE_F16: From 3a653fea932e6a026e14cc0fb6507a986679f688 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Sat, 22 Aug 2026 11:31:30 +0300 Subject: [PATCH 11/16] ci : add older, min and dry-run options to ccache-clear (#27504) * ci : add older, min and dry-run options to ccache-clear Assisted-by: pi:llama.cpp/Qwen3.8-27B * pi : add note about not wrapping lines in PR descriptions [no ci] Assisted-by: pi:llama.cpp/Qwen3.8-27B --- .github/actions/ccache-clear/action.yml | 77 +++++++++++++++++++++++-- .github/workflows/build-cpu.yml | 12 ++++ .pi/gg/SYSTEM.md | 1 + 3 files changed, 84 insertions(+), 6 deletions(-) diff --git a/.github/actions/ccache-clear/action.yml b/.github/actions/ccache-clear/action.yml index d20f6bf84..2420045ed 100644 --- a/.github/actions/ccache-clear/action.yml +++ b/.github/actions/ccache-clear/action.yml @@ -1,23 +1,88 @@ # note: place this as the last step of the job, so the new cache is saved by "Post ccache" right after the old one is cleared name: "ccache-clear" -description: "Delete all GitHub Actions caches matching a key prefix" +description: "Delete GitHub Actions caches matching a key prefix, oldest first" inputs: key: description: "Cache key prefix to match and delete" required: true + older: + description: "Only delete caches created more than this long ago (e.g. 90m, 1h, 1d). By default all matching caches are deleted" + required: false + default: "" + min: + description: "Stop deleting if fewer than this many caches would remain (e.g. 1). By default there is no minimum" + required: false + default: "0" + dry-run: + description: "Only print the caches that would be deleted, without deleting them" + required: false + default: "false" runs: using: "composite" steps: - name: Clear caches shell: bash + env: + CLEAR_KEY: ${{ inputs.key }} + CLEAR_OLDER: ${{ inputs.older }} + CLEAR_MIN: ${{ inputs.min }} + CLEAR_DRY_RUN: ${{ inputs.dry-run }} run: | - CACHES=$(gh cache list --key "ccache-${{ inputs.key }}" --json id,key --jq '.[] | "\(.id) \(.key)"' 2>/dev/null) + # Convert a duration (e.g. 90m, 1h, 1d, plain seconds) to seconds + to_seconds() { + local val="$1" + [[ "$val" =~ ^[0-9]+$ ]] && { echo "$val"; return 0; } + local num="${val%?}" unit="${val: -1}" mult + [[ "$num" =~ ^[0-9]+$ ]] || return 1 + case "$unit" in + s) mult=1 ;; + m) mult=60 ;; + h) mult=3600 ;; + d) mult=86400 ;; + *) return 1 ;; + esac + echo $((num * mult)) + } + + [[ "$CLEAR_MIN" =~ ^[0-9]+$ ]] || { echo "Invalid min value: $CLEAR_MIN" >&2; exit 1; } + [[ "$CLEAR_DRY_RUN" =~ ^(true|false)$ ]] || { echo "Invalid dry-run value: $CLEAR_DRY_RUN" >&2; exit 1; } + + CACHES=$(gh cache list --key "ccache-$CLEAR_KEY" --json id,key,createdAt --jq '.[] | [.createdAt, .id, .key] | @tsv' 2>/dev/null | LC_ALL=C sort) if [ -z "$CACHES" ]; then - echo "No caches found with key prefix: ${{ inputs.key }}" + echo "No caches found with key prefix: $CLEAR_KEY" exit 0 fi - while read -r id key; do - echo "Deleting cache: $id ($key)" - gh cache delete "$id" + + TOTAL=$(( $(wc -l <<< "$CACHES") )) + + echo "Found $TOTAL cache(s) with key prefix: $CLEAR_KEY (oldest first):" + while IFS=$'\t' read -r CREATED ID KEY; do + printf ' %s %s %s\n' "$CREATED" "$ID" "$KEY" + done <<< "$CACHES" + + CUTOFF="" + if [ -n "$CLEAR_OLDER" ]; then + OLDER_SECONDS=$(to_seconds "$CLEAR_OLDER") || { echo "Invalid older value: $CLEAR_OLDER (expected e.g. 90m, 1h, 1d)" >&2; exit 1; } + CUTOFF=$(( $(date +%s) - OLDER_SECONDS )) + fi + + # Caches are sorted oldest first + DELETED=0 + while IFS=$'\t' read -r CREATED ID KEY; do + if [ -n "$CUTOFF" ] && [ "$(date -d "$CREATED" +%s)" -ge "$CUTOFF" ]; then + echo "Rest are not older than $CLEAR_OLDER, stopping" + break + fi + if [ $((TOTAL - DELETED - 1)) -lt "$CLEAR_MIN" ]; then + echo "Keeping at least $CLEAR_MIN cache(s), stopping" + break + fi + if [ "$CLEAR_DRY_RUN" = "true" ]; then + echo "Would delete cache: $ID ($KEY)" + else + echo "Deleting cache: $ID ($KEY)" + gh cache delete "$ID" + fi + DELETED=$((DELETED + 1)) done <<< "$CACHES" diff --git a/.github/workflows/build-cpu.yml b/.github/workflows/build-cpu.yml index ffced86c6..f39304cab 100644 --- a/.github/workflows/build-cpu.yml +++ b/.github/workflows/build-cpu.yml @@ -117,6 +117,18 @@ jobs: ./bin/llama-convert-llama2c-to-ggml --copy-vocab-from-model ./tok512.bin --llama2c-model stories260K.bin --llama2c-output-model stories260K.gguf ./bin/llama-completion -m stories260K.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256 + # note: real deletion only on push to master (same condition as the ccache save), + # dry-run otherwise (the token is read-only on PRs from forks) + - name: ccache-clear + uses: ./.github/actions/ccache-clear + env: + GH_TOKEN: ${{ github.token }} + with: + key: cpu-${{ matrix.os }} + older: 1h + min: 1 + dry-run: ${{ github.event_name != 'push' || github.ref != 'refs/heads/master' }} + windows: name: windows / ${{ matrix.build }} runs-on: windows-2025 diff --git a/.pi/gg/SYSTEM.md b/.pi/gg/SYSTEM.md index 6a757c869..47883081c 100644 --- a/.pi/gg/SYSTEM.md +++ b/.pi/gg/SYSTEM.md @@ -17,6 +17,7 @@ Coding: Pull requests (PRs): - New branch names are prefixed with "gg/" - Before opening a pull request, ask the user to confirm the description +- Don't explicitly wrap lines in the PR description (each paragraph and bullet is a single line) - When creating a pull request, look for the repository's PR template and follow it - For the AI usage disclosure section, write "YES. pi:llama.cpp/[MODEL]" - Ask the user to tell you what model was used and write it in place of [MODEL] From 54ee5ee643f29abba6852903ddfdb688c2361b5b Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Sat, 22 Aug 2026 10:35:50 +0200 Subject: [PATCH 12/16] mtmd: support dots3-note vision+audio (#27524) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * text: conversion * init impl * mtmd: conversion * impl mtmd cpp * Update gguf-py/gguf/tensor_mapping.py Co-authored-by: Sigbjørn Skjæret --------- Co-authored-by: Sigbjørn Skjæret --- conversion/__init__.py | 2 + conversion/dots3.py | 132 +++++++++++++++++++++++++++++++- gguf-py/gguf/constants.py | 19 +++++ gguf-py/gguf/gguf_writer.py | 6 ++ gguf-py/gguf/tensor_mapping.py | 55 ++++++++++++- tools/mtmd/CMakeLists.txt | 1 + tools/mtmd/clip-graph.h | 6 ++ tools/mtmd/clip-impl.h | 11 ++- tools/mtmd/clip-model.h | 8 ++ tools/mtmd/clip.cpp | 129 +++++++++++++++++++++++++++++-- tools/mtmd/models/dots3note.cpp | 61 +++++++++++++++ tools/mtmd/models/models.h | 5 ++ tools/mtmd/mtmd-audio.cpp | 94 +++++++++++++++++++++++ tools/mtmd/mtmd-audio.h | 9 +++ tools/mtmd/mtmd.cpp | 8 ++ 15 files changed, 535 insertions(+), 11 deletions(-) create mode 100644 tools/mtmd/models/dots3note.cpp diff --git a/conversion/__init__.py b/conversion/__init__.py index b4afdf0f8..aaaebba97 100644 --- a/conversion/__init__.py +++ b/conversion/__init__.py @@ -282,6 +282,8 @@ MMPROJ_MODEL_MAP: dict[str, str] = { "CogVLMForCausalLM": "cogvlm", "DeepseekOCR2ForCausalLM": "deepseek", "DeepseekOCRForCausalLM": "deepseek", + "Dots3NoteForCausalLM": "dots3", + "Dots3NoteForConditionalGeneration": "dots3", "DotsOCRForCausalLM": "dotsocr", "Exaone4_5_ForConditionalGeneration": "exaone", "Gemma3ForConditionalGeneration": "gemma", diff --git a/conversion/dots3.py b/conversion/dots3.py index 7c36b8482..c7ac2319e 100644 --- a/conversion/dots3.py +++ b/conversion/dots3.py @@ -3,12 +3,14 @@ from __future__ import annotations import math import re -from typing import TYPE_CHECKING, Callable, Iterable +import torch + +from typing import TYPE_CHECKING, Any, Callable, Iterable if TYPE_CHECKING: from torch import Tensor -from .base import ModelBase, gguf +from .base import MmprojModel, ModelBase, gguf from .deepseek import DeepseekV2Model @@ -193,3 +195,129 @@ class Dots3NoteModel(DeepseekV2Model): return yield from super().modify_tensors(data_torch, name, bid) + + +@ModelBase.register("Dots3NoteForCausalLM", "Dots3NoteForConditionalGeneration") +class Dots3NoteMmprojModel(MmprojModel): + has_vision_encoder = True + has_audio_encoder = True + + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + assert self.hparams_vision is not None + assert self.hparams_audio is not None + + # preprocessor_config.json nests the image params under vision_config + self.preprocessor_config = {**self.preprocessor_config, **self.preprocessor_config.get("vision_config", {})} + + vis = self.hparams_vision + # in this config, hidden_size is the adapter output width; embed_dim is the tower width + vis["hidden_size"] = vis["embed_dim"] + vis["image_size"] = 0 # dynamic resolution + self.pyramid = [max(0, n) for n in vis["pyramid_num_routed"]] + + if vis.get("adapter_type") != "patch_merger" or not vis.get("pre_pixel_shuffle"): + raise ValueError("dots3-note vision conversion requires adapter_type=patch_merger and pre_pixel_shuffle") + if vis.get("router_scoring_func", "sigmoid") != "sigmoid" or vis.get("router_scale", 1.0) != 1.0: + raise ValueError("dots3-note vision conversion only supports sigmoid routing with router_scale=1.0") + if vis.get("temporal_patch_size", 1) != 1 or vis.get("use_bias") or not vis.get("use_qk_norm"): + raise ValueError("unsupported dots3-note vision config variant") + + aud = self.hparams_audio + if not aud.get("use_conv2d_stem") or not aud.get("use_rope") or not aud.get("use_rms_norm") or aud.get("use_causal"): + raise ValueError("unsupported dots3-note audio config variant") + if aud["whisper_config"].get("activation_function") != "swiglu": + raise ValueError("dots3-note audio conversion requires the swiglu activation") + if aud.get("merge_factor", 1) != 1 or aud.get("chunk_seconds") != 60: + raise ValueError("unsupported dots3-note audio chunking config") + # the graph hard-codes these rope parameters + rope = aud.get("rope_parameters", {}) + if rope.get("partial_rotary_factor") != 0.5 or rope.get("rope_theta") != 10000.0: + raise ValueError("unsupported dots3-note audio rope config") + + def get_audio_config(self) -> dict[str, Any] | None: + cfg = self.global_config.get("audio_config") + if cfg is not None: + # aliases so MmprojModel.find_aparam() / n_block_keys can resolve them + whisper = cfg["whisper_config"] + cfg["hidden_size"] = whisper["d_model"] + cfg["intermediate_size"] = whisper["encoder_ffn_dim"] + cfg["num_attention_heads"] = whisper["encoder_attention_heads"] + cfg["num_hidden_layers"] = whisper["encoder_layers"] + return cfg + + def set_gguf_parameters(self): + super().set_gguf_parameters() + assert self.hparams_vision is not None + assert self.hparams_audio is not None + + self.gguf_writer.add_clip_vision_projector_type(gguf.VisionProjectorType.DOTS3NOTE_V) + self.gguf_writer.add_vision_use_silu(True) + self.gguf_writer.add_vision_attention_layernorm_eps(self.hparams_vision["rms_norm_eps"]) + self.gguf_writer.add_vision_spatial_merge_size(self.hparams_vision["spatial_merge_size"]) + self.gguf_writer.add_vision_min_pixels(self.preprocessor_config["min_pixels"]) + self.gguf_writer.add_vision_max_pixels(self.preprocessor_config["max_pixels"]) + # pyramid MoE: per-block routed expert count, 0 = dense block + self.gguf_writer.add_vision_expert_count_per_layer(self.pyramid) + self.gguf_writer.add_vision_expert_used_count(int(self.hparams_vision["capacity_factor"])) + + self.gguf_writer.add_clip_audio_projector_type(gguf.VisionProjectorType.DOTS3NOTE_A) + self.gguf_writer.add_audio_num_mel_bins(self.hparams_audio["whisper_config"]["num_mel_bins"]) + self.gguf_writer.add_audio_attention_layernorm_eps(1e-6) # Dots3NoteAudioRMSNorm default + + @classmethod + def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None: + name, _ = item + if not name.startswith(("vision_encoder.", "audio_encoder.")): + return None + return super().filter_tensors(item) + + _vis_experts: dict[int, dict[str, Tensor]] | None = None + + def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: + # router params have no .weight suffix in the checkpoint, but gguf tools expect one + if name.endswith((".gate_weight", ".router_bias")): + name += ".weight" + + # audio fc1 fuses gate and up for swiglu; split it + if ".speech_encoder.layers." in name and ".fc1." in name: + gate, up = data_torch.chunk(2, dim=0) + yield from super().modify_tensors(gate, name.replace(".fc1.", ".fc1_gate."), bid) + yield from super().modify_tensors(up, name.replace(".fc1.", ".fc1_up."), bid) + return + + # vision MoE: stack per-expert weights into a single 3D tensor per block + if ".mlp.experts." in name: + assert bid is not None + n_expert = self.pyramid[bid] + if self._vis_experts is None: + self._vis_experts = {} + buf = self._vis_experts.setdefault(bid, {}) + buf[name] = data_torch + + if len(buf) >= n_expert * 3: + for w_name in ("fc1", "fc2", "fc3"): + datas: list[Tensor] = [] + for xid in range(n_expert): + ename = f"vision_encoder.blocks.{bid}.mlp.experts.{xid}.{w_name}.weight" + datas.append(buf.pop(ename)) + merged = torch.stack(datas, dim=0) + yield from super().modify_tensors(merged, f"vision_encoder.blocks.{bid}.mlp.experts.{w_name}.weight", bid) + return + + yield from super().modify_tensors(data_torch, name, bid) + + def prepare_tensors(self): + super().prepare_tensors() + if self._vis_experts is not None: + leftover = [k for d in self._vis_experts.values() for k in d.keys()] + if leftover: + raise ValueError(f"unprocessed vision experts: {leftover}") + + def tensor_force_quant(self, name, new_name, bid, n_dims): + # FP32 routing is load-bearing for the vision MoE (near-tied expert scores) + if ".ffn_gate_inp." in new_name or ".exp_probs_b." in new_name: + return gguf.GGMLQuantizationType.F32 + if ".conv2d" in new_name or "a.conv_out" in new_name: + return gguf.GGMLQuantizationType.F32 + return super().tensor_force_quant(name, new_name, bid, n_dims) diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 886253d88..8f6f55519 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -364,6 +364,8 @@ class Keys: IMAGE_MEAN = "clip.vision.image_mean" IMAGE_STD = "clip.vision.image_std" SPATIAL_MERGE_SIZE = "clip.vision.spatial_merge_size" + EXPERT_COUNT_PER_LAYER = "clip.vision.expert_count_per_layer" # dots3note pyramid MoE, 0 = dense layer + EXPERT_USED_COUNT = "clip.vision.expert_used_count" USE_GELU = "clip.use_gelu" USE_SILU = "clip.use_silu" N_WA_PATTERN = "clip.vision.n_wa_pattern" # used by qwen2.5vl @@ -874,6 +876,11 @@ class MODEL_TENSOR(IntEnum): V_ENC_FFN_UP = auto() V_ENC_FFN_GATE = auto() V_ENC_FFN_DOWN = auto() + V_ENC_FFN_GATE_INP = auto() # dots3note vision MoE router + V_ENC_FFN_GATE_EXPS = auto() + V_ENC_FFN_UP_EXPS = auto() + V_ENC_FFN_DOWN_EXPS = auto() + V_ENC_FFN_EXP_PROBS_B = auto() V_ENC_ATTN_POST_NORM = auto() # gemma4 V_ENC_FFN_POST_NORM = auto() V_LAYER_SCALE_1 = auto() @@ -1591,6 +1598,11 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = { MODEL_TENSOR.V_ENC_FFN_UP: "v.blk.{bid}.ffn_up", MODEL_TENSOR.V_ENC_FFN_GATE: "v.blk.{bid}.ffn_gate", MODEL_TENSOR.V_ENC_FFN_DOWN: "v.blk.{bid}.ffn_down", + MODEL_TENSOR.V_ENC_FFN_GATE_INP: "v.blk.{bid}.ffn_gate_inp", + MODEL_TENSOR.V_ENC_FFN_GATE_EXPS: "v.blk.{bid}.ffn_gate_exps", + MODEL_TENSOR.V_ENC_FFN_UP_EXPS: "v.blk.{bid}.ffn_up_exps", + MODEL_TENSOR.V_ENC_FFN_DOWN_EXPS: "v.blk.{bid}.ffn_down_exps", + MODEL_TENSOR.V_ENC_FFN_EXP_PROBS_B: "v.blk.{bid}.exp_probs_b", MODEL_TENSOR.V_ENC_ATTN_POST_NORM: "v.blk.{bid}.attn_post_norm", MODEL_TENSOR.V_ENC_FFN_POST_NORM: "v.blk.{bid}.ffn_post_norm", MODEL_TENSOR.V_LAYER_SCALE_1: "v.blk.{bid}.ls1", @@ -1913,6 +1925,11 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = { MODEL_TENSOR.V_ENC_FFN_UP, MODEL_TENSOR.V_ENC_FFN_GATE, MODEL_TENSOR.V_ENC_FFN_DOWN, + MODEL_TENSOR.V_ENC_FFN_GATE_INP, + MODEL_TENSOR.V_ENC_FFN_GATE_EXPS, + MODEL_TENSOR.V_ENC_FFN_UP_EXPS, + MODEL_TENSOR.V_ENC_FFN_DOWN_EXPS, + MODEL_TENSOR.V_ENC_FFN_EXP_PROBS_B, MODEL_TENSOR.V_ENC_ATTN_POST_NORM, MODEL_TENSOR.V_ENC_FFN_POST_NORM, MODEL_TENSOR.V_LAYER_SCALE_1, @@ -5497,6 +5514,8 @@ class VisionProjectorType: COGVLM = "cogvlm" JANUS_PRO = "janus_pro" DOTSOCR = "dots_ocr" + DOTS3NOTE_V = "dots3note_v" + DOTS3NOTE_A = "dots3note_a" # audio DEEPSEEKOCR = "deepseekocr" DEEPSEEKOCR2 = "deepseekocr2" LFM2A = "lfm2a" # audio diff --git a/gguf-py/gguf/gguf_writer.py b/gguf-py/gguf/gguf_writer.py index 496882025..d8a96a27b 100644 --- a/gguf-py/gguf/gguf_writer.py +++ b/gguf-py/gguf/gguf_writer.py @@ -1327,6 +1327,12 @@ class GGUFWriter: def add_vision_spatial_merge_size(self, value: int) -> None: self.add_uint32(Keys.ClipVision.SPATIAL_MERGE_SIZE, value) + def add_vision_expert_count_per_layer(self, value: Sequence[int]) -> None: + self.add_array(Keys.ClipVision.EXPERT_COUNT_PER_LAYER, value) + + def add_vision_expert_used_count(self, value: int) -> None: + self.add_uint32(Keys.ClipVision.EXPERT_USED_COUNT, value) + def add_vision_use_gelu(self, value: bool) -> None: self.add_bool(Keys.ClipVision.USE_GELU, value) diff --git a/gguf-py/gguf/tensor_mapping.py b/gguf-py/gguf/tensor_mapping.py index 1ff4b61d9..ef580518e 100644 --- a/gguf-py/gguf/tensor_mapping.py +++ b/gguf-py/gguf/tensor_mapping.py @@ -1454,6 +1454,7 @@ class TensorNameMap: "mlp_AR.linear_{bid}", # PaddleOCR-VL "merger.mlp.{bid}", "vision_tower.merger.mlp.{bid}", # dots.ocr + "vision_encoder.adapter.mlp.{bid}", # dots3note "vit.perceive.proj.{bid}", # HunyuanVL (proj.0 = conv1, proj.2 = conv2) ), @@ -1504,6 +1505,7 @@ class TensorNameMap: "vision_model.radio_model.model.patch_generator.embedder", # Nemotron Nano v2 VL "model.vision_tower.patch_embedder.input_proj", # gemma4 "vision_tower.patch_embed.patchifier.proj", # dots.ocr + "vision_encoder.patch_embed.proj", # dots3note "vision_model.conv1", # Step3-VL "model.vision_embedder.patch_dense", # gemma4 unified "model.vision_tower.patch_embedder.patch_embedding", # muse-glimmer @@ -1512,6 +1514,7 @@ class TensorNameMap: MODEL_TENSOR.V_ENC_EMBD_NORM: ( "visual.post_conv_layernorm", # glm4v "vision_tower.patch_embed.patchifier.norm", # dots.ocr + "vision_encoder.patch_embed.norm", # dots3note ), MODEL_TENSOR.V_ENC_EMBD_PATCH_NORM: ( @@ -1551,6 +1554,7 @@ class TensorNameMap: MODEL_TENSOR.V_ENC_ATTN_QKV: ( "visual.blocks.{bid}.attn.qkv", # qwen3vl "vision_tower.blocks.{bid}.attn.qkv", # dots.ocr + "vision_encoder.blocks.{bid}.attn.qkv", # dots3note "model.vision.transformer.layers.{bid}.attention.query_key_value", # cogvlm "model.vision_model.transformer.layers.{bid}.self_attn.qkv_proj", # Deepseek-OCR CLIP "vision_tower.encoder.blocks.{bid}.wqkv", # Kimi-K2.5 @@ -1579,6 +1583,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_ATTN_Q_NORM: ( + "vision_encoder.blocks.{bid}.attn.q_norm", # dots3note "vision_tower.vision_model.encoder.layers.{bid}.attn.q_norm", # InternVL "model.vision_tower.encoder.layer.{bid}.attention.q_norm", # Intern-S1 "visual.blocks.{bid}.attn.q_norm", # GLM-OCR @@ -1606,6 +1611,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_ATTN_K_NORM: ( + "vision_encoder.blocks.{bid}.attn.k_norm", # dots3note "vision_tower.vision_model.encoder.layers.{bid}.attn.k_norm", # InternVL "model.vision_tower.encoder.layer.{bid}.attention.k_norm", # Intern-S1 "visual.blocks.{bid}.attn.k_norm", # GLM-OCR @@ -1651,6 +1657,7 @@ class TensorNameMap: "siglip2.vision_model.encoder.layers.{bid}.layer_norm1", "vision_model.radio_model.model.blocks.{bid}.norm1", # Nemotron Nano v2 VL "vision_tower.blocks.{bid}.norm1", # dots.ocr + "vision_encoder.blocks.{bid}.norm_1", # dots3note "vision_model.transformer.resblocks.{bid}.ln_1", # Step3-VL "model.qwen2_model.model.model.layers.{bid}.input_layernorm", # Deepseek-OCR-2 qwen2 "model.vision_tower.layers.{bid}.norm1", # muse-glimmer @@ -1678,6 +1685,7 @@ class TensorNameMap: "model.qwen2_model.model.model.layers.{bid}.self_attn.o_proj", # Deepseek-OCR-2 qwen2 "vision_model.model.layers.{bid}.self_attn.o_proj.linear", # gemma4 "vision_tower.blocks.{bid}.attn.proj", # dots.ocr + "vision_encoder.blocks.{bid}.attn.proj", # dots3note "vision_model.transformer.resblocks.{bid}.attn.out_proj", # Step3-VL "model.vision_tower.layers.{bid}.attn.proj", # muse-glimmer ), @@ -1706,12 +1714,14 @@ class TensorNameMap: "vision_model.radio_model.model.blocks.{bid}.norm2", # Nemotron Nano v2 VL "vision_model.model.layers.{bid}.pre_feedforward_layernorm", # gemma4 "vision_tower.blocks.{bid}.norm2", # dots.ocr + "vision_encoder.blocks.{bid}.norm_2", # dots3note "vision_model.transformer.resblocks.{bid}.ln_2", # Step3-VL "model.qwen2_model.model.model.layers.{bid}.post_attention_layernorm", # Deepseek-OCR-2 qwen2 "model.vision_tower.layers.{bid}.norm2", # muse-glimmer ), MODEL_TENSOR.V_ENC_FFN_UP: ( + "vision_encoder.blocks.{bid}.mlp.fc3", # dots3note "model.vision_tower.vision_model.encoder.layers.{bid}.mlp.fc1", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.mlp.fc1", "model.vision_tower.encoder.layers.{bid}.mlp.fc1", # minicpmv4_6 @@ -1737,6 +1747,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_FFN_GATE: ( + "vision_encoder.blocks.{bid}.mlp.fc1", # dots3note "vision_tower.transformer.layers.{bid}.feed_forward.gate_proj", # pixtral-hf "vision_encoder.transformer.layers.{bid}.feed_forward.w1", # pixtral "visual.blocks.{bid}.mlp.gate_proj", # qwen2.5vl @@ -1745,6 +1756,7 @@ class TensorNameMap: ), MODEL_TENSOR.V_ENC_FFN_DOWN: ( + "vision_encoder.blocks.{bid}.mlp.fc2", # dots3note "model.vision_tower.vision_model.encoder.layers.{bid}.mlp.fc2", # Granite4Vision "vision_tower.vision_model.encoder.layers.{bid}.mlp.fc2", "model.vision_tower.encoder.layers.{bid}.mlp.fc2", # minicpmv4_6 @@ -1769,6 +1781,29 @@ class TensorNameMap: "model.vision_tower.layers.{bid}.mlp.fc2", # muse-glimmer ), + + MODEL_TENSOR.V_ENC_FFN_GATE_INP: ( + "vision_encoder.blocks.{bid}.mlp.gate_weight", # dots3note + ), + + MODEL_TENSOR.V_ENC_FFN_EXP_PROBS_B: ( + "vision_encoder.blocks.{bid}.mlp.router_bias", # dots3note + ), + + # note: expert weights are stacked into a single 3D tensor in conversion code, + # which emits the pseudo-names below + MODEL_TENSOR.V_ENC_FFN_GATE_EXPS: ( + "vision_encoder.blocks.{bid}.mlp.experts.fc1", # dots3note + ), + + MODEL_TENSOR.V_ENC_FFN_UP_EXPS: ( + "vision_encoder.blocks.{bid}.mlp.experts.fc3", # dots3note + ), + + MODEL_TENSOR.V_ENC_FFN_DOWN_EXPS: ( + "vision_encoder.blocks.{bid}.mlp.experts.fc2", # dots3note + ), + MODEL_TENSOR.V_ENC_ATTN_POST_NORM: ( "vision_model.model.layers.{bid}.post_attention_layernorm", # gemma4 ), @@ -1800,6 +1835,7 @@ class TensorNameMap: "vision_model.layernorm_pre", # llama4 "model.vision_model.pre_layrnorm", # Deepseek-OCR CLIP "vision_tower.patch_embed.patchifier.norm", # dots.ocr + "vision_encoder.patch_embed.norm", # dots3note "vision_model.ln_pre", # Step3-VL "model.vision_tower.ln_pre", # muse-glimmer ), @@ -1821,6 +1857,7 @@ class TensorNameMap: MODEL_TENSOR.V_MM_POST_NORM: ( "visual.merger.post_projection_norm", # glm4v "vision_tower.post_trunk_norm", # dots.ocr + "vision_encoder.post_trunk_norm", # dots3note "vit.perceive.after_rms", # HunyuanVL ), @@ -1838,6 +1875,7 @@ class TensorNameMap: "mlp_AR.pre_norm", # PaddleOCR-VL "merger.ln_q", "vision_tower.merger.ln_q", # dots.ocr + "vision_encoder.adapter.ln_q", # dots3note "model.merger.mlp.0.pre_norm", # minicpmv4_6 ), @@ -2173,10 +2211,12 @@ class TensorNameMap: MODEL_TENSOR.A_ENC_CONV2D: ( "audio_tower.conv2d{bid}", # qwen3omni + "audio_encoder.dots_encoder.speech_encoder.conv2d{bid}", # dots3note ), MODEL_TENSOR.A_ENC_CONV_OUT: ( "audio_tower.conv_out", # qwen3omni + "audio_encoder.dots_encoder.speech_encoder.conv_out", # dots3note "speaker_encoder.mfa.conv", # qwen3tts speaker encoder: multi-layer feature aggregation ), @@ -2184,12 +2224,14 @@ class TensorNameMap: MODEL_TENSOR.A_POST_NORM: ( "audio_tower.layer_norm", # ultravox + "audio_encoder.dots_encoder.speech_encoder.layer_norm", # dots3note "audio_tower.ln_post", # qwen2omni "encoder.layer_norm", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_ATTN_Q: ( "audio_tower.layers.{bid}.self_attn.q_proj", # ultravox + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.self_attn.q_proj", # dots3note "conformer.layers.{bid}.self_attn.linear_q", # lfm2 "conformer.layers.{bid}.attention.attn.q_proj", # gemma3n "conformer.layers.{bid}.self_attn.q_proj", # gemma4 @@ -2200,6 +2242,7 @@ class TensorNameMap: MODEL_TENSOR.A_ENC_ATTN_K: ( "audio_tower.layers.{bid}.self_attn.k_proj", # ultravox + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.self_attn.k_proj", # dots3note "conformer.layers.{bid}.self_attn.linear_k", # lfm2 "conformer.layers.{bid}.attention.attn.k_proj", # gemma3n "conformer.layers.{bid}.self_attn.k_proj", # gemma4 @@ -2210,6 +2253,7 @@ class TensorNameMap: MODEL_TENSOR.A_ENC_ATTN_V: ( "audio_tower.layers.{bid}.self_attn.v_proj", # ultravox + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.self_attn.v_proj", # dots3note "conformer.layers.{bid}.self_attn.linear_v", # lfm2 "conformer.layers.{bid}.attention.attn.v_proj", # gemma3n "conformer.layers.{bid}.self_attn.v_proj", # gemma4 @@ -2241,6 +2285,7 @@ class TensorNameMap: MODEL_TENSOR.A_ENC_INPUT_NORM: ( "audio_tower.layers.{bid}.self_attn_layer_norm", # ultravox + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.self_attn_layer_norm", # dots3note "conformer.layers.{bid}.norm_self_att", # lfm2 "conformer.layers.{bid}.attention.pre_attn_norm", # gemma3n "sound_encoder.encoder.layers.{bid}.norm_self_att", # parakeet @@ -2250,6 +2295,7 @@ class TensorNameMap: MODEL_TENSOR.A_ENC_OUTPUT: ( "audio_tower.layers.{bid}.self_attn.out_proj", # ultravox + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.self_attn.out_proj", # dots3note "conformer.layers.{bid}.self_attn.linear_out", # lfm2 "conformer.layers.{bid}.attention.post", # gemma3n "conformer.layers.{bid}.self_attn.post", # gemma4 @@ -2260,6 +2306,7 @@ class TensorNameMap: MODEL_TENSOR.A_ENC_OUTPUT_NORM: ( "audio_tower.layers.{bid}.final_layer_norm", # ultravox + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.final_layer_norm", # dots3note "conformer.layers.{bid}.norm_out", # lfm2 "conformer.layers.{bid}.attention.post_norm", # gemma3n "sound_encoder.encoder.layers.{bid}.norm_out", # parakeet @@ -2285,6 +2332,7 @@ class TensorNameMap: ), MODEL_TENSOR.A_ENC_FFN_UP: ( + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.fc1_up", # dots3note (split from fc1 in conversion code) "audio_tower.layers.{bid}.fc1", # ultravox "conformer.layers.{bid}.feed_forward1.linear1", # lfm2 "conformer.layers.{bid}.ffw_layer_start.ffw_layer_1", # gemma3n @@ -2294,9 +2342,12 @@ class TensorNameMap: "encoder.layers.{bid}.fc1", # mimo-audio-tokenizer ), - MODEL_TENSOR.A_ENC_FFN_GATE: (), + MODEL_TENSOR.A_ENC_FFN_GATE: ( + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.fc1_gate", # dots3note (split from fc1 in conversion code) + ), MODEL_TENSOR.A_ENC_FFN_DOWN: ( + "audio_encoder.dots_encoder.speech_encoder.layers.{bid}.fc2", # dots3note "audio_tower.layers.{bid}.fc2", # ultravox "conformer.layers.{bid}.feed_forward1.linear2", # lfm2 "conformer.layers.{bid}.ffw_layer_start.ffw_layer_2", # gemma3n @@ -2380,6 +2431,7 @@ class TensorNameMap: MODEL_TENSOR.A_MMPROJ: ( "audio.multi_modal_projector.linear_{bid}", # ultravox, meralion + "audio_encoder.audio_adapter.proj.{bid}", # dots3note (proj.1, proj.3) "audio_adapter.model.{bid}", # lfm2 "audio_tower.proj{bid}", # qwen3omni "sound_projection.linear{bid}", # parakeet (linear1, linear2) @@ -2394,6 +2446,7 @@ class TensorNameMap: MODEL_TENSOR.A_MM_NORM_PRE: ( "audio.multi_modal_projector.ln_pre", # ultravox + "audio_encoder.audio_adapter.proj.0", # dots3note "sound_projection.norm", # parakeet ), diff --git a/tools/mtmd/CMakeLists.txt b/tools/mtmd/CMakeLists.txt index 63a8b17a5..e60c9c878 100644 --- a/tools/mtmd/CMakeLists.txt +++ b/tools/mtmd/CMakeLists.txt @@ -30,6 +30,7 @@ add_library(mtmd models/models.h models/cogvlm.cpp models/conformer.cpp + models/dots3note.cpp models/dotsocr.cpp models/exaone4_5.cpp models/gemma4a.cpp diff --git a/tools/mtmd/clip-graph.h b/tools/mtmd/clip-graph.h index e12140ba0..2cf1b683a 100644 --- a/tools/mtmd/clip-graph.h +++ b/tools/mtmd/clip-graph.h @@ -120,6 +120,12 @@ struct clip_graph { ffn_op_type type_op, int il) const; + ggml_tensor * build_moe_ffn( + ggml_tensor * cur, + const clip_layer & layer, + ffn_op_type type_op, + int il) const; + ggml_tensor * build_attn( ggml_tensor * wo, ggml_tensor * wo_b, diff --git a/tools/mtmd/clip-impl.h b/tools/mtmd/clip-impl.h index ea8549776..f6045093c 100644 --- a/tools/mtmd/clip-impl.h +++ b/tools/mtmd/clip-impl.h @@ -75,6 +75,7 @@ #define KEY_SAM_N_HEAD "clip.vision.sam.head_count" #define KEY_SAM_N_BLOCK "clip.vision.sam.block_count" #define KEY_SAM_N_EMBD "clip.vision.sam.embedding_length" +#define KEY_VISION_N_EXPERT_USED "clip.vision.expert_used_count" // audio-specific #define KEY_AUDIO_PROJ_TYPE "clip.audio.projector_type" // for models with mixed modalities #define KEY_A_NUM_MEL_BINS "clip.audio.num_mel_bins" @@ -119,7 +120,11 @@ #define TN_FFN_DOWN "%s.blk.%d.ffn_down.%s" #define TN_FFN_GATE "%s.blk.%d.ffn_gate.%s" #define TN_FFN_UP "%s.blk.%d.ffn_up.%s" -#define TN_FFN_GATE "%s.blk.%d.ffn_gate.%s" +#define TN_FFN_GATE_INP "%s.blk.%d.ffn_gate_inp.%s" // MoE router (dots3note) +#define TN_FFN_GATE_EXPS "%s.blk.%d.ffn_gate_exps.%s" +#define TN_FFN_UP_EXPS "%s.blk.%d.ffn_up_exps.%s" +#define TN_FFN_DOWN_EXPS "%s.blk.%d.ffn_down_exps.%s" +#define TN_FFN_EXP_PROBS_B "%s.blk.%d.exp_probs_b.%s" #define TN_LN_1 "%s.blk.%d.ln1.%s" // layer norm #define TN_LN_2 "%s.blk.%d.ln2.%s" // layer norm #define TN_LS_1 "%s.blk.%d.ls1.%s" // layer scale @@ -471,6 +476,8 @@ enum projector_type { PROJECTOR_TYPE_COGVLM, PROJECTOR_TYPE_JANUS_PRO, PROJECTOR_TYPE_DOTS_OCR, + PROJECTOR_TYPE_DOTS3NOTE_V, + PROJECTOR_TYPE_DOTS3NOTE_A, PROJECTOR_TYPE_DEEPSEEKOCR, PROJECTOR_TYPE_DEEPSEEKOCR2, PROJECTOR_TYPE_LFM2A, @@ -533,6 +540,8 @@ static std::map PROJECTOR_TYPE_NAMES = { { PROJECTOR_TYPE_COGVLM, "cogvlm"}, { PROJECTOR_TYPE_JANUS_PRO, "janus_pro"}, { PROJECTOR_TYPE_DOTS_OCR, "dots_ocr"}, + { PROJECTOR_TYPE_DOTS3NOTE_V, "dots3note_v"}, + { PROJECTOR_TYPE_DOTS3NOTE_A, "dots3note_a"}, { PROJECTOR_TYPE_DEEPSEEKOCR, "deepseekocr"}, { PROJECTOR_TYPE_DEEPSEEKOCR2, "deepseekocr2"}, { PROJECTOR_TYPE_LFM2A, "lfm2a"}, diff --git a/tools/mtmd/clip-model.h b/tools/mtmd/clip-model.h index ad25c008e..fcdabd633 100644 --- a/tools/mtmd/clip-model.h +++ b/tools/mtmd/clip-model.h @@ -93,6 +93,7 @@ struct clip_hparams { float eps = 1e-6; float rope_theta = 0.0; + int32_t n_expert_used = 0; std::vector feature_layers; int32_t attn_window_size = 0; int32_t n_wa_pattern = 0; @@ -259,6 +260,13 @@ struct clip_layer { ggml_tensor * ff_down_w = nullptr; ggml_tensor * ff_down_b = nullptr; + // MoE FFN (dots3note vision pyramid blocks) + ggml_tensor * ff_gate_inp_w = nullptr; + ggml_tensor * ff_gate_exps_w = nullptr; + ggml_tensor * ff_up_exps_w = nullptr; + ggml_tensor * ff_down_exps_w = nullptr; + ggml_tensor * ff_exp_probs_b = nullptr; + // layernorm 2 (or pre-FFN norm) ggml_tensor * ln_2_w = nullptr; ggml_tensor * ln_2_b = nullptr; diff --git a/tools/mtmd/clip.cpp b/tools/mtmd/clip.cpp index 45e33042d..9977ed490 100644 --- a/tools/mtmd/clip.cpp +++ b/tools/mtmd/clip.cpp @@ -514,11 +514,13 @@ ggml_tensor * clip_graph::build_vit( cb(cur, "ffn_inp_normed", il); // ffn - cur = build_ffn(cur, - layer.ff_up_w, layer.ff_up_b, - layer.ff_gate_w, layer.ff_gate_b, - layer.ff_down_w, layer.ff_down_b, - ffn_t, il); + cur = layer.ff_gate_exps_w + ? build_moe_ffn(cur, layer, ffn_t, il) + : build_ffn(cur, + layer.ff_up_w, layer.ff_up_b, + layer.ff_gate_w, layer.ff_gate_b, + layer.ff_down_w, layer.ff_down_b, + ffn_t, il); cb(cur, "ffn_out", il); @@ -699,6 +701,50 @@ ggml_tensor * clip_graph::build_ffn( return cur; } +// MoE FFN with sigmoid router and normalized top-k weights (dots3note vision) +// the router runs in fp32; exp_probs_b only affects expert selection, not the weights +ggml_tensor * clip_graph::build_moe_ffn(ggml_tensor * cur, const clip_layer & layer, ffn_op_type type_op, int il) const { + const int64_t n_tokens = cur->ne[1]; + const int64_t n_expert = layer.ff_gate_exps_w->ne[2]; + const int64_t n_expert_used = std::min((int64_t) hparams.n_expert_used, n_expert); + GGML_ASSERT(n_expert_used > 0); + GGML_ASSERT(type_op == FFN_SILU); + + ggml_tensor * probs = ggml_sigmoid(ctx0, build_mm(layer.ff_gate_inp_w, cur)); // [n_expert, n_tokens] + cb(probs, "ffn_moe_probs", il); + + ggml_tensor * sel = layer.ff_exp_probs_b + ? ggml_add(ctx0, probs, layer.ff_exp_probs_b) + : probs; + ggml_tensor * selected = ggml_top_k(ctx0, sel, n_expert_used); // [n_expert_used, n_tokens] + + ggml_tensor * weights = ggml_get_rows(ctx0, + ggml_reshape_3d(ctx0, probs, 1, n_expert, n_tokens), selected); + weights = ggml_reshape_2d(ctx0, weights, n_expert_used, n_tokens); + weights = ggml_div(ctx0, weights, ggml_sum_rows(ctx0, weights)); + weights = ggml_reshape_3d(ctx0, weights, 1, n_expert_used, n_tokens); + cb(weights, "ffn_moe_weights", il); + + cur = ggml_reshape_3d(ctx0, cur, cur->ne[0], 1, n_tokens); + ggml_tensor * gate = ggml_mul_mat_id(ctx0, layer.ff_gate_exps_w, cur, selected); // [n_ff, n_expert_used, n_tokens] + ggml_tensor * up = ggml_mul_mat_id(ctx0, layer.ff_up_exps_w, cur, selected); + cur = ggml_mul(ctx0, ggml_silu(ctx0, gate), up); + cur = ggml_mul_mat_id(ctx0, layer.ff_down_exps_w, cur, selected); // [n_embd, n_expert_used, n_tokens] + cur = ggml_mul(ctx0, cur, weights); + + // sum over the selected experts + ggml_tensor * out = nullptr; + for (int64_t i = 0; i < n_expert_used; i++) { + ggml_tensor * v = ggml_view_2d(ctx0, cur, cur->ne[0], n_tokens, cur->nb[2], i * cur->nb[1]); + out = out ? ggml_add(ctx0, out, v) : v; + } + if (n_expert_used == 1) { + out = ggml_cont(ctx0, out); + } + cb(out, "ffn_moe_out", il); + return out; +} + ggml_tensor * clip_graph::build_attn( ggml_tensor * wo, ggml_tensor * wo_b, @@ -933,9 +979,14 @@ static std::unique_ptr clip_get_graph_builder(clip_ctx * ctx, const builder = std::make_unique(ctx, img); } break; case PROJECTOR_TYPE_DOTS_OCR: + case PROJECTOR_TYPE_DOTS3NOTE_V: // same ViT + merger; pyramid MoE is handled by build_vit { builder = std::make_unique(ctx, img); } break; + case PROJECTOR_TYPE_DOTS3NOTE_A: + { + builder = std::make_unique(ctx, img); + } break; case PROJECTOR_TYPE_QWEN2VL: case PROJECTOR_TYPE_QWEN25VL: { @@ -1510,6 +1561,25 @@ struct clip_model_loader { get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels); hparams.set_warmup_n_tokens(46*46); // avoid OOM on warmup } break; + case PROJECTOR_TYPE_DOTS3NOTE_V: + { + hparams.rope_theta = 10000.0f; + hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW; + get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge); + get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels); + get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels); + get_u32(KEY_VISION_N_EXPERT_USED, hparams.n_expert_used); + hparams.set_warmup_n_tokens(46*46); // avoid OOM on warmup + } break; + case PROJECTOR_TYPE_DOTS3NOTE_A: + { + hparams.rope_theta = 10000.0f; + hparams.audio_chunk_len = 60; // in seconds + hparams.audio_sample_rate = 16000; + hparams.audio_n_fft = 400; + hparams.audio_window_len = 400; + hparams.audio_hop_len = 160; + } break; case PROJECTOR_TYPE_KIMIVL: { hparams.image_resize_algo = RESIZE_ALGO_BILINEAR; @@ -2190,12 +2260,20 @@ struct clip_model_loader { layer.ln_1_b = get_tensor(string_format(TN_LN_1, prefix, il, "bias"), false); layer.ln_2_b = get_tensor(string_format(TN_LN_2, prefix, il, "bias"), false); + // MoE ffn (dots3note vision pyramid blocks); replaces the dense ffn when present + layer.ff_gate_inp_w = get_tensor(string_format(TN_FFN_GATE_INP, prefix, il, "weight"), false); + layer.ff_gate_exps_w = get_tensor(string_format(TN_FFN_GATE_EXPS, prefix, il, "weight"), false); + layer.ff_up_exps_w = get_tensor(string_format(TN_FFN_UP_EXPS, prefix, il, "weight"), false); + layer.ff_down_exps_w = get_tensor(string_format(TN_FFN_DOWN_EXPS, prefix, il, "weight"), false); + layer.ff_exp_probs_b = get_tensor(string_format(TN_FFN_EXP_PROBS_B, prefix, il, "weight"), false); + const bool is_moe = layer.ff_gate_exps_w != nullptr; + // ffn - layer.ff_up_w = get_tensor(string_format(TN_FFN_UP, prefix, il, "weight")); + layer.ff_up_w = get_tensor(string_format(TN_FFN_UP, prefix, il, "weight"), !is_moe); layer.ff_up_b = get_tensor(string_format(TN_FFN_UP, prefix, il, "bias"), false); layer.ff_gate_w = get_tensor(string_format(TN_FFN_GATE, prefix, il, "weight"), false); layer.ff_gate_b = get_tensor(string_format(TN_FFN_GATE, prefix, il, "bias"), false); - layer.ff_down_w = get_tensor(string_format(TN_FFN_DOWN, prefix, il, "weight")); + layer.ff_down_w = get_tensor(string_format(TN_FFN_DOWN, prefix, il, "weight"), !is_moe); layer.ff_down_b = get_tensor(string_format(TN_FFN_DOWN, prefix, il, "bias"), false); // mimovl per-head attention sink bias @@ -2677,6 +2755,7 @@ struct clip_model_loader { model.mm_patch_merger_w = get_tensor(string_format(TN_MM_PATCH_MERGER, "weight"), false); } break; case PROJECTOR_TYPE_DOTS_OCR: + case PROJECTOR_TYPE_DOTS3NOTE_V: { model.mm_0_w = get_tensor(string_format(TN_LLAVA_PROJ, 0, "weight")); model.mm_0_b = get_tensor(string_format(TN_LLAVA_PROJ, 0, "bias")); @@ -2687,6 +2766,23 @@ struct clip_model_loader { // post_trunk_norm: applied after all ViT blocks, before the merger model.post_ln_w = get_tensor(string_format(TN_MM_POST_NORM, "weight")); } break; + case PROJECTOR_TYPE_DOTS3NOTE_A: + { + model.conv2d_1_w = get_tensor(string_format(TN_CONV2D, 1, "weight")); + model.conv2d_1_b = get_tensor(string_format(TN_CONV2D, 1, "bias")); + model.conv2d_2_w = get_tensor(string_format(TN_CONV2D, 2, "weight")); + model.conv2d_2_b = get_tensor(string_format(TN_CONV2D, 2, "bias")); + model.conv2d_3_w = get_tensor(string_format(TN_CONV2D, 3, "weight")); + model.conv2d_3_b = get_tensor(string_format(TN_CONV2D, 3, "bias")); + model.conv_out_w = get_tensor(string_format(TN_CONV_OUT, "weight")); // no bias + // adapter: LayerNorm -> Linear -> GELU -> Linear + model.mm_norm_pre_w = get_tensor(string_format(TN_MM_NORM_PRE, "weight")); + model.mm_norm_pre_b = get_tensor(string_format(TN_MM_NORM_PRE, "bias")); + model.mm_1_w = get_tensor(string_format(TN_MM_AUDIO_MLP, 1, "weight")); + model.mm_1_b = get_tensor(string_format(TN_MM_AUDIO_MLP, 1, "bias")); + model.mm_2_w = get_tensor(string_format(TN_MM_AUDIO_MLP, 3, "weight")); + model.mm_2_b = get_tensor(string_format(TN_MM_AUDIO_MLP, 3, "bias")); + } break; case PROJECTOR_TYPE_ULTRAVOX: { model.conv1d_1_w = get_tensor(string_format(TN_CONV1D, 1, "weight")); @@ -4075,12 +4171,18 @@ int clip_n_output_tokens(const clip_ctx * ctx, const clip_image_f32 * img) { } break; case PROJECTOR_TYPE_PADDLEOCR: case PROJECTOR_TYPE_DOTS_OCR: + case PROJECTOR_TYPE_DOTS3NOTE_V: { // dynamic size int n_merge = ctx->model.hparams.n_merge; int stride = n_merge * n_merge; n_patches = CLIP_ALIGN(n_patches, stride) / stride; } break; + case PROJECTOR_TYPE_DOTS3NOTE_A: + { + // 3x stride-2 conv2d over mel frames + n_patches = (img->nx() + 7) / 8; + } break; case PROJECTOR_TYPE_PIXTRAL: case PROJECTOR_TYPE_LIGHTONOCR: { @@ -4727,6 +4829,7 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) { set_input_i32("minimax_pos_w", pos_w); } break; case PROJECTOR_TYPE_DOTS_OCR: + case PROJECTOR_TYPE_DOTS3NOTE_V: { const int pw = image_size_width / patch_size; const int ph = image_size_height / patch_size; @@ -5217,6 +5320,16 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) { } set_input_i32("pos_w", pos_data); } break; + case PROJECTOR_TYPE_DOTS3NOTE_A: + { + GGML_ASSERT(imgs.entries.size() == 1); + const int n_pos = (imgs.entries.front().nx() + 7) / 8; // 3x stride-2 conv2d + std::vector positions(n_pos); + for (int i = 0; i < n_pos; i++) { + positions[i] = i; + } + set_input_i32("positions", positions); + } break; case PROJECTOR_TYPE_GEMMA4A: { GGML_ASSERT(imgs.entries.size() == 1); @@ -5713,6 +5826,8 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) { case PROJECTOR_TYPE_PIXTRAL: case PROJECTOR_TYPE_LIGHTONOCR: case PROJECTOR_TYPE_DOTS_OCR: + case PROJECTOR_TYPE_DOTS3NOTE_V: + case PROJECTOR_TYPE_DOTS3NOTE_A: return ctx->model.mm_2_w->ne[1]; case PROJECTOR_TYPE_MLP_NORM: return ctx->model.mm_3_b->ne[0]; diff --git a/tools/mtmd/models/dots3note.cpp b/tools/mtmd/models/dots3note.cpp new file mode 100644 index 000000000..93c14fc79 --- /dev/null +++ b/tools/mtmd/models/dots3note.cpp @@ -0,0 +1,61 @@ +#include "models.h" + +ggml_cgraph * clip_graph_dots3note_a::build() { + // inp_raw: [n_frames, n_mel, 1], one 60s chunk, mel frames not padded + // the reference impl zero-masks conv inputs beyond the valid length at each stage; + // running on exactly the valid frames with the convs' zero padding is equivalent + ggml_tensor * inp = build_inp_raw(1); + GGML_ASSERT(inp->type == GGML_TYPE_F32); + + // 3x conv2d (k=3, s=2, p=1) + gelu + { + auto conv_block = [&](ggml_tensor * x, ggml_tensor * w, ggml_tensor * b) { + x = ggml_conv_2d(ctx0, w, x, 2, 2, 1, 1, 1, 1); + x = ggml_add(ctx0, x, ggml_reshape_4d(ctx0, b, 1, 1, x->ne[2], 1)); + return ggml_gelu_erf(ctx0, x); + }; + + inp = conv_block(inp, model.conv2d_1_w, model.conv2d_1_b); + inp = conv_block(inp, model.conv2d_2_w, model.conv2d_2_b); + inp = conv_block(inp, model.conv2d_3_w, model.conv2d_3_b); + // inp: [OW=n_frames/8, OH=n_mel/8, OC=480, 1] + cb(inp, "after_conv_stem", -1); + } + + // [OW, OH, OC, 1] -> [OH*OC, OW], feature index f + OH*c (matches the reference permute+reshape) + inp = ggml_cont(ctx0, ggml_permute(ctx0, inp, 2, 0, 1, 3)); + inp = ggml_reshape_2d(ctx0, inp, inp->ne[0] * inp->ne[1], inp->ne[2]); + + // project to d_model (no bias) + inp = ggml_mul_mat(ctx0, model.conv_out_w, inp); + cb(inp, "after_conv_out", -1); + + const int64_t n_pos = inp->ne[1]; + + ggml_tensor * positions = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_pos); + ggml_set_name(positions, "positions"); + ggml_set_input(positions); + + // partial rotary: first half of each head, NEOX style + auto add_pos = [&](ggml_tensor * cur, const clip_layer &) { + return ggml_rope_ext(ctx0, cur, positions, nullptr, d_head/2, + GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f); + }; + + ggml_tensor * cur = build_vit(inp, n_pos, + NORM_TYPE_RMS, hparams.ffn_op, + nullptr, add_pos); + cb(cur, "after_transformer", -1); + + // adapter: LayerNorm -> Linear -> GELU -> Linear + cur = build_norm(cur, model.mm_norm_pre_w, model.mm_norm_pre_b, NORM_TYPE_NORMAL, 1e-5, -1); + cur = build_ffn(cur, + model.mm_1_w, model.mm_1_b, + nullptr, nullptr, + model.mm_2_w, model.mm_2_b, + FFN_GELU_ERF, -1); + cb(cur, "projected", -1); + + ggml_build_forward_expand(gf, cur); + return gf; +} diff --git a/tools/mtmd/models/models.h b/tools/mtmd/models/models.h index 3631d849b..10546fa5d 100644 --- a/tools/mtmd/models/models.h +++ b/tools/mtmd/models/models.h @@ -119,6 +119,11 @@ struct clip_graph_dotsocr : clip_graph { ggml_cgraph * build() override; }; +struct clip_graph_dots3note_a : clip_graph { + clip_graph_dots3note_a(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {} + ggml_cgraph * build() override; +}; + struct clip_graph_cogvlm : clip_graph { clip_graph_cogvlm(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {} ggml_cgraph * build() override; diff --git a/tools/mtmd/mtmd-audio.cpp b/tools/mtmd/mtmd-audio.cpp index 98a8c11ee..ce08f9e93 100644 --- a/tools/mtmd/mtmd-audio.cpp +++ b/tools/mtmd/mtmd-audio.cpp @@ -723,6 +723,100 @@ bool mtmd_audio_preprocessor_qwen3a::preprocess(const float * sa return true; } +// +// mtmd_audio_preprocessor_dots3note +// +// Matches Dots3NoteFeatureExtractor: the waveform is split into 60s chunks and each chunk gets +// its own whisper-style log-mel (center=True, log10 + (max-8)/4). Only sample_length//hop frames +// per chunk are valid; the reference masks everything beyond them, so we emit exactly that many. +// + +void mtmd_audio_preprocessor_dots3note::initialize() { + cache.fill_sin_cos_table(hparams.audio_n_fft); + cache.fill_hann_window(hparams.audio_window_len, true); + cache.fill_mel_filterbank_matrix(hparams.n_mel_bins, hparams.audio_n_fft, hparams.audio_sample_rate); +} + +bool mtmd_audio_preprocessor_dots3note::preprocess(const float * samples, + size_t n_samples, + std::vector & output) { + if (n_samples == 0) { + return false; + } + + GGML_ASSERT(!cache.sin_vals.empty()); + GGML_ASSERT(!cache.cos_vals.empty()); + GGML_ASSERT(!cache.filters.data.empty()); + + const int pad = hparams.audio_n_fft / 2; // center=True padding + const int hop = hparams.audio_hop_len; + const size_t chunk_samples = (size_t) hparams.audio_chunk_len * hparams.audio_sample_rate; + + for (size_t start = 0; start < n_samples; start += chunk_samples) { + const size_t n_chunk = std::min(chunk_samples, n_samples - start); + const float * chunk = samples + start; + + const int64_t n_valid = n_chunk / hop; + if (n_valid == 0) { + continue; // sub-hop tail, contributes no frames + } + + // reflect-pad the start; the reference zero-pads partial chunks to 60s before the STFT, + // so a partial chunk sees zeros past its end while a full chunk reflects its own tail + std::vector padded(n_chunk + 2 * pad, 0.0f); + for (int i = 0; i < pad; i++) { + int src = pad - i; + padded[i] = (src < (int) n_chunk) ? chunk[src] : 0.0f; + } + std::copy(chunk, chunk + n_chunk, padded.begin() + pad); + if (n_chunk == chunk_samples) { + for (int i = 0; i < pad; i++) { + int src = (int) n_chunk - 2 - i; + padded[n_chunk + pad + i] = (src >= 0) ? chunk[src] : 0.0f; + } + } + + filter_params params; + params.n_mel = hparams.n_mel_bins; + params.n_fft_bins = 1 + (hparams.audio_n_fft / 2); + params.hann_window_size = hparams.audio_window_len; + params.hop_length = hop; + params.sample_rate = hparams.audio_sample_rate; + params.no_padding = true; // padding already applied above + params.use_natural_log = false; + + mtmd_audio_mel mel_full; + if (!log_mel_spectrogram(padded.data(), (int) padded.size(), 4, params, cache, mel_full)) { + return false; + } + GGML_ASSERT(mel_full.n_len >= n_valid); + + // per-chunk whisper-style normalization, then keep only the valid frames + mtmd_audio_mel out; + out.n_mel = mel_full.n_mel; + out.n_len = n_valid; + out.n_len_org = n_valid; + out.data.resize((size_t) out.n_mel * (size_t) out.n_len); + + double mmax = -1e20; + for (int64_t m = 0; m < out.n_mel; m++) { + for (int64_t t = 0; t < n_valid; t++) { + mmax = std::max(mmax, (double) mel_full.data[(size_t) m * mel_full.n_len + t]); + } + } + mmax -= 8.0; + for (int64_t m = 0; m < out.n_mel; m++) { + for (int64_t t = 0; t < n_valid; t++) { + const double v = std::max((double) mel_full.data[(size_t) m * mel_full.n_len + t], mmax); + out.data[(size_t) m * n_valid + t] = (float) ((v + 4.0) / 4.0); + } + } + + output.push_back(std::move(out)); + } + return !output.empty(); +} + // // mtmd_audio_preprocessor_mimo_audio // diff --git a/tools/mtmd/mtmd-audio.h b/tools/mtmd/mtmd-audio.h index 44ad098ae..0f47d4502 100644 --- a/tools/mtmd/mtmd-audio.h +++ b/tools/mtmd/mtmd-audio.h @@ -111,6 +111,15 @@ struct mtmd_audio_preprocessor_qwen3a : mtmd_audio_preprocessor { mtmd_audio_cache cache; }; +struct mtmd_audio_preprocessor_dots3note : mtmd_audio_preprocessor { + mtmd_audio_preprocessor_dots3note(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {} + void initialize() override; + bool preprocess(const float * samples, size_t n_samples, std::vector & output) override; + + private: + mtmd_audio_cache cache; +}; + struct mtmd_audio_preprocessor_mimo_audio : mtmd_audio_preprocessor { mtmd_audio_preprocessor_mimo_audio(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {} void initialize() override; diff --git a/tools/mtmd/mtmd.cpp b/tools/mtmd/mtmd.cpp index 95f17f7af..5b306180d 100644 --- a/tools/mtmd/mtmd.cpp +++ b/tools/mtmd/mtmd.cpp @@ -825,6 +825,7 @@ struct mtmd_context { image_preproc = std::make_unique(ctx_v); } break; case PROJECTOR_TYPE_DOTS_OCR: + case PROJECTOR_TYPE_DOTS3NOTE_V: { // <|img|> ... (image embeddings) ... <|endofimg|> img_beg = "<|img|>"; @@ -976,6 +977,13 @@ struct mtmd_context { aud_end = ""; audio_preproc = std::make_unique(ctx_a); } break; + case PROJECTOR_TYPE_DOTS3NOTE_A: + { + // <|audio_comp_start|> ... (embeddings) ... <|audio_comp_end|> + aud_beg = "<|audio_comp_start|>"; + aud_end = "<|audio_comp_end|>"; + audio_preproc = std::make_unique(ctx_a); + } break; case PROJECTOR_TYPE_MIMO_AUDIO: { aud_beg = "<|mimo_audio_start|>"; From 2115b73d8ebdbd659075cce66c609506863bc826 Mon Sep 17 00:00:00 2001 From: Tiwei Bie Date: Sat, 22 Aug 2026 17:19:48 +0800 Subject: [PATCH 13/16] model : support DSpark for bailingmoe3 (#27508) --- conversion/__init__.py | 1 + conversion/qwen.py | 8 +++++++- src/llama-arch.cpp | 1 + src/models/bailingmoe3.cpp | 41 +++++++++++++++++++++++--------------- 4 files changed, 34 insertions(+), 17 deletions(-) diff --git a/conversion/__init__.py b/conversion/__init__.py index aaaebba97..8de97e959 100644 --- a/conversion/__init__.py +++ b/conversion/__init__.py @@ -58,6 +58,7 @@ TEXT_MODEL_MAP: dict[str, str] = { "DSparkDraftModel": "qwen", "DSparkSpeculator": "qwen", "Lfm2DSparkDraftModel": "qwen", + "LingDSparkModel": "qwen", "DeepseekV4ForCausalLM": "deepseek", "DeepseekV4DSparkModel": "deepseek", "DistilBertForMaskedLM": "bert", diff --git a/conversion/qwen.py b/conversion/qwen.py index 355365763..cdba8a63e 100644 --- a/conversion/qwen.py +++ b/conversion/qwen.py @@ -709,7 +709,13 @@ class DFlashModel(Qwen3Model): yield from super().modify_tensors(data_torch, name, bid) -@ModelBase.register("Qwen3DSparkModel", "DSparkDraftModel", "DSparkSpeculator", "Lfm2DSparkDraftModel") +@ModelBase.register( + "Qwen3DSparkModel", + "DSparkDraftModel", + "DSparkSpeculator", + "Lfm2DSparkDraftModel", + "LingDSparkModel", +) @ModelBase.example("satgeze/Qwen3.6-27B-DSpark") class DSparkModel(DFlashModel): # DSpark = DFlash + a semi-autoregressive Markov head. diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index 60b8e461c..025f9fb54 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -1038,6 +1038,7 @@ bool llm_arch_supports_rs_rollback(const llm_arch & arch) { case LLM_ARCH_NEMOTRON_H_MOE: case LLM_ARCH_LFM2: case LLM_ARCH_LFM2MOE: + case LLM_ARCH_BAILINGMOE3: return true; default: return false; diff --git a/src/models/bailingmoe3.cpp b/src/models/bailingmoe3.cpp index f5855696e..0637931cc 100644 --- a/src/models/bailingmoe3.cpp +++ b/src/models/bailingmoe3.cpp @@ -1,6 +1,8 @@ #include "models.h" #include "llama-memory-recurrent.h" +#include + void llama_model_bailingmoe3::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_ATTENTION_KEY_LENGTH_MLA, hparams.n_embd_head_k_mla_impl); @@ -179,7 +181,9 @@ static ggml_tensor * bailingmoe3_causal_conv1d( int64_t n_seq_tokens, int64_t n_seqs, int64_t n_tokens, - int64_t cache_head) { + int64_t cache_head, + uint32_t mem_size, + uint32_t n_rs_seq) { const int64_t d_inner = head_dim * n_head; const int64_t conv_state_size = (d_conv - 1) * d_inner; const int64_t total_state_size = 3 * conv_state_size; @@ -193,13 +197,18 @@ static ggml_tensor * bailingmoe3_causal_conv1d( x_proj = ggml_reshape_3d(ctx0, x_proj, d_inner, n_seq_tokens, n_seqs); ggml_tensor * conv_x = ggml_concat(ctx0, conv_state, ggml_transpose(ctx0, x_proj), 0); - ggml_tensor * last_conv_x = ggml_view_3d(ctx0, conv_x, d_conv - 1, d_inner, n_seqs, - conv_x->nb[1], conv_x->nb[2], n_seq_tokens * conv_x->nb[0]); - ggml_build_forward_expand(gf, ggml_cpy(ctx0, last_conv_x, - ggml_view_3d(ctx0, conv_states_all, d_conv - 1, d_inner, n_seqs, - (d_conv - 1) * ggml_element_size(conv_states_all), - total_state_size * ggml_element_size(conv_states_all), - (cache_head * total_state_size + qkv * conv_state_size) * ggml_element_size(conv_states_all)))); + const int64_t K = (int64_t) n_rs_seq + 1; + const int64_t n_written = std::min(n_seq_tokens, K); + + for (int64_t slot = 0; slot < n_written; ++slot) { + ggml_tensor * conv_snap = ggml_view_3d(ctx0, conv_x, d_conv - 1, d_inner, n_seqs, + conv_x->nb[1], conv_x->nb[2], (conv_x->ne[0] - (d_conv - 1) - slot) * conv_x->nb[0]); + ggml_build_forward_expand(gf, ggml_cpy(ctx0, conv_snap, + ggml_view_3d(ctx0, conv_states_all, d_conv - 1, d_inner, n_seqs, + (d_conv - 1) * ggml_element_size(conv_states_all), + total_state_size * ggml_element_size(conv_states_all), + ((slot * mem_size + cache_head) * total_state_size + qkv * conv_state_size) * ggml_element_size(conv_states_all)))); + } ggml_tensor * conv_weight = ggml_reshape_2d(ctx0, conv_w, d_conv, d_inner); ggml_tensor * out = ggml_ssm_conv(ctx0, conv_x, conv_weight); @@ -237,6 +246,8 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph GGML_ASSERT(ubatch.n_tokens == n_seq_tokens * n_seqs); for (int il = 0; il < n_layer; ++il) { + res->t_layer_inp[il] = inpL; + const auto & layer = model.layers[il]; ggml_tensor * inpSA = inpL; ggml_tensor * cur = build_norm(inpL, layer.attn_norm, nullptr, LLM_NORM_RMS, il); @@ -245,18 +256,19 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph if (hparams.is_recr(il)) { const auto * mctx_cur = inp_rs->mctx; const auto cache_head = mctx_cur->get_head(); + const auto mem_size = mctx_cur->get_size(); ggml_tensor * conv_states_all = mctx_cur->get_r_l(il); ggml_tensor * conv_state_all = build_rs(inp_rs, conv_states_all, hparams.n_embd_r(), n_seqs); ggml_tensor * q = bailingmoe3_causal_conv1d( gf, ctx0, conv_states_all, conv_state_all, 0, cur, layer.wq, layer.ssm_q_conv, - d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head); + d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head, mem_size, cparams.n_rs_seq); ggml_tensor * k = bailingmoe3_causal_conv1d( gf, ctx0, conv_states_all, conv_state_all, 1, cur, layer.wk, layer.ssm_k_conv, - d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head); + d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head, mem_size, cparams.n_rs_seq); ggml_tensor * v = bailingmoe3_causal_conv1d( gf, ctx0, conv_states_all, conv_state_all, 2, cur, layer.wv, layer.ssm_v_conv, - d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head); + d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head, mem_size, cparams.n_rs_seq); ggml_tensor * gate = ggml_mul_mat(ctx0, layer.ssm_f_a, cur); gate = ggml_add(ctx0, gate, layer.ssm_dt_b); @@ -276,11 +288,8 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph ggml_tensor * state = build_rs(inp_rs, states_all, hparams.n_embd_s(), n_seqs); state = ggml_reshape_4d(ctx0, state, head_dim, head_dim, n_head, n_seqs); - auto result = build_delta_net(q, k, v, gate, beta, state, il); - ggml_tensor * out = ggml_cont(ctx0, result.first); - ggml_build_forward_expand(gf, ggml_cpy(ctx0, result.second, - ggml_view_1d(ctx0, states_all, hparams.n_embd_s() * n_seqs, - cache_head * hparams.n_embd_s() * ggml_element_size(states_all)))); + ggml_tensor * out = ggml_cont(ctx0, build_recurrent_attn( + inp_rs, states_all, q, k, v, gate, beta, state, il)); ggml_tensor * out_gate = ggml_mul_mat(ctx0, layer.ssm_g_a, cur); out_gate = ggml_reshape_3d(ctx0, out_gate, head_dim, n_head, n_tokens); From e85caa81ea2b65797396018c179b87ad61fa38ab Mon Sep 17 00:00:00 2001 From: Mario Limonciello Date: Sat, 22 Aug 2026 05:28:30 -0500 Subject: [PATCH 14/16] ci : Restore ROCm job for Ubuntu (#27399) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Revert "ci : disable ubuntu-rocm (#26969)" This reverts commit 9558fa44c92746a58dd07ad1bf0c889715b938a6. * ci: set ccache compiler_check=content for ROCm build The ROCm toolchain is pip-installed fresh on every run, so the clang binary's mtime changes each time. With ccache's default compiler_check=mtime that invalidates the whole cache and warm builds only reached ~70% hits. Hash the compiler contents instead so the cache survives toolchain reinstalls. * Update ccache size to 1GB We're waivering with so many architectures built, we need a bigger ccache limit. * merge fix --------- Co-authored-by: Jim Wu Co-authored-by: Sigbjørn Skjæret --- .github/workflows/release.yml | 212 ++++++++++++++++++---------------- 1 file changed, 112 insertions(+), 100 deletions(-) diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index 90aab37e8..8b7777bc0 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -774,6 +774,7 @@ jobs: with: key: windows-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }} evict-old-files: 1d + max-size: "1G" # - name: Cache ROCm Installation # id: cache-rocm @@ -1286,123 +1287,134 @@ jobs: with: key: release-ubuntu-24.04-sycl-${{ matrix.build }} -# ubuntu-22-rocm: -# needs: [check-release, get-version] -# if: ${{ needs.check-release.outputs.should_release == 'true' }} + ubuntu-22-rocm: + needs: [check-release, get-version] + if: ${{ needs.check-release.outputs.should_release == 'true' }} -# runs-on: ubuntu-22.04 + runs-on: ubuntu-22.04 -# permissions: -# actions: write + permissions: + actions: write -# strategy: -# matrix: -# include: -# - ROCM_VERSION: "7.14.0" -# gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201" -# build: 'x64' + strategy: + matrix: + include: + - ROCM_VERSION: "7.14.0" + gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201" + build: 'x64' -# steps: -# - name: Clone -# id: checkout -# uses: actions/checkout@v6 -# with: -# fetch-depth: 0 + steps: + - name: Clone + id: checkout + uses: actions/checkout@v6 + with: + fetch-depth: 0 -# - name: Setup Node.js -# uses: actions/setup-node@v6 -# with: -# node-version: "24" -# cache: "npm" -# cache-dependency-path: "tools/ui/package-lock.json" + - name: Setup Node.js + uses: actions/setup-node@v6 + with: + node-version: "24" + cache: "npm" + cache-dependency-path: "tools/ui/package-lock.json" -# - name: Free up disk space -# uses: ggml-org/free-disk-space@v1.3.1 -# with: -# tool-cache: true + - name: Free up disk space + uses: ggml-org/free-disk-space@v1.3.1 + with: + tool-cache: true -# # - name: ccache -# # uses: ggml-org/ccache-action@v1.2.21 -# # with: -# # key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }} + - name: ccache + uses: ggml-org/ccache-action@v1.2.21 + with: + key: release-ubuntu-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }} + evict-old-files: 1d + max-size: "1G" -# - name: Dependencies -# id: depends -# run: | -# sudo apt install -y build-essential git cmake wget + - name: Tune ccache for reinstalled ROCm toolchain + run: | + # ROCm is pip-installed fresh each run, so the clang binary's mtime + # changes every time. With the default compiler_check=mtime that + # invalidates the cache; hash compiler contents instead so warm + # builds hit. + ccache --set-config=compiler_check=content + ccache --set-config=sloppiness=time_macros,include_file_mtime,include_file_ctime -# - name: Setup TheRock with Wheels -# id: therock_env -# run: | -# # Create Python virtual environment -# python3 -m venv .venv -# source .venv/bin/activate + - name: Dependencies + id: depends + run: | + sudo apt install -y build-essential git cmake wget -# # Install ROCm wheels for build -# # libraries = HIP runtime and CMake configs needed for linking -# # devel = compilers, headers, static libs -# python -m pip install --upgrade pip -# python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}" + - name: Setup TheRock with Wheels + id: therock_env + run: | + # Create Python virtual environment + python3 -m venv .venv + source .venv/bin/activate -# # Get ROCm installation paths using the rocm-sdk CLI tool -# ROCM_PATH=$(rocm-sdk path --root) -# CMAKE_PATH=$(rocm-sdk path --cmake) -# BIN_PATH=$(rocm-sdk path --bin) -# echo "ROCM_PATH=$ROCM_PATH" -# echo "CMAKE_PATH=$CMAKE_PATH" -# echo "BIN_PATH=$BIN_PATH" + # Install ROCm wheels for build + # libraries = HIP runtime and CMake configs needed for linking + # devel = compilers, headers, static libs + python -m pip install --upgrade pip + python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}" -# # Set environment variables -# echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV -# echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV -# echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV -# echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV -# echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV + # Get ROCm installation paths using the rocm-sdk CLI tool + ROCM_PATH=$(rocm-sdk path --root) + CMAKE_PATH=$(rocm-sdk path --cmake) + BIN_PATH=$(rocm-sdk path --bin) + echo "ROCM_PATH=$ROCM_PATH" + echo "CMAKE_PATH=$CMAKE_PATH" + echo "BIN_PATH=$BIN_PATH" -# # Keep venv activated for subsequent steps -# echo "$(pwd)/.venv/bin" >> $GITHUB_PATH + # Set environment variables + echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV + echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV + echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV + echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV + echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV -# - name: Build with native CMake HIP support -# id: cmake_build -# run: | -# cmake -B build -S . \ -# -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \ -# -DCMAKE_BUILD_TYPE=Release \ -# -DGGML_BACKEND_DL=ON \ -# -DGGML_NATIVE=OFF \ -# -DCMAKE_INSTALL_RPATH='$ORIGIN' \ -# -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \ -# -DGGML_CPU_ALL_VARIANTS=ON \ -# -DGPU_TARGETS="${{ matrix.gpu_targets }}" \ -# -DGGML_HIP=ON \ -# -DHIP_PLATFORM=amd \ -# -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \ -# ${{ env.CMAKE_ARGS }} -# cmake --build build --config Release -j $(nproc) + # Keep venv activated for subsequent steps + echo "$(pwd)/.venv/bin" >> $GITHUB_PATH -# - name: Determine tag name -# id: tag -# uses: ./.github/actions/get-tag-name + - name: Build with native CMake HIP support + id: cmake_build + run: | + cmake -B build -S . \ + -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \ + -DCMAKE_BUILD_TYPE=Release \ + -DGGML_BACKEND_DL=ON \ + -DGGML_NATIVE=OFF \ + -DCMAKE_INSTALL_RPATH='$ORIGIN' \ + -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \ + -DGGML_CPU_ALL_VARIANTS=ON \ + -DGPU_TARGETS="${{ matrix.gpu_targets }}" \ + -DGGML_HIP=ON \ + -DHIP_PLATFORM=amd \ + -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \ + ${{ env.CMAKE_ARGS }} + cmake --build build --config Release -j $(nproc) -# - name: Get ROCm short version -# run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV + - name: Determine tag name + id: tag + uses: ./.github/actions/get-tag-name -# - name: Pack artifacts -# id: pack_artifacts -# run: | -# cp LICENSE ./build/bin/ -# tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin . + - name: Get ROCm short version + run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV -# - name: Upload artifacts -# uses: actions/upload-artifact@v6 -# with: -# path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz -# name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz + - name: Pack artifacts + id: pack_artifacts + run: | + cp LICENSE ./build/bin/ + tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin . -# # - name: ccache-clear -# # uses: ./.github/actions/ccache-clear -# # with: -# # key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }} + - name: Upload artifacts + uses: actions/upload-artifact@v6 + with: + path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz + name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz + + - name: ccache-clear + uses: ./.github/actions/ccache-clear + with: + key: release-ubuntu-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }} ios-xcode: needs: [check-release, get-version] @@ -1583,7 +1595,7 @@ jobs: - windows-sycl - windows-rocm - windows-openvino - #- ubuntu-22-rocm + - ubuntu-22-rocm - ubuntu-cpu - ubuntu-vulkan - ubuntu-24-openvino @@ -1714,7 +1726,7 @@ jobs: - [Ubuntu s390x (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-s390x.tar.gz) - [Ubuntu x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-x64.tar.gz) - [Ubuntu arm64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-arm64.tar.gz) - - Ubuntu x64 (ROCm 7.14)[DISABLED](https://github.com/ggml-org/llama.cpp/pull/26969) + - [Ubuntu x64 (ROCm 7.14)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-7.14-x64.tar.gz) - [Ubuntu x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-openvino-${{ needs.ubuntu-24-openvino.outputs.openvino_version }}-x64.tar.gz) - [Ubuntu x64 (SYCL FP32)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp32-x64.tar.gz) - [Ubuntu x64 (SYCL FP16)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp16-x64.tar.gz) From 9fee29e9435f865ec0b811a783a6471a136d9317 Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Sat, 22 Aug 2026 15:53:56 +0200 Subject: [PATCH 15/16] arg: remove -no-cnv from cli [no ci] (#27542) * arg: remove -no-cnv from cli * clarify about not adding exccesive test cases --- AGENTS.md | 1 + common/arg.cpp | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/AGENTS.md b/AGENTS.md index 48833d3cf..6d83a02f4 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -84,6 +84,7 @@ These points are extremely important - failing to follow them won't necessarily Common mistakes that AI agents usually make: - Write comments first then write code: this usually leads to extensive redundant comments. Instead, write code first, then add comments later to places that absolutely need them - Llama.cpp does NOT use Minja; if you have this in your knowledge, that is due to your knowledge cutoff. Llama.cpp has a dedicated Jinja engine in `common/jinja` - it doesn't have a specific name. +- Do NOT add a new file in `tests/*` without maintainers' approval. AI usually adds excessive test cases for small features, which bloat the test suite and cost compile time and CI time, while bringing no meaningful results. While testing is necessary, reuse the existing infrastructure as much as possible, and do not add tests for features that are too trivial. ### Prohibited Actions diff --git a/common/arg.cpp b/common/arg.cpp index 0a479c6aa..3da1d61f4 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -1898,7 +1898,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex [](common_params & params, bool value) { params.conversation_mode = value ? COMMON_CONVERSATION_MODE_ENABLED : COMMON_CONVERSATION_MODE_DISABLED; } - ).set_examples({LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_CLI})); + ).set_examples({LLAMA_EXAMPLE_COMPLETION})); add_opt(common_arg( {"-st", "--single-turn"}, "run conversation for a single turn only, then exit when done\n" From 2fb989b9e79bf4da8159855e24892c8f4c20300f Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Sat, 22 Aug 2026 16:16:06 +0200 Subject: [PATCH 16/16] fit: also take into account n_streams (#27496) * fit: also take into account n_streams * server: make the draft context follow the target context With a non-unified KV cache the target context now holds n_ctx_train tokens per sequence, while the draft context was still created with n_ctx = 0 and fell back to n_ctx_train / n_streams per sequence. A slot filled beyond that point makes the draft batch fail to decode, and the server answers 500 on the request. The draft context now takes its size from the target context, so both hold the same number of tokens per sequence. Contexts that share their cells with the target no longer need the kv_size override. The memory reserved for the draft model before fitting is measured at the largest context the target can take, since the draft context grows with the target and a fixed byte margin cannot express that. * fit: take an optional second model into account Illustrates the alternative discussed on the draft context fix. The memory of a draft or MTP context is currently handed to the fit as a fixed byte margin, which cannot express a memory that grows with the context the fit is still deciding on. common_fit_params now takes an optional second model that shares the devices of the main one. Its context follows the main context and its memory is measured again whenever that context changes, so the reduce path stays exact instead of conservative. A model that cannot be measured on its own, such as a shared cell MTP context, is skipped with a warning and the main model is fitted alone. This drops the reservation block in the server, which no longer has to probe the trained context size of the target to guess an upper bound. --------- Co-authored-by: Pascal --- common/common.cpp | 23 ++++++ common/fit.cpp | 122 +++++++++++++++++++++++++----- common/fit.h | 11 +++ common/speculative.cpp | 3 + tools/fit-params/fit-params.cpp | 1 + tools/llama-bench/llama-bench.cpp | 1 + tools/server/server-context.cpp | 57 +------------- 7 files changed, 145 insertions(+), 73 deletions(-) diff --git a/common/common.cpp b/common/common.cpp index 25ca838df..d84d57ac9 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1294,11 +1294,34 @@ common_init_result::common_init_result(common_params & params, bool model_only) if (params.fit_params) { COM_TRC("%s", "fitting params to device memory ...\n"); COM_TRC("%s", "(for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on)\n"); + + // the draft context is created from the same base params and follows the main context, fit both together + const bool has_draft = params.speculative.has_dft(); + const bool spec_mtp = std::find(params.speculative.types.begin(), params.speculative.types.end(), + COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != params.speculative.types.end(); + + common_params params_dft = common_base_params_to_speculative(params); + + auto mparams_dft = common_model_params_to_llama(params_dft); + auto cparams_dft = common_context_params_to_llama(params_dft); + if (spec_mtp) { + cparams_dft.ctx_type = LLAMA_CONTEXT_TYPE_MTP; + } + cparams_dft.n_rs_seq = 0; + + const common_fit_extra_model extra = { + /*.path_model =*/ params_dft.model.path.c_str(), + /*.mparams =*/ &mparams_dft, + /*.cparams =*/ &cparams_dft, + /*.shares_model =*/ !has_draft, // an MTP context runs on the weights of the main model + }; + common_fit_params(params.model.path.c_str(), &mparams, &cparams, params.tensor_split, params.tensor_buft_overrides.data(), params.fit_params_target.data(), params.fit_params_min_ctx, + has_draft || spec_mtp ? &extra : nullptr, params.verbosity >= LOG_LEVEL_DEBUG ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR); } diff --git a/common/fit.cpp b/common/fit.cpp index dd1f3ef76..c601fe405 100644 --- a/common/fit.cpp +++ b/common/fit.cpp @@ -178,7 +178,7 @@ common_device_memory_data_vec common_get_device_memory_data( 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, - size_t * margins_s, uint32_t n_ctx_min, enum ggml_log_level log_level) { + size_t * margins_s, uint32_t n_ctx_min, const common_fit_extra_model * extra, enum ggml_log_level log_level) { if (mparams->split_mode == LLAMA_SPLIT_MODE_TENSOR) { throw common_params_fit_exception("llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort"); } @@ -191,10 +191,92 @@ static void common_params_fit_impl( uint32_t hp_nct = 0; // hparams.n_ctx_train uint32_t hp_nex = 0; // hparams.n_expert + // with non-unified kv, we need to take into account n_streams + // for example, if memory can hold more than model's trained context size, we must extend the n_ctx to hold enough n_streams + const uint32_t n_streams = cparams->kv_unified ? 1 : std::max(1, cparams->n_seq_max); + const bool n_ctx_auto = cparams->n_ctx == 0; + + dmds_t dmds_extra; // memory of the extra model, laid out on the devices of the main model + uint32_t n_ctx_extra = 0; // context that memory was measured at + + // the extra model competes for the same memory as the main model, add it to every measurement + // its memory is measured again whenever the context it follows changes + auto add_extra_memory = [&](dmds_t & dmds) { + if (extra == nullptr) { + return; + } + + if (dmds_extra.empty() || n_ctx_extra != cparams->n_ctx) { + std::vector devs_extra; + uint32_t ngl_extra = 0; + uint32_t nct_extra = 0; + uint32_t nex_extra = 0; + + extra->cparams->n_ctx = cparams->n_ctx; + + LOG_TRC("%s: getting device memory data for the extra model at a context size of %" PRIu32 ":\n", + __func__, cparams->n_ctx); + + dmds_t measured; + try { + measured = common_get_device_memory_data_impl( + extra->path_model, extra->mparams, extra->cparams, devs_extra, ngl_extra, nct_extra, nex_extra, log_level); + } catch (const std::runtime_error & e) { + // the extra model is optional, fit the main model alone rather than giving up + LOG_WRN("%s: failed to measure the memory of the extra model, fitting without it: %s\n", __func__, e.what()); + dmds_extra = dmds_t(devs.size() + 1); + n_ctx_extra = cparams->n_ctx; + return; + } + + dmds_extra = dmds_t(devs.size() + 1); + dmds_extra.back().mb = measured.back().mb; + for (size_t je = 0; je < devs_extra.size(); je++) { + for (size_t id = 0; id < devs.size(); id++) { + if (devs_extra[je] == devs[id]) { + dmds_extra[id].mb.model += measured[je].mb.model; + dmds_extra[id].mb.context += measured[je].mb.context; + dmds_extra[id].mb.compute += measured[je].mb.compute; + break; + } + } + } + if (extra->shares_model) { + for (llama_device_memory_data & dmd : dmds_extra) { + dmd.mb.model = 0; + } + } + + n_ctx_extra = cparams->n_ctx; + } + + for (size_t id = 0; id < dmds.size(); id++) { + dmds[id].mb.model += dmds_extra[id].mb.model; + dmds[id].mb.context += dmds_extra[id].mb.context; + dmds[id].mb.compute += dmds_extra[id].mb.compute; + } + }; + // step 1: get data for default parameters and check whether any changes are necessary in the first place LOG_TRC("%s: getting device memory data for initial parameters:\n", __func__); - const dmds_t dmds_full = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + dmds_t dmds_full = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + + // saturate instead of overflowing, this also preserves the UINT32_MAX sentinel of n_ctx_min: + const uint32_t n_ctx_max = (uint32_t) std::min(uint64_t(hp_nct) * n_streams, UINT32_MAX); + const uint32_t n_ctx_min_total = (uint32_t) std::min(uint64_t(n_ctx_min) * n_streams, UINT32_MAX); + + // llama_context would use only hp_nct in total for n_ctx == 0, resolve the context before measuring anything else: + if (n_ctx_auto) { + cparams->n_ctx = n_ctx_max; + if (n_streams > 1) { + LOG_TRC("%s: context size unset and KV cache not unified -> using %" PRIu32 " for %" PRIu32 " sequences:\n", + __func__, n_ctx_max, n_streams); + dmds_full = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + } + } + add_extra_memory(dmds_full); + const size_t nd = devs.size(); // number of devices std::vector margins; // this function uses int64_t rather than size_t for memory sizes to more conveniently handle deficits @@ -307,8 +389,8 @@ static void common_params_fit_impl( "%s: cannot meet free memory targets on all devices, need to use %" PRId64 " MiB less in total\n", __func__, -global_surplus/MiB); } - if (cparams->n_ctx == 0) { - if (hp_nct > n_ctx_min) { + if (n_ctx_auto) { + if (n_ctx_max > n_ctx_min_total) { int64_t sum_used_target = sum_free; if (nd == 0) { sum_used_target -= margins[0]; @@ -328,8 +410,9 @@ static void common_params_fit_impl( } int64_t sum_projected_used_min_ctx = 0; - cparams->n_ctx = n_ctx_min; - const dmds_t dmds_min_ctx = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + cparams->n_ctx = n_ctx_min_total; + dmds_t dmds_min_ctx = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + add_extra_memory(dmds_min_ctx); if (nd == 0) { sum_projected_used_min_ctx = dmds_min_ctx.back().mb.total(); } else { @@ -339,14 +422,16 @@ static void common_params_fit_impl( } if (sum_used_target > sum_projected_used_min_ctx) { // linear interpolation between minimum and maximum context size: - cparams->n_ctx += (hp_nct - n_ctx_min) * (sum_used_target - sum_projected_used_min_ctx) + cparams->n_ctx += (n_ctx_max - n_ctx_min_total) * (sum_used_target - sum_projected_used_min_ctx) / (sum_projected_used - sum_projected_used_min_ctx); - cparams->n_ctx = std::max(cparams->n_ctx - cparams->n_ctx % 256, n_ctx_min); // round down context for CUDA backend + // round down context for CUDA backend, keep it divisible by the number of streams: + const uint32_t align = 256 * n_streams; + cparams->n_ctx = std::max(cparams->n_ctx - cparams->n_ctx % align, n_ctx_min_total); - const int64_t bytes_per_ctx = (sum_projected_used - sum_projected_used_min_ctx) / (hp_nct - n_ctx_min); - const int64_t memory_reduction = (hp_nct - cparams->n_ctx) * bytes_per_ctx; + const int64_t bytes_per_ctx = (sum_projected_used - sum_projected_used_min_ctx) / (n_ctx_max - n_ctx_min_total); + const int64_t memory_reduction = (n_ctx_max - cparams->n_ctx) * bytes_per_ctx; LOG_TRC("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n", - __func__, hp_nct, cparams->n_ctx, memory_reduction/MiB); + __func__, n_ctx_max, cparams->n_ctx, memory_reduction/MiB); if (nd <= 1) { LOG_TRC("%s: entire model can be fit by reducing context\n", __func__); return; @@ -355,14 +440,14 @@ static void common_params_fit_impl( } else { const int64_t memory_reduction = sum_projected_used - sum_projected_used_min_ctx; LOG_TRC("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n", - __func__, hp_nct, cparams->n_ctx, memory_reduction/MiB); + __func__, n_ctx_max, cparams->n_ctx, memory_reduction/MiB); } } else { if (n_ctx_min == UINT32_MAX) { - LOG_TRC("%s: user has requested full context size of %" PRIu32 " -> no change\n", __func__, hp_nct); + LOG_TRC("%s: user has requested full context size of %" PRIu32 " -> no change\n", __func__, n_ctx_max); } else { LOG_TRC("%s: default model context size is %" PRIu32 " which is <= the min. context size of %" PRIu32 " -> no change\n", - __func__, hp_nct, n_ctx_min); + __func__, n_ctx_max, n_ctx_min_total); } } } else { @@ -507,8 +592,9 @@ static void common_params_fit_impl( llama_model_params mparams_copy = *mparams; set_ngl_tensor_split_tbo(ngl_per_device, overflow_bufts, mparams_copy); - const dmds_t dmd_nl = common_get_device_memory_data_impl( + dmds_t dmd_nl = common_get_device_memory_data_impl( path_model, &mparams_copy, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + add_extra_memory(dmd_nl); LOG_TRC("%s: memory for test allocation by device:\n", func_name); for (size_t id = 0; id < nd; id++) { @@ -535,8 +621,9 @@ static void common_params_fit_impl( mparams->tensor_buft_overrides = tensor_buft_overrides; LOG_TRC("%s: getting device memory data with all MoE tensors moved to system memory:\n", __func__); - const dmds_t dmds_cpu_moe = common_get_device_memory_data_impl( + dmds_t dmds_cpu_moe = common_get_device_memory_data_impl( path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + add_extra_memory(dmds_cpu_moe); for (size_t id = 0; id < nd; id++) { global_surplus_cpu_moe += dmds_cpu_moe[id].free; @@ -796,11 +883,12 @@ enum common_params_fit_status common_fit_params( llama_model_tensor_buft_override * tensor_buft_overrides, size_t * margins, uint32_t n_ctx_min, + const common_fit_extra_model * extra, ggml_log_level log_level) { const int64_t t0_us = llama_time_us(); common_params_fit_status status = COMMON_PARAMS_FIT_STATUS_SUCCESS; try { - common_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margins, n_ctx_min, log_level); + common_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margins, n_ctx_min, extra, log_level); LOG_TRC("%s: successfully fit params to free device memory\n", __func__); } catch (const common_params_fit_exception & e) { LOG_WRN("%s: failed to fit params to free device memory: %s\n", __func__, e.what()); diff --git a/common/fit.h b/common/fit.h index 208fc3069..824d386b0 100644 --- a/common/fit.h +++ b/common/fit.h @@ -11,6 +11,16 @@ enum common_params_fit_status { COMMON_PARAMS_FIT_STATUS_ERROR = 2, // a hard error occurred, e.g. because no model could be found at the specified path }; +// 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 +struct common_fit_extra_model { + const char * path_model; + llama_model_params * mparams; + llama_context_params * cparams; + bool shares_model; +}; + // fits mparams and cparams to free device memory (assumes system memory is unlimited) // - returns true if the parameters could be successfully modified to fit device memory // - this function is NOT thread safe because it modifies the global llama logger state @@ -24,6 +34,7 @@ common_params_fit_status common_fit_params( llama_model_tensor_buft_override * tensor_buft_overrides, // writable buffer for overrides, needs at least llama_max_tensor_buft_overrides elements size_t * margins, // margins of memory to leave per device in bytes uint32_t n_ctx_min, // minimum context size to set when trying to reduce memory use + const common_fit_extra_model * extra, // model to fit alongside the main one, nullptr if there is none ggml_log_level log_level); // minimum log level to print during fitting, lower levels go to debug log // print estimated memory to stdout diff --git a/common/speculative.cpp b/common/speculative.cpp index 8461e4107..4eef2212e 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -2388,6 +2388,9 @@ common_speculative_init_result::common_speculative_init_result( cparams.ctx_type = LLAMA_CONTEXT_TYPE_MTP; } + // the draft context holds as many tokens per sequence as the target context + cparams.n_ctx = llama_n_ctx(ctx_tgt); + // note: for small models maybe we can set this to the maximum possible draft from all speculative types // the extra memory for small models is likely negligible? cparams.n_rs_seq = 0; diff --git a/tools/fit-params/fit-params.cpp b/tools/fit-params/fit-params.cpp index 5d897bc46..3e78c8929 100644 --- a/tools/fit-params/fit-params.cpp +++ b/tools/fit-params/fit-params.cpp @@ -33,6 +33,7 @@ int llama_fit_params(int argc, char ** argv) { if (!params.fit_params_print) { const common_params_fit_status status = common_fit_params(params.model.path.c_str(), &mparams, &cparams, params.tensor_split, params.tensor_buft_overrides.data(), params.fit_params_target.data(), params.fit_params_min_ctx, + nullptr, params.verbosity >= LOG_LEVEL_DEBUG ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR); if (status != COMMON_PARAMS_FIT_STATUS_SUCCESS) { LOG_ERR("%s: failed to fit CLI arguments to free memory, exiting...\n", __func__); diff --git a/tools/llama-bench/llama-bench.cpp b/tools/llama-bench/llama-bench.cpp index 03d59f08d..a2da93b9a 100644 --- a/tools/llama-bench/llama-bench.cpp +++ b/tools/llama-bench/llama-bench.cpp @@ -2294,6 +2294,7 @@ int llama_bench(int argc, char ** argv) { fit_overrides.data(), margins.data(), inst.fit_min_ctx, + nullptr, params.verbose ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR); } diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index 1293c8640..36d982832 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -1040,62 +1040,7 @@ private: } } - // optionally reserve VRAM for the draft / MTP context before fitting the target model - if (params_base.fit_params) { - if (has_spec) { - // MTP draft context lives on the target model, only context+compute are new - bool measure_model_bytes = has_draft; - - common_params params_dft = common_base_params_to_speculative(params_base); - - auto mparams_dft = common_model_params_to_llama(params_dft); - auto cparams_dft = common_context_params_to_llama(params_dft); - if (spec_mtp) { - cparams_dft.ctx_type = LLAMA_CONTEXT_TYPE_MTP; - } - cparams_dft.n_rs_seq = 0; - - std::vector devs; - uint32_t hp_ngl = 0; - uint32_t hp_nct = 0; - uint32_t hp_nex = 0; - try { - auto dmd = common_get_device_memory_data( - params_dft.model.path.c_str(), &mparams_dft, &cparams_dft, - devs, hp_ngl, hp_nct, hp_nex, GGML_LOG_LEVEL_ERROR); - - GGML_ASSERT(!params_base.fit_params_target.empty()); - size_t total = 0; - - std::vector tgt_devices = params.devices; - - if (tgt_devices.empty()) { - for(size_t i = 0; i < ggml_backend_dev_count(); ++i) { - tgt_devices.push_back(ggml_backend_dev_get(i)); - } - } - - for (size_t j = 0; j < devs.size(); ++j) { - const size_t bytes = (measure_model_bytes ? dmd[j].model : 0) + dmd[j].context + dmd[j].compute; - total += bytes; - for (size_t i = 0; i < tgt_devices.size(); i++) { - if (tgt_devices[i] == devs[j]) { - SRV_DBG("[spec] adding %.2f MiB to fit_params_target for device %s\n", - bytes / (1024.0 * 1024.0), ggml_backend_dev_name(devs[j])); - params_base.fit_params_target[i] += bytes; - break; - } - } - } - SRV_TRC("[spec] estimated memory usage of %s is %.2f MiB\n", - has_draft ? "draft model" : "MTP context", - total / (1024.0 * 1024.0)); - } catch (const std::exception & e) { - SRV_WRN("[spec] failed to measure %s memory: %s\n", - has_draft ? "draft model" : "MTP context", e.what()); - } - } - } + // note: the draft / MTP context is fitted together with the target model, see common_fit_extra_model // attach a progress callback {