Merge branch 'upstream' into concedo_experimental

# Conflicts:
#	.github/workflows/build-cmake-pkg.yml
#	.github/workflows/build-cpu.yml
#	.github/workflows/server-sanitize.yml
#	CMakeLists.txt
#	docs/ops.md
#	docs/ops/CUDA.csv
#	docs/ops/SYCL.csv
#	examples/model-conversion/Makefile
#	ggml/CMakeLists.txt
#	ggml/include/ggml-sycl.h
#	ggml/src/ggml-hexagon/ggml-hexagon.cpp
#	ggml/src/ggml-hexagon/htp/dma-queue.h
#	ggml/src/ggml-hexagon/htp/flash-attn-ops.c
#	ggml/src/ggml-opencl/ggml-opencl.cpp
#	ggml/src/ggml-sycl/fusion.cpp
#	ggml/src/ggml-sycl/getrows.cpp
#	ggml/src/ggml-sycl/getrows.hpp
#	ggml/src/ggml-sycl/ggml-sycl.cpp
#	ggml/src/ggml-sycl/mmvq.cpp
#	ggml/src/ggml-sycl/mmvq.hpp
#	ggml/src/ggml-sycl/norm.cpp
#	ggml/src/ggml-sycl/norm.hpp
#	ggml/src/ggml-vulkan/CMakeLists.txt
#	scripts/make-release-desc.sh
#	scripts/make-release-summary.txt
#	scripts/sync-ggml.last
#	tests/test-backend-ops.cpp
#	tests/test-jinja.cpp
#	tests/test-llama-archs.cpp
#	tests/test-save-load-state.cpp
#	tools/mtmd/CMakeLists.txt

common/hf-cache.cpp

@@ -30,8 +30,8 @@ namespace hf_cache {
 
 namespace fs = std::filesystem;
 
-std::string get_cache_path() {
+static fs::path get_cache_directory() {
-    static const std::string cache = []() {
+    static const fs::path cache = []() {
         struct {
             const char * var;
             fs::path path;
@@ -46,14 +46,14 @@ std::string get_cache_path() {
         for (const auto & entry : entries) {
             if (auto * p = std::getenv(entry.var); p && *p) {
                 fs::path base(p);
-                return (entry.path.empty() ? base : base / entry.path).string();
+                return entry.path.empty() ? base : base / entry.path;
             }
         }
 #ifndef _WIN32
         const struct passwd * pw = getpwuid(getuid());
 
         if (pw && pw->pw_dir && *pw->pw_dir) {
-            return (fs::path(pw->pw_dir) / ".cache" / "huggingface" / "hub").string();
+            return fs::path(pw->pw_dir) / ".cache" / "huggingface" / "hub";
         }
 #endif
         throw std::runtime_error("Failed to determine HF cache directory");
@@ -62,6 +62,15 @@ std::string get_cache_path() {
     return cache;
 }
 
+std::string get_cache_path() {
+#if defined(__cpp_lib_char8_t)
+    const std::u8string u8str = get_cache_directory().u8string();
+    return std::string(reinterpret_cast<const char *>(u8str.data()), u8str.size());
+#else
+    return get_cache_directory().u8string();
+#endif
+}
+
 static std::string folder_name_to_repo(const std::string & folder) {
     constexpr std::string_view prefix = "models--";
     if (folder.rfind(prefix, 0)) {
@@ -80,7 +89,7 @@ static std::string repo_to_folder_name(const std::string & repo_id) {
 }
 
 static fs::path get_repo_path(const std::string & repo_id) {
-    return fs::path(get_cache_path()) / repo_to_folder_name(repo_id);
+    return get_cache_directory() / repo_to_folder_name(repo_id);
 }
 
 static bool is_hex_char(const char c) {
@@ -393,7 +402,7 @@ static std::string get_cached_ref(const fs::path & repo_path) {
 }
 
 hf_files get_cached_files(const std::string & repo_id) {
-    const fs::path cache_path = get_cache_path();
+    const fs::path cache_path = get_cache_directory();
     if (!fs::exists(cache_path)) {
         return {};
     }

common/jinja/parser.cpp

@@ -440,7 +440,8 @@ private:
     }
 
     statement_ptr parse_filter_expression() {
-        auto operand = parse_call_member_expression();
+        // Filters/tests bind outside unary so -n|abs is (-n)|abs, not -(n|abs).
+        auto operand = parse_unary_expression();
         while (is(token::pipe)) {
             size_t start_pos = current;
             ++current; // consume pipe
@@ -451,6 +452,15 @@ private:
         return operand;
     }
 
+    statement_ptr parse_unary_expression() {
+        if (is(token::unary_operator)) {
+            size_t start_pos = current;
+            auto op = next();
+            return mk_stmt<unary_expression>(start_pos, op, parse_unary_expression());
+        }
+        return parse_call_member_expression();
+    }
+
     statement_ptr parse_call_member_expression() {
         // Handle member expressions recursively
         auto member = parse_member_expression(parse_primary_expression());

common/jinja/runtime.cpp

@@ -453,6 +453,11 @@ value unary_expression::execute_impl(context & ctx) const {
         } else {
             throw std::runtime_error("Unary - operator requires numeric operand");
         }
+    } else if (op.value == "+") {
+        if (is_val<value_int>(operand_val) || is_val<value_float>(operand_val)) {
+            return operand_val;
+        }
+        throw std::runtime_error("Unary + operator requires numeric operand");
     }
 
     throw std::runtime_error("Unknown unary operator '" + op.value + "'");
@@ -485,14 +490,8 @@ value for_statement::execute_impl(context & ctx) const {
     const jinja::select_expression * select_expr = cast_stmt<select_expression>(iterable);
     statement_ptr test_expr_nullptr;
 
-    const statement_ptr & iter_expr = [&]() -> const statement_ptr & {
+    const statement_ptr & iter_expr = select_expr ? select_expr->lhs : iterable;
-        auto tmp = cast_stmt<select_expression>(iterable);
+    const statement_ptr & test_expr = select_expr ? select_expr->test : test_expr_nullptr;
-        return tmp ? tmp->lhs : iterable;
-    }();
-    const statement_ptr & test_expr = [&]() -> const statement_ptr & {
-        auto tmp = cast_stmt<select_expression>(iterable);
-        return tmp ? tmp->test : test_expr_nullptr;
-    }();
 
     JJ_DEBUG("Executing for statement, iterable type: %s", iter_expr->type().c_str());
 

common/speculative.cpp

@@ -228,9 +228,7 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
             common_params_sampling params;
             params.no_perf = false;
             params.top_k = 10;
-            params.samplers = {
+            params.samplers.assign(1, COMMON_SAMPLER_TYPE_TOP_K);
-                COMMON_SAMPLER_TYPE_TOP_K,
-            };
 
             smpl.reset(common_sampler_init(llama_get_model(ctx_dft), params));
         }

conversion/__init__.py

@@ -28,6 +28,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
     "BailingMoeForCausalLM": "bailingmoe",
     "BailingMoeV2ForCausalLM": "bailingmoe",
     "BailingMoeV3ForCausalLM": "bailingmoe3",
+    "BailingMoeV3VLForConditionalGeneration": "bailingmoe3",
     "BambaForCausalLM": "granite",
     "BertForMaskedLM": "bert",
     "BertForSequenceClassification": "bert",
@@ -94,6 +95,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
     "Gemma3nForCausalLM": "gemma",
     "Gemma3nForConditionalGeneration": "gemma",
     "Gemma4AssistantForCausalLM": "gemma",
+    "Gemma4DSparkModel": "gemma",
     "Gemma4ForConditionalGeneration": "gemma",
     "Gemma4ForCausalLM": "gemma",
     "Gemma4UnifiedForConditionalGeneration": "gemma",
@@ -301,6 +303,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
     "Gemma4ForConditionalGeneration": "gemma",
     "Gemma4UnifiedForConditionalGeneration": "gemma",
     "Glm4vForConditionalGeneration": "qwen3vl",
+    "BailingMoeV3VLForConditionalGeneration": "bailingmoe3",
     "Glm4vMoeForConditionalGeneration": "qwen3vl",
     "Glm5vForConditionalGeneration": "kimivl",
     "GlmOcrForConditionalGeneration": "qwen3vl",

conversion/bailingmoe3.py

@@ -9,7 +9,9 @@ import torch
 if TYPE_CHECKING:
     from torch import Tensor
 
-from .base import ModelBase, TextModel, gguf
+from .base import ModelBase, MmprojModel, TextModel, gguf
+
+from .qwen3vl import Qwen3VLVisionModel
 
 
 @ModelBase.register("BailingMoeV3ForCausalLM")
@@ -74,7 +76,7 @@ class BailingMoeV3Model(TextModel):
 
         self.gguf_writer.add_expert_feed_forward_length(self.hparams["moe_intermediate_size"])
         self.gguf_writer.add_expert_shared_feed_forward_length(self.hparams["moe_shared_expert_intermediate_size"])
-        self.gguf_writer.add_expert_shared_count(self.hparams["num_shared_experts"])
+        self.gguf_writer.add_expert_shared_count(self.hparams.get("num_shared_experts", 1))
         self.gguf_writer.add_leading_dense_block_count(self.hparams["first_k_dense_replace"])
         self.gguf_writer.add_expert_weights_scale(self.hparams["routed_scaling_factor"])
         self.gguf_writer.add_expert_weights_norm(self.hparams["norm_topk_prob"])
@@ -191,3 +193,111 @@ class BailingMoeV3Model(TextModel):
             experts = [name for layer in self._experts for name in layer]
             if experts:
                 raise ValueError(f"Unprocessed experts: {experts}")
+
+
+@ModelBase.register("BailingMoeV3VLForConditionalGeneration")
+@ModelBase.example("inclusionAI/Ling-3.0-flash-VL")
+class BailingMoeV3VLModel(BailingMoeV3Model):
+    model_arch = gguf.MODEL_ARCH.BAILINGMOE3
+
+    def index_tensors(self, remote_hf_model_id: str | None = None):
+        # hoist text_config before the shared BailingMoeV3 logic runs:
+        # ModelBase.__init__ calls this with the raw VL config, where the text
+        # dims still live under text_config
+        if "text_config" in self.hparams:
+            self.hparams = {**self.hparams, **self.hparams["text_config"]}
+        return super().index_tensors(remote_hf_model_id=remote_hf_model_id)
+
+    def set_gguf_parameters(self):
+        super().set_gguf_parameters()
+        mrope_section = self.hparams.get("mrope_section")
+        if mrope_section is None:
+            raise ValueError("BailingMoeV3VL requires mrope_section in the config")
+        if sum(mrope_section[:3]) * 2 != self.hparams["qk_rope_head_dim"]:
+            raise ValueError(
+                f"mrope_section {mrope_section[:3]} counts rope pairs and must sum to"
+                f" qk_rope_head_dim / 2 = {self.hparams['qk_rope_head_dim'] // 2}"
+            )
+        # mrope_section is [t, h, w]; pad to the 4-wide sections array
+        self.gguf_writer.add_rope_dimension_sections(list(mrope_section[:3]) + [0])
+
+    @classmethod
+    def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
+        name, gen = item
+
+        # Skip projector tensors; the vision tower is skipped by TextModel.filter_tensors
+        if name.startswith("linear_proj"):
+            return None
+
+        return super().filter_tensors(item)
+
+
+@ModelBase.register("BailingMoeV3VLForConditionalGeneration")
+@ModelBase.example("inclusionAI/Ling-3.0-flash-VL")
+class BailingMoeV3VLVisionModel(Qwen3VLVisionModel):
+    def __init__(self, *args, **kwargs):
+        super().__init__(*args, **kwargs)
+        assert self.hparams_vision is not None
+
+        if self.hparams_vision.get("disable_merger_proj") is not True:
+            raise ValueError("BailingMoeV3VL requires disable_merger_proj=true")
+
+        # out_hidden_size is the vision encoder output (post spatial merge, pre linear_proj)
+        self.image_emb_dim = self.hparams_vision.get("out_hidden_size")
+        if self.image_emb_dim is None:
+            raise ValueError("BailingMoeV3VL vision config requires out_hidden_size")
+
+    def set_gguf_parameters(self):
+        assert self.hparams_vision is not None
+        MmprojModel.set_gguf_parameters(self)  # skip Qwen3VLVisionModel parameters
+        self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.LING3VL)
+        self.gguf_writer.add_vision_use_gelu(True)
+
+        merge_size = self.hparams_vision.get("spatial_merge_size")
+        if merge_size is not None:
+            self.gguf_writer.add_vision_spatial_merge_size(int(merge_size))
+
+        rms_norm_eps = self.global_config.get("text_config", {}).get("rms_norm_eps", 1e-6)
+        self.gguf_writer.add_vision_attention_layernorm_eps(rms_norm_eps)
+
+    @classmethod
+    def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
+        name, gen = item
+
+        if name.startswith("lm_head."):
+            return None
+
+        if name.startswith("linear_proj"):
+            # top-level projector MLP: linear_proj.0 -> mm.0, linear_proj.2 -> mm.2
+            parts = name.split(".")
+            if len(parts) != 3:
+                raise ValueError(f"Unexpected linear_proj tensor: {name}")
+            idx, suffix = int(parts[1]), parts[2]
+            name = f"mm.{idx}.{suffix}"
+            # the qwen3vl filter keeps only visual.*; skip it for the renamed projector tensors
+            return MmprojModel.filter_tensors((name, gen))
+
+        if name.startswith("model.visual."):
+            name = name.replace("model.visual.", "visual.", 1)
+
+        if not name.startswith("visual."):
+            return None
+
+        return super().filter_tensors((name, gen))
+
+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        assert self.hparams_vision is not None
+
+        if name.startswith("mm.0.") or name.startswith("mm.2."):
+            # top-level projector MLP (linear_proj.0 / linear_proj.2, renamed by filter_tensors)
+            yield (name, data_torch)
+            return
+
+        if name == "visual.merger.norm.weight" or name == "visual.merger.norm.bias":
+            # the merger is norm-only for Ling: per-patch LayerNorm before the spatial merge
+            new_name = f"mm.input_norm.{name.split('.')[-1]}"
+            yield (new_name, data_torch)
+            return
+
+        # Ling has no patch bias; the Conv3D split below matches the stock qwen3vl path
+        yield from Qwen3VLVisionModel.modify_tensors(self, data_torch, name, bid)

conversion/gemma.py

@@ -11,6 +11,7 @@ if TYPE_CHECKING:
     from torch import Tensor
 
 from .base import MmprojModel, ModelBase, TextModel, gguf, logger
+from .qwen import DFlashModel
 
 
 @ModelBase.register("GemmaForCausalLM")
@@ -809,6 +810,105 @@ class Gemma4Model(Gemma3Model):
         yield from super().modify_tensors(data_torch, name, bid)
 
 
+@ModelBase.register("Gemma4DSparkModel")
+class Gemma4DSparkModel(DFlashModel):
+    model_arch = gguf.MODEL_ARCH.DFLASH
+
+    def __init__(self, *args, **kwargs):
+        super().__init__(*args, **kwargs)
+
+        if not self.hparams.get("attention_k_eq_v", False):
+            raise ValueError("Gemma4 DSpark currently requires attention_k_eq_v")
+        if self.hparams.get("layer_types") != ["full_attention"] * self.block_count:
+            raise ValueError("Gemma4 DSpark currently requires uniform full_attention layer types")
+        if self.hparams.get("hidden_activation", "gelu_pytorch_tanh") != "gelu_pytorch_tanh":
+            raise ValueError("Gemma4 DSpark currently requires hidden_activation=gelu_pytorch_tanh")
+        if self.hparams.get("attention_bias", False) or self.hparams.get("enable_moe_block", False):
+            raise ValueError("Gemma4 DSpark attention bias and MoE are not supported")
+        if (self.hparams.get("draft_vocab_size") or self.hparams["vocab_size"]) != self.hparams["vocab_size"]:
+            raise ValueError("Gemma4 DSpark currently requires a full draft vocabulary")
+        if "model.lm_head.weight" not in self.model_tensors and self.hparams.get("tie_word_embeddings") is not True:
+            raise ValueError("Gemma4 DSpark requires lm_head.weight unless tie_word_embeddings is true")
+
+        self.dflash_config = self.hparams.get("dflash_config", {})
+        markov_type = self.dflash_config.get("markov_head_type", self.hparams.get("markov_head_type", "vanilla"))
+        if markov_type != "vanilla":
+            raise ValueError("Gemma4 DSpark currently requires a vanilla Markov head")
+
+        # Gemma4TextConfig supplies these defaults when rope_parameters is absent.
+        rope = self.hparams.get("rope_parameters") or {
+            "full_attention": {"rope_type": "proportional", "partial_rotary_factor": 0.25, "rope_theta": 1000000.0},
+        }
+        self.rope_parameters = rope.get("full_attention", rope)
+        if self.rope_parameters.get("rope_type") not in ("default", "proportional"):
+            raise ValueError("Gemma4 DSpark requires default or proportional RoPE")
+
+    def set_vocab(self):
+        super().set_vocab()
+        mask_id = self.dflash_config.get("mask_token_id", self.hparams.get("mask_token_id"))
+        if mask_id is None:
+            raise ValueError("Gemma4 DSpark requires mask_token_id")
+        if "mask_token_id" not in self.dflash_config:
+            self.gguf_writer.add_mask_token_id(mask_id)
+
+    def set_gguf_parameters(self):
+        super().set_gguf_parameters()
+        head_dim = int(self.hparams["global_head_dim"])
+        self.gguf_writer.add_head_count_kv(self.hparams["num_global_key_value_heads"])
+        self.gguf_writer.add_key_length(head_dim)
+        self.gguf_writer.add_value_length(head_dim)
+        self.gguf_writer.add_rope_dimension_count(head_dim)
+        self.gguf_writer.add_embedding_scale(self.hparams["hidden_size"] ** 0.5)
+        self.gguf_writer.add_attention_scale(1.0)
+        self.gguf_writer.add_hidden_act("gelu_pytorch_tanh")
+
+        self.gguf_writer.add_sample_from_anchor(self.hparams.get("sample_from_anchor", True))
+        target_layers = self.dflash_config.get("target_layer_ids", self.hparams.get("target_layer_ids"))
+        if not target_layers:
+            raise ValueError("Gemma4 DSpark requires target_layer_ids")
+        self.gguf_writer.add_has_confidence_head(any("confidence_head.proj" in name for name in self.model_tensors))
+
+        if self.hparams.get("final_logit_softcapping"):
+            raise ValueError("Gemma4 DSpark logit softcapping is not supported")
+        # The top-level sliding_window is inert unless the draft enables SWA.
+        if self.dflash_config.get("use_swa", False):
+            window = self.dflash_config["swa_window_size"]
+            if window <= 0:
+                raise ValueError("Gemma4 DSpark swa_window_size must be positive")
+
+    @classmethod
+    def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
+        name, gen = item
+        if not name.startswith("model."):
+            name = "model." + name
+        if name.endswith(".layer_scalar"):
+            name += ".weight"
+        name = name.replace("model.confidence_proj.", "model.confidence_head.proj.")
+        return super().filter_tensors((name, gen))
+
+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        # The shared DFlash map assigns this name to Qwen's pre-FFN norm.
+        if name.endswith(".post_attention_layernorm.weight"):
+            name = self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_POST_NORM, bid)
+        elif name.endswith(".pre_feedforward_layernorm.weight"):
+            name = self.format_tensor_name(gguf.MODEL_TENSOR.FFN_NORM, bid)
+        yield from super().modify_tensors(data_torch, name, bid)
+
+    def generate_extra_tensors(self) -> Iterable[tuple[str, Tensor]]:
+        if self.rope_parameters["rope_type"] == "proportional":
+            # Keep the unrotated dimensions in place, as in the Gemma4 converter.
+            head_dim = int(self.hparams["global_head_dim"])
+            fraction_value = self.rope_parameters.get("partial_rotary_factor", 0.25)
+            if not isinstance(fraction_value, (int, float)):
+                raise ValueError("Gemma4 DSpark partial_rotary_factor must be numeric")
+            fraction = float(fraction_value)
+            n_rot = int(head_dim * fraction / 2)
+            if not 0 < fraction <= 1 or head_dim * fraction != 2 * n_rot:
+                raise ValueError("Gemma4 DSpark rotary dimension count must be positive and even")
+            factors = torch.tensor([1.0] * n_rot + [1e30] * (head_dim // 2 - n_rot), dtype=torch.float32)
+            yield self.format_tensor_name(gguf.MODEL_TENSOR.ROPE_FREQS), factors
+
+
 @ModelBase.register("Gemma4UnifiedForConditionalGeneration")
 @ModelBase.example("hf-tiny-v2/tiny-random-Gemma4UnifiedForConditionalGeneration")
 class Gemma4UnifiedModel(Gemma4Model):

conversion/qwen.py

@@ -10,7 +10,7 @@ import torch
 if TYPE_CHECKING:
     from torch import Tensor
 
-from .base import LazyTorchTensor, ModelBase, TextModel, gguf, logger
+from .base import LazyTorchTensor, ModelBase, ModelType, TextModel, get_model_architecture, gguf, logger
 
 
 @ModelBase.register("QWenLMHeadModel")
@@ -666,7 +666,7 @@ class DFlashModel(Qwen3Model):
         from . import get_model_class
         with open(self.target_model_dir / "config.json", "r", encoding="utf-8") as f:
             target_hparams = json.load(f)
-            target_arch = target_hparams["architectures"][0]
+        target_arch = get_model_architecture(target_hparams, ModelType.TEXT)
         target_cls = get_model_class(target_arch)
 
         if target_cls is not type(self):
@@ -711,7 +711,7 @@ class DFlashModel(Qwen3Model):
         if embedding_scale is not None:
             self.gguf_writer.add_embedding_scale(float(embedding_scale))
 
-        target_layer_ids = dflash_config.get("target_layer_ids", [])
+        target_layer_ids = dflash_config.get("target_layer_ids", self.hparams.get("target_layer_ids", []))
         if target_layer_ids:
             extract_layer_ids = [i + 1 for i in target_layer_ids]
             self.gguf_writer.add_target_layers(extract_layer_ids)
@@ -719,8 +719,9 @@ class DFlashModel(Qwen3Model):
         use_sliding_window = self.hparams.get("use_sliding_window", False) or dflash_config.get("use_swa", False)
         sliding_window = dflash_config.get("swa_window_size") or self.hparams.get("sliding_window")
         layer_types = self.hparams.get("layer_types")
-        if use_sliding_window and sliding_window and layer_types:
+        if use_sliding_window and sliding_window:
-            is_swa = [lt == "sliding_attention" for lt in layer_types]
+            is_swa = ([True] * self.block_count if dflash_config.get("use_swa", False)
+                      else [lt == "sliding_attention" for lt in layer_types or []])
             self.gguf_writer.add_sliding_window(sliding_window)
             self.gguf_writer.add_sliding_window_pattern(is_swa)
 
@@ -840,13 +841,6 @@ class DSparkModel(DFlashModel):
             return None
         return super().filter_tensors(item)
 
-    _ROPE_PERMUTE_SUFFIXES = (
-        "self_attn.q_proj.weight",
-        "self_attn.k_proj.weight",
-        "self_attn.q_norm.weight",
-        "self_attn.k_norm.weight",
-    )
-
     def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
         if name == "model.d2t":
             self._d2t = data_torch
@@ -855,12 +849,6 @@ class DSparkModel(DFlashModel):
         if self._n_vocab_draft == self.hparams["vocab_size"] and name.endswith("lm_head.weight"):
             return
 
-        # interleaved-rope checkpoints (rope_is_neox_style = false) -> NeoX layout: per head, even dims first then odd
-        if not self.hparams.get("rope_is_neox_style", True) and name.endswith(self._ROPE_PERMUTE_SUFFIXES):
-            head_dim = self.hparams["head_dim"]
-            shape = data_torch.shape
-            data_torch = data_torch.reshape(-1, head_dim // 2, 2, *shape[1:]).transpose(1, 2).reshape(shape)
-
         yield from super().modify_tensors(data_torch, name, bid)
 
     def prepare_tensors(self):

ggml/src/ggml-backend-meta.cpp

@@ -1200,11 +1200,6 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor_impl(ggml_backend_m
         ggml_context          * simple_ctx = stc.ctxs[j].get();
         ggml_backend_buffer_t   simple_buf = buf_ctx->bufs[j].get();
 
-        if ((simple_buf != nullptr) && ggml_backend_buffer_is_multi_buffer(simple_buf)) {
-            // see https://github.com/ggml-org/llama.cpp/issues/22197
-            GGML_ABORT("multi buffers are not supported by the meta backend");
-        }
-
         if (split_dim >= 0 && split_dim < GGML_MAX_DIMS) {
             // TODO: the following assert fails for llama-parallel even though the results are correct:
             // GGML_ASSERT(ggml_is_contiguously_allocated(tensor));
@@ -1252,16 +1247,27 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor_impl(ggml_backend_m
                 }
             }
         }
+        // TODO: revisit once the graph allocator has been refactored, see https://github.com/ggml-org/llama.cpp/pull/25051#issuecomment-4842873396
+        ggml_backend_buffer_t init_buf = simple_buf;
         if (t_ij->view_src != nullptr) {
             t_ij->data = (char *) t_ij->view_src->data + t_ij->view_offs;
+            // views inherit the source slice's concrete sub-buffer (issue 22197)
+            if (tensor->view_src != nullptr && ggml_backend_buffer_is_meta(tensor->view_src->buffer)
+                    && t_ij->view_src->buffer != nullptr) {
+                t_ij->buffer = t_ij->view_src->buffer;
+                init_buf     = t_ij->view_src->buffer;
+            }
         } else if (simple_buf != nullptr) {
+            if (ggml_backend_buffer_is_multi_buffer(simple_buf)) {
+                GGML_ABORT("multi buffers are not supported by the meta backend");
+            }
             t_ij->data = (char *) ggml_backend_buffer_get_base(simple_buf)
                 + size_t(tensor->data) - size_t(ggml_backend_buffer_get_base(tensor->buffer));
         }
 
-        if (simple_buf) {
+        if (init_buf) {
             // the backend that owns the buffer will set .extra
-            ggml_backend_buffer_init_tensor(simple_buf, t_ij);
+            ggml_backend_buffer_init_tensor(init_buf, t_ij);
         } else {
             t_ij->extra = tensor->extra;
         }

ggml/src/ggml-cuda/conv3d.cu

@@ -0,0 +1,362 @@
+#include "conv3d.cuh"
+#include "convert.cuh"
+#include "mma.cuh"
+
+struct conv3d_params {
+    int64_t IW, IH, ID;
+    int64_t OW, OH, OD;
+    int64_t KW, KH, KD;
+    int64_t ST_X, ST_Y, ST_Z;
+    int64_t PD_X, PD_Y, PD_Z;
+    int64_t DL_X, DL_Y, DL_Z;
+    int64_t IC, OC, B, TOTAL;
+};
+
+template <typename T>
+static __global__ void conv3d_kernel(const float * input, const T * weight, float * output, const conv3d_params P) {
+    const int64_t spatial = P.OW * P.OH * P.OD;
+    for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x; i < P.TOTAL;
+         i += int64_t(gridDim.x) * blockDim.x) {
+        const int64_t x = i % P.OW, y = i / P.OW % P.OH, z = i / (P.OW * P.OH) % P.OD;
+        const int64_t co = i / spatial % P.OC, n = i / (spatial * P.OC);
+        float         sum = 0.0f;
+        for (int64_t ci = 0; ci < P.IC; ++ci) {
+            for (int64_t kz = 0; kz < P.KD; ++kz) {
+                const int64_t iz = z * P.ST_Z + kz * P.DL_Z - P.PD_Z;
+                if (iz < 0 || iz >= P.ID) {
+                    continue;
+                }
+                for (int64_t ky = 0; ky < P.KH; ++ky) {
+                    const int64_t iy = y * P.ST_Y + ky * P.DL_Y - P.PD_Y;
+                    if (iy < 0 || iy >= P.IH) {
+                        continue;
+                    }
+                    for (int64_t kx = 0; kx < P.KW; ++kx) {
+                        const int64_t ix = x * P.ST_X + kx * P.DL_X - P.PD_X;
+                        if (ix >= 0 && ix < P.IW) {
+                            const int64_t xi = (((n * P.IC + ci) * P.ID + iz) * P.IH + iy) * P.IW + ix;
+                            const int64_t wi = (((co * P.IC + ci) * P.KD + kz) * P.KH + ky) * P.KW + kx;
+                            sum += input[xi] * ggml_cuda_cast<float>(weight[wi]);
+                        }
+                    }
+                }
+            }
+        }
+        output[i] = sum;
+    }
+}
+
+static __global__ void conv3d_pad_f16(const float * input,
+                                      half *        output,
+                                      int           iw,
+                                      int           ih,
+                                      int           id,
+                                      int           pw,
+                                      int           ph,
+                                      int           pd,
+                                      int           px,
+                                      int           py,
+                                      int           pz,
+                                      int           total) {
+    const int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= total) {
+        return;
+    }
+    const int x = i % pw - px, y = i / pw % ph - py, z = i / (pw * ph) % pd - pz;
+    const int nc = i / (pw * ph * pd);
+    output[i] =
+        __float2half((unsigned) x < (unsigned) iw && (unsigned) y < (unsigned) ih && (unsigned) z < (unsigned) id ?
+                         input[((nc * id + z) * ih + y) * iw + x] :
+                         0.0f);
+}
+
+template <int KW, int KH, int KD, bool use_mma>
+static __global__ void conv3d_implicit_gemm_f16(const half * __restrict__ input,
+                                                const half * __restrict__ weight,
+                                                float * __restrict__ output,
+                                                const conv3d_params P,
+                                                const int           split_k,
+                                                const bool          aligned_weights) {
+    using namespace ggml_cuda_mma;
+    constexpr int warp_size = ggml_cuda_get_physical_warp_size();
+    constexpr int nthreads  = 4 * warp_size;
+    constexpr int BM = 64, BN = 64, BK = 64;
+    constexpr int AS = BK / 2 + 4;
+    constexpr int BS = BN / 2 + 4;
+    static_assert(AS * sizeof(half2) % sizeof(int4) == 0, "shared weight rows must be 16-byte aligned");
+    __shared__ __align__(16) half2 a_s[BM][AS];
+    __shared__ __align__(16) half2 b_s[BK][BS];
+    const int                      tid = threadIdx.y * warp_size + threadIdx.x;
+    const int                      iw = int(P.IW), ih = int(P.IH), id = int(P.ID);
+    const int                      ow = int(P.OW), oh = int(P.OH), od = int(P.OD);
+    const int                      kw = KW ? KW : int(P.KW), kh = KH ? KH : int(P.KH), kd = KD ? KD : int(P.KD);
+    const int                      ic = int(P.IC), oc = int(P.OC);
+    const int                      sx = int(P.ST_X), sy = int(P.ST_Y), sz = int(P.ST_Z);
+    const int                      dx = int(P.DL_X), dy = int(P.DL_Y), dz = int(P.DL_Z);
+    const int                      n = blockIdx.z / split_k, split = blockIdx.z % split_k;
+    const int                      m0 = blockIdx.y * BM, n0 = blockIdx.x * BN;
+    const int                      k_total   = ic * kw * kh * kd;
+    const int                      load_lane = warp_size == 32 ? threadIdx.x : threadIdx.x % (BN / 2);
+    const int load_row = threadIdx.y * (warp_size / (BN / 2)) + (warp_size == 32 ? 0 : threadIdx.x / (BN / 2));
+    const int spatial  = n0 + 2 * load_lane;
+    const int spatial0 = min(spatial, ow * oh * od - 1), spatial1 = min(spatial + 1, ow * oh * od - 1);
+    const int z0 = spatial0 / (ow * oh), y0 = spatial0 / ow % oh, x0 = spatial0 % ow;
+    const int z1 = spatial1 / (ow * oh), y1 = spatial1 / ow % oh, x1 = spatial1 % ow;
+    const int pos0                = (z0 * sz * ih + y0 * sy) * iw + x0 * sx;
+    const int pos1                = (z1 * sz * ih + y1 * sy) * iw + x1 * sx;
+    [[maybe_unused]] const int wm = threadIdx.y / 2 * 32, wn = threadIdx.y % 2 * 32;
+    using tile_ab = tile<16, 8, half2, get_input_data_layout()>;
+#if defined(AMD_WMMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE)
+    // AMD accumulator fragments transpose the input fragment's row/column mapping.
+    using tile_c = tile<16, 16, float, DATA_LAYOUT_J_MAJOR>;
+#else
+    using tile_c = tile<16, 16, float>;
+#endif
+    [[maybe_unused]] tile_c    c[2][2];
+    constexpr int              RM = 4, RN = BM * BN / (nthreads * RM);
+    [[maybe_unused]] const int simt_m = tid / (BN / RN) * RM, simt_n = tid % (BN / RN) * RN;
+    [[maybe_unused]] float     c_simt[RM][RN] = {};
+    const int                  tiles          = (k_total + BK - 1) / BK;
+    const int                  begin          = int(int64_t(tiles) * split / split_k) * BK;
+    const int                  end            = int(int64_t(tiles) * (split + 1) / split_k) * BK;
+    for (int k0 = begin; k0 < end; k0 += BK) {
+        if (aligned_weights) {
+#pragma unroll
+            for (int i0 = 0; i0 < BM * BK / 8; i0 += nthreads) {
+                const int  i   = i0 + tid;
+                const int  row = i / (BK / 8), col = 8 * (i % (BK / 8));
+                const int4 v                 = m0 + row < oc && k0 + col < k_total ?
+                                                   ((const int4 *) weight)[((m0 + row) * k_total + k0 + col) / 8] :
+                                                   make_int4(0, 0, 0, 0);
+                *(int4 *) &a_s[row][col / 2] = v;
+            }
+        } else {
+#pragma unroll
+            for (int i0 = 0; i0 < BM * BK / 2; i0 += nthreads) {
+                const int i   = i0 + tid;
+                const int row = i / (BK / 2), col = 2 * (i % (BK / 2));
+                half      lo = __float2half(0.0f), hi = lo;
+                if (m0 + row < oc && k0 + col < k_total) {
+                    lo = weight[(m0 + row) * k_total + k0 + col];
+                    if (k0 + col + 1 < k_total) {
+                        hi = weight[(m0 + row) * k_total + k0 + col + 1];
+                    }
+                }
+                a_s[row][col / 2] = __halves2half2(lo, hi);
+            }
+        }
+#pragma unroll
+        for (int kb = 0; kb < BK; kb += nthreads / (BN / 2)) {
+            const int k  = kb + load_row;
+            const int ki = k0 + k;
+            const int ci = ki / (kw * kh * kd), kz = ki / (kw * kh) % kd, ky = ki / kw % kh, kx = ki % kw;
+            const int offset = ki < k_total ? ((n * ic + ci) * id + kz * dz) * ih * iw + ky * dy * iw + kx * dx : 0;
+            half      lo = __float2half(0.0f), hi = lo;
+            if (ki < k_total && spatial < ow * oh * od) {
+                lo = input[offset + pos0];
+            }
+            if (ki < k_total && spatial + 1 < ow * oh * od) {
+                hi = input[offset + pos1];
+            }
+            b_s[k][load_lane] = __halves2half2(lo, hi);
+        }
+        __syncthreads();
+        if constexpr (use_mma) {
+#pragma unroll
+            for (int k = 0; k < BK; k += 16) {
+                tile_ab a[2], b[2];
+#pragma unroll
+                for (int i = 0; i < 2; ++i) {
+                    load_ldmatrix(a[i], &a_s[wm + 16 * i][k / 2], AS);
+                    load_ldmatrix_trans(b[i], &b_s[k][(wn + 16 * i) / 2], BS);
+                }
+#pragma unroll
+                for (int i = 0; i < 2; ++i) {
+#pragma unroll
+                    for (int j = 0; j < 2; ++j) {
+                        mma(c[i][j], a[i], b[j]);
+                    }
+                }
+            }
+        } else {
+#pragma unroll 4
+            for (int k = 0; k < BK; ++k) {
+                float a[RM], b[RN];
+#pragma unroll
+                for (int i = 0; i < RM; ++i) {
+                    a[i] = __half2float(((const half *) a_s[simt_m + i])[k]);
+                }
+#pragma unroll
+                for (int j = 0; j < RN; ++j) {
+                    b[j] = __half2float(((const half *) b_s[k])[simt_n + j]);
+                }
+#pragma unroll
+                for (int i = 0; i < RM; ++i) {
+#pragma unroll
+                    for (int j = 0; j < RN; ++j) {
+                        c_simt[i][j] += a[i] * b[j];
+                    }
+                }
+            }
+        }
+        __syncthreads();
+    }
+    if constexpr (use_mma) {
+#pragma unroll
+        for (int i = 0; i < 2; ++i) {
+#pragma unroll
+            for (int j = 0; j < 2; ++j) {
+#pragma unroll
+                for (int l = 0; l < c[i][j].ne; ++l) {
+                    const int co  = m0 + wm + 16 * i + c[i][j].get_i(l);
+                    const int pos = n0 + wn + 16 * j + c[i][j].get_j(l);
+                    if (co < oc && pos < ow * oh * od) {
+                        output[(int64_t(blockIdx.z) * oc + co) * ow * oh * od + pos] = c[i][j].x[l];
+                    }
+                }
+            }
+        }
+    } else {
+#pragma unroll
+        for (int i = 0; i < RM; ++i) {
+#pragma unroll
+            for (int j = 0; j < RN; ++j) {
+                const int co = m0 + simt_m + i, pos = n0 + simt_n + j;
+                if (co < oc && pos < ow * oh * od) {
+                    output[(int64_t(blockIdx.z) * oc + co) * ow * oh * od + pos] = c_simt[i][j];
+                }
+            }
+        }
+    }
+}
+
+static __global__ void conv3d_reduce_split_k(const float * __restrict__ partial,
+                                             float * __restrict__ output,
+                                             const int total,
+                                             const int per_batch,
+                                             const int split_k) {
+    const int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= total) {
+        return;
+    }
+    const int     n   = i / per_batch;
+    // Partial slices are ordered as [batch, split, output channel, spatial position].
+    const float * src = partial + int64_t(n) * (split_k - 1) * per_batch + i;
+    float         sum = 0.0f;
+    for (int k = 0; k < split_k; ++k) {
+        sum += src[int64_t(k) * per_batch];
+    }
+    output[i] = sum;
+}
+
+template <bool use_mma>
+static void conv3d_launch_implicit_gemm(const half *          input,
+                                        const half *          weight,
+                                        float *               output,
+                                        const conv3d_params & params,
+                                        int                   split_k,
+                                        dim3                  grid,
+                                        dim3                  block,
+                                        cudaStream_t          stream) {
+    // Vector loads require both the base pointer and each weight row to be 16-byte aligned.
+    const bool aligned_weights = uintptr_t(weight) % sizeof(int4) == 0 &&
+                                 (params.IC * params.KW * params.KH * params.KD) % (sizeof(int4) / sizeof(half)) == 0;
+    if (params.KW == 3 && params.KH == 3 && params.KD == 3) {
+        conv3d_implicit_gemm_f16<3, 3, 3, use_mma>
+            <<<grid, block, 0, stream>>>(input, weight, output, params, split_k, aligned_weights);
+    } else if (params.KW == 1 && params.KH == 1 && params.KD == 3) {
+        conv3d_implicit_gemm_f16<1, 1, 3, use_mma>
+            <<<grid, block, 0, stream>>>(input, weight, output, params, split_k, aligned_weights);
+    } else if (params.KW == 1 && params.KH == 1 && params.KD == 1) {
+        conv3d_implicit_gemm_f16<1, 1, 1, use_mma>
+            <<<grid, block, 0, stream>>>(input, weight, output, params, split_k, aligned_weights);
+    } else {
+        conv3d_implicit_gemm_f16<0, 0, 0, use_mma>
+            <<<grid, block, 0, stream>>>(input, weight, output, params, split_k, aligned_weights);
+    }
+}
+
+void ggml_cuda_op_conv3d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+    const ggml_tensor * kernel = dst->src[0];
+    const ggml_tensor * input  = dst->src[1];
+    GGML_ASSERT(input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32);
+    GGML_ASSERT(kernel->type == GGML_TYPE_F16 || kernel->type == GGML_TYPE_F32);
+    GGML_ASSERT(ggml_is_contiguous(input) && ggml_is_contiguous(kernel) && ggml_is_contiguous(dst));
+
+    const int32_t * p  = dst->op_params;
+    const int64_t   IW = input->ne[0], IH = input->ne[1], ID = input->ne[2];
+    const int64_t   OW = dst->ne[0], OH = dst->ne[1], OD = dst->ne[2];
+    const int64_t   KW = kernel->ne[0], KH = kernel->ne[1], KD = kernel->ne[2];
+    const int64_t   IC = p[9], B = p[10], OC = p[11];
+    GGML_ASSERT(IC > 0 && B > 0 && OC > 0 && input->ne[3] == IC * B && kernel->ne[3] == IC * OC);
+    GGML_ASSERT(dst->ne[3] == OC * B && p[0] > 0 && p[1] > 0 && p[2] > 0 && p[6] > 0 && p[7] > 0 && p[8] > 0);
+    const int64_t       total  = ggml_nelements(dst);
+    const conv3d_params params = { IW,   IH,   ID,   OW,   OH,   OD,   KW,   KH, KD, p[0], p[1],
+                                   p[2], p[3], p[4], p[5], p[6], p[7], p[8], IC, OC, B,    total };
+    const float *       x      = (const float *) input->data;
+    const half *        w      = (const half *) kernel->data;
+    float *             y      = (float *) dst->data;
+    cudaStream_t        stream = ctx.stream();
+    const auto &        device = ggml_cuda_info().devices[ctx.device];
+    const bool          use_mma =
+        turing_mma_available(device.cc) || amd_wmma_available(device.cc) || amd_mfma_available(device.cc);
+    const bool pointwise =
+        KW == 1 && KH == 1 && KD == 1 && p[0] == 1 && p[1] == 1 && p[2] == 1 && p[3] == 0 && p[4] == 0 && p[5] == 0;
+    const bool    use_blas = pointwise && fast_fp16_hardware_available(device.cc);
+    const int64_t limit    = INT_MAX - 256;
+    const int64_t pw = IW + 2 * int64_t(p[3]), ph = IH + 2 * int64_t(p[4]), pd = ID + 2 * int64_t(p[5]);
+    const bool    padded_fits = pw > 0 && pw <= limit && ph > 0 && ph <= limit && pd > 0 && pd <= limit &&
+                             pw * ph <= limit / pd && IC * B <= limit / (pw * ph * pd);
+    if (kernel->type == GGML_TYPE_F16 && KW > 0 && KH > 0 && KD > 0 && ggml_nelements(input) <= limit &&
+        ggml_nelements(kernel) <= limit && total <= limit && padded_fits && p[3] >= 0 && p[4] >= 0 && p[5] >= 0 &&
+        (OW - 1) * p[0] + (KW - 1) * p[6] < pw && (OH - 1) * p[1] + (KH - 1) * p[7] < ph &&
+        (OD - 1) * p[2] + (KD - 1) * p[8] < pd && (OC + 63) / 64 <= 65535 && B <= 65535) {
+        const int                  padded_total = int(pw * ph * pd * IC * B);
+        ggml_cuda_pool_alloc<half> x_half(ctx.pool(), padded_total);
+        // Match im2col's F16 input precision without materializing all patches in global memory.
+        if (p[3] == 0 && p[4] == 0 && p[5] == 0) {
+            ggml_get_to_fp16_cuda(input->type)(x, x_half.get(), padded_total, stream);
+        } else {
+            conv3d_pad_f16<<<(padded_total + 255) / 256, 256, 0, stream>>>(
+                x, x_half.get(), int(IW), int(IH), int(ID), int(pw), int(ph), int(pd), p[3], p[4], p[5], padded_total);
+        }
+        const conv3d_params padded_params = { pw,   ph, pd, OW, OH,   OD,   KW,   KH, KD, p[0], p[1],
+                                              p[2], 0,  0,  0,  p[6], p[7], p[8], IC, OC, B,    total };
+        const int           positions     = int(OW * OH * OD);
+        if (use_blas) {
+            const float    alpha = 1.0f, beta = 0.0f;
+            cublasHandle_t cublas_h = ctx.cublas_handle();
+            for (int n = 0; n < B; ++n) {
+                CUBLAS_CHECK(cublasGemmEx(cublas_h, CUBLAS_OP_N, CUBLAS_OP_N, positions, int(OC), int(IC), &alpha,
+                                          x_half.get() + int64_t(n) * IC * positions, CUDA_R_16F, positions, w,
+                                          CUDA_R_16F, int(IC), &beta, y + int64_t(n) * OC * positions, CUDA_R_32F,
+                                          positions, CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP));
+            }
+            return;
+        }
+        const int64_t blocks  = ((positions + 63) / 64) * ((OC + 63) / 64) * B;
+        const int     target  = 8 * device.nsm;
+        const int     split_k = int(std::min({ int64_t(32), int64_t(65535) / B, (IC * KW * KH * KD + 63) / 64,
+                                               std::max(int64_t(1), (target + blocks - 1) / blocks) }));
+        ggml_cuda_pool_alloc<float> partial(ctx.pool());
+        float *                     result = split_k == 1 ? y : partial.alloc(total * split_k);
+        const dim3                  block(device.warp_size, 4);
+        const dim3 grid(unsigned((positions + 63) / 64), unsigned((OC + 63) / 64), unsigned(B * split_k));
+        if (use_mma) {
+            conv3d_launch_implicit_gemm<true>(x_half.get(), w, result, padded_params, split_k, grid, block, stream);
+        } else {
+            conv3d_launch_implicit_gemm<false>(x_half.get(), w, result, padded_params, split_k, grid, block, stream);
+        }
+        if (split_k > 1) {
+            conv3d_reduce_split_k<<<unsigned((total + 255) / 256), 256, 0, stream>>>(result, y, int(total),
+                                                                                     int(OC * positions), split_k);
+        }
+        return;
+    }
+    const int blocks = int(std::min(int64_t(65535), (total + 255) / 256));
+    if (kernel->type == GGML_TYPE_F16) {
+        conv3d_kernel<<<blocks, 256, 0, stream>>>(x, w, y, params);
+    } else {
+        conv3d_kernel<<<blocks, 256, 0, stream>>>(x, (const float *) kernel->data, y, params);
+    }
+}

ggml/src/ggml-cuda/conv3d.cuh

@@ -0,0 +1,8 @@
+#ifndef GGML_CUDA_CONV3D_CUH
+#define GGML_CUDA_CONV3D_CUH
+
+#include "common.cuh"
+
+void ggml_cuda_op_conv3d(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
+
+#endif

ggml/src/ggml-cuda/fattn-mma-f16.cuh

@@ -2011,7 +2011,7 @@ static __global__ void flash_attn_ext_f16(
 #endif // defined(FLASH_ATTN_AVAILABLE) && (defined(VOLTA_MMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE))
 }
 
-bool ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(const int cc, const ggml_tensor * dst, const int ncols1);
+bool ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(const int cc, const ggml_tensor * dst, const int ncols1, const int ncols2);
 
 template <int DKQ, int DV, int ncols1, int ncols2>
 void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
@@ -2065,7 +2065,7 @@ void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml
         constexpr bool use_logit_softcap = false;
 #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
         if constexpr (ggml_cuda_flash_attn_ext_mma_f16_may_use_sparse(DKQ, DV, ncols1, ncols2)) {
-            if (ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(cc, dst, ncols1)) {
+            if (ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(cc, dst, ncols1, ncols2)) {
                 constexpr bool use_sparse_kernel = true;
                 fattn_kernel = flash_attn_ext_f16<DKQ, DV, ncols1, ncols2, use_logit_softcap, V_is_K_view, use_sparse_kernel>;
                 use_sparse = true;

ggml/src/ggml-cuda/fattn.cu

@@ -65,10 +65,11 @@
 
 #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
 // one list per group of ncols1 queries: a column is selected if any query of the group can see it
+template <int ncols1, bool oob>
 __launch_bounds__(256, 1)
 static __global__ void flash_attn_mask_to_sparse_indices(
         const half * mask_ptr, int32_t * indices_ptr, int32_t * counts_ptr, const int ne30, const int n_queries,
-        const int ncols1, const int n_kv_max, const int64_t s31, const int64_t s33) {
+        const int n_kv_max, const int64_t s31, const int64_t s33) {
     ggml_cuda_pdl_sync();
 
     constexpr int values_per_lane = 8;
@@ -100,8 +101,11 @@ static __global__ void flash_attn_mask_to_sparse_indices(
         for (int item = 0; item < values_per_lane; ++item) {
             const int i = i0 + (warp*values_per_lane + item)*WARP_SIZE + lane;
             bool selected = false;
-            for (int q = 0; q < q1 - q0 && !selected; ++q) {
+            if (i < ne30) {
-                selected = i < ne30 && isfinite(__half2float(mask[q*s31 + i]));
+#pragma unroll
+                for (int q = 0; q < ncols1; ++q) {
+                    selected |= (!oob || q < q1 - q0) && isfinite(__half2float(mask[q*s31 + i]));
+                }
             }
             selected_warp[item] = __ballot_sync(0xFFFFFFFF, selected);
             warp_count += __popc(selected_warp[item]);
@@ -168,15 +172,20 @@ void ggml_cuda_flash_attn_ext_compact_mask(
     const dim3 blocks_num((n_queries + ncols1 - 1)/ncols1, mask->ne[3], 1);
     const dim3 block_dim(256, 1, 1);
     const ggml_cuda_kernel_launch_params launch_params(blocks_num, block_dim, 0, stream);
-    ggml_cuda_kernel_launch(flash_attn_mask_to_sparse_indices, launch_params,
+    // the last group of queries is partial only if ncols1 does not divide n_queries
-        (const half *) mask->data, indices, counts, int(mask->ne[0]), n_queries, ncols1, n_kv_max, s31, s33);
+    GGML_ASSERT(ncols1 == 1 || ncols1 == 8);
+    const auto kernel = ncols1 == 1       ? flash_attn_mask_to_sparse_indices<1, false> :
+                        n_queries % 8 != 0 ? flash_attn_mask_to_sparse_indices<8, true>  :
+                                             flash_attn_mask_to_sparse_indices<8, false>;
+    ggml_cuda_kernel_launch(kernel, launch_params,
+        (const half *) mask->data, indices, counts, int(mask->ne[0]), n_queries, n_kv_max, s31, s33);
     CUDA_CHECK(cudaGetLastError());
 #endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
 }
 
-bool ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(const int cc, const ggml_tensor * dst, const int ncols1) {
+bool ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(const int cc, const ggml_tensor * dst, const int ncols1, const int ncols2) {
 #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA)
-    GGML_UNUSED_VARS(cc, dst, ncols1);
+    GGML_UNUSED_VARS(cc, dst, ncols1, ncols2);
     return false;
 #else
     const ggml_tensor * Q    = dst->src[0];
@@ -190,7 +199,8 @@ bool ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(const int cc, const ggml_
 
     const int32_t n_kv_max = ggml_get_op_params_i32(dst, 4);
 
-    const int64_t n_gather = (ncols1 == 1 ? Q->ne[1] : ncols1) * (int64_t) n_kv_max;
+    // the dense kernel handles up to 64/ncols2 queries per K/V pass, the single-query gather has to beat that
+    const int64_t n_gather = (ncols1 == 1 ? std::min<int64_t>(Q->ne[1], 64/ncols2) : ncols1) * (int64_t) n_kv_max;
 
     return GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) &&
         mask != nullptr && n_kv_max > 0 && max_bias == 0.0f && logit_softcap == 0.0f &&
@@ -206,7 +216,9 @@ static void ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1(ggml_backend_cuda_con
 
 #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
     if constexpr (ggml_cuda_flash_attn_ext_mma_f16_may_use_sparse(DKQ, DV, 1, ncols2)) {
-        if (ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(cc, dst, 1)) {
+        // a sparse variant at the full tile width gathers the union of its queries once, prefer it for large batches
+        constexpr bool has_wide_sparse = ggml_cuda_flash_attn_ext_mma_f16_may_use_sparse(DKQ, DV, 64/ncols2, ncols2);
+        if (!(has_wide_sparse && Q->ne[1] > 32/ncols2) && ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(cc, dst, 1, ncols2)) {
             ggml_cuda_flash_attn_ext_mma_f16_case<DKQ, DV, 1, ncols2>(ctx, dst);
             return;
         }
@@ -687,7 +699,7 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
                 // the sparse gather exists only in the MMA kernel: (DKQ, DV, 1, 8) with GQA > 4
                 const bool sparse_decode = gqa_opt_applies && gqa_ratio > 4 &&
                     ggml_cuda_flash_attn_ext_mma_f16_may_use_sparse(K->ne[0], V->ne[0], 1, 8) &&
-                    ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(cc, dst, 1);
+                    ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(cc, dst, 1, 8);
                 if (!sparse_decode && cc >= GGML_CUDA_CC_ADA_LOVELACE && Q->ne[1] == 1 && Q->ne[3] == 1 &&
                         !(gqa_ratio > 4 && (Q->ne[0] >= 256 || K->ne[1] >= 8192))) {
                     return BEST_FATTN_KERNEL_VEC;

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

@@ -19,6 +19,7 @@ bool g_mul_mat_q = true;
 #include "ggml-cuda/conv2d.cuh"
 #include "ggml-cuda/conv2d-dw.cuh"
 #include "ggml-cuda/conv2d-transpose.cuh"
+#include "ggml-cuda/conv3d.cuh"
 #include "ggml-cuda/convert.cuh"
 #include "ggml-cuda/count-equal.cuh"
 #include "ggml-cuda/cpy.cuh"
@@ -2327,6 +2328,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
         case GGML_OP_CONV_2D:
             ggml_cuda_op_conv2d(ctx, dst);
             break;
+        case GGML_OP_CONV_3D:
+            ggml_cuda_op_conv3d(ctx, dst);
+            break;
         case GGML_OP_CONV_2D_DW:
             ggml_cuda_op_conv2d_dw(ctx, dst);
             break;
@@ -3479,6 +3483,7 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
         ggml_cuda_topk_moe_args args;
         const bool              can_fuse = ggml_cuda_topk_moe_fusion(cgraph, i, args);
         std::vector<ggml_op>    ops;
+        ops.reserve(13);  // max ops; avoids gcc -Wstringop-overflow false positive
 
         if (can_fuse) {
             const ggml_tensor * logits  = node->src[0];
@@ -4519,24 +4524,91 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph
     ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) backend->context;
 
     static const bool disable_fusion = getenv("GGML_CUDA_DISABLE_FUSION") != nullptr && std::atoi(getenv("GGML_CUDA_DISABLE_FUSION"));
+
+    auto add_alloc_deps = [&](size_t start, size_t last_node) {
+
+        for (size_t i = start; i < last_node; ++i) {
+            params->add_alloc_dep(params->user_data, cgraph->nodes[i], cgraph->nodes[last_node]);
+
+            for (int j = 0; j < GGML_MAX_SRC; ++j) {
+                if (cgraph->nodes[i]->src[j]) {
+                    params->add_alloc_dep(params->user_data, cgraph->nodes[i]->src[j], cgraph->nodes[last_node]);
+                }
+            }
+        }
+    };
+
     if (!disable_fusion) {
+        // add alloc deps for performance positive fusions. This may increase the overall compute buffer size.
+        // TODO: consolidate fusion paths in graph_optimize and graph_compute
         for (int i = 0; i < cgraph->n_nodes; ++i) {
-            if (cgraph->nodes[i]->op != GGML_OP_MUL) {
-                continue;
-            }
-
             ggml_cuda_moe_weighted_reduction_match match;
-            if (!ggml_cuda_match_moe_weighted_reduction(cgraph, i, match)) {
+            if (ggml_cuda_match_moe_weighted_reduction(cgraph, i, match)) {
-                continue;
+                params->add_alloc_dep(params->user_data, const_cast<ggml_tensor *>(match.experts), match.dst);
+                params->add_alloc_dep(params->user_data, const_cast<ggml_tensor *>(match.weights), match.dst);
+                if (match.expert_scale != nullptr) {
+                    params->add_alloc_dep(
+                        params->user_data, const_cast<ggml_tensor *>(match.expert_scale), match.dst);
+                }
+                i += match.node_count - 1;
             }
 
-            params->add_alloc_dep(params->user_data, const_cast<ggml_tensor *>(match.experts), match.dst);
+            if (cgraph->nodes[i]->op == GGML_OP_UNARY || cgraph->nodes[i]->op == GGML_OP_SOFT_MAX ||
-            params->add_alloc_dep(params->user_data, const_cast<ggml_tensor *>(match.weights), match.dst);
+                    cgraph->nodes[i]->op == GGML_OP_ARGSORT) {
-            if (match.expert_scale != nullptr) {
+                ggml_cuda_topk_moe_args args;
-                params->add_alloc_dep(
+                const bool              can_fuse = ggml_cuda_topk_moe_fusion(cgraph, i, args);
-                    params->user_data, const_cast<ggml_tensor *>(match.expert_scale), match.dst);
+                std::vector<ggml_op>    ops;
+                ops.reserve(13);  // max ops; avoids gcc -Wstringop-overflow false positive
+
+                const ggml_tensor * node = cgraph->nodes[i];
+
+                if (can_fuse) {
+                    const ggml_tensor * logits  = node->src[0];
+                    ggml_tensor *       weights = nullptr;
+                    ggml_tensor *       ids     = nullptr;
+
+                    if (!args.delayed_softmax) {
+                        int out_nodes[2];  // nodes which can't be elided
+
+                        if (args.sigmoid) {
+                            ops.insert(ops.end(), { GGML_OP_UNARY });
+                        } else if (args.sqrt_softplus) {
+                            ops.insert(ops.end(), { GGML_OP_UNARY, GGML_OP_SQRT });
+                        } else {
+                            ops.insert(ops.end(), { GGML_OP_SOFT_MAX });
+                        }
+                        const int i_probs = i + (int) ops.size() - 1;  // last node of the gating activation
+
+                        if (args.prob_bias) {
+                            ops.insert(ops.end(), { GGML_OP_RESHAPE, GGML_OP_ADD, GGML_OP_ARGSORT, GGML_OP_VIEW,
+                                                    GGML_OP_GET_ROWS });
+                            out_nodes[0] = i_probs + 4;
+                        } else {
+                            ops.insert(ops.end(), { GGML_OP_RESHAPE, GGML_OP_ARGSORT, GGML_OP_VIEW, GGML_OP_GET_ROWS });
+                            out_nodes[0] = i_probs + 3;
+                        }
+                        ids = cgraph->nodes[out_nodes[0]];
+
+                        if (args.norm) {
+                            ops.insert(ops.end(),
+                                       { GGML_OP_RESHAPE, GGML_OP_SUM_ROWS, GGML_OP_CLAMP, GGML_OP_DIV, GGML_OP_RESHAPE });
+                        }
+                        if (args.scale) {
+                            ops.insert(ops.end(), { GGML_OP_SCALE });
+                        }
+
+                        weights      = cgraph->nodes[i + ops.size() - 1];
+                        out_nodes[1] = i + ops.size() - 1;
+
+                        if (ggml_can_fuse_subgraph(cgraph, i, ops.size(), ops.data(), out_nodes, 2) &&
+                                ggml_cuda_should_use_topk_moe(node, logits, weights, ids)) {
+
+                            add_alloc_deps(i, i + ops.size());
+                            i += ops.size() - 1;
+                        }
+                    }
+                }
             }
-            i += match.node_count - 1;
         }
     }
 
@@ -5458,6 +5530,10 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
         case GGML_OP_IM2COL_3D:
         case GGML_OP_CONV_2D:
             return (ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]));
+        case GGML_OP_CONV_3D:
+            return (op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_F32) &&
+                   op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&
+                   ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]) && ggml_is_contiguous(op);
         case GGML_OP_CONV_2D_DW:
             return op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_CONV_TRANSPOSE_2D:

ggml/src/ggml-cuda/vendors/hip.h

@@ -277,7 +277,15 @@ static __device__ __forceinline__ int __vsubss4(const int a, const int b) {
 }
 
 static __device__ __forceinline__ int __vsub4(const int a, const int b) {
-    return __vsubss4(a, b);
+    // do some small modifications to a and b to make the subtraction not underflow
+    const unsigned int a_large = a | 0x80808080;
+    const unsigned int b_small = b & 0x7f7f7f7f;
+    const unsigned int result_low_7bits = a_large - b_small;
+
+    // if two ops share the same high bit, we should flip the high bit of the result
+    const unsigned int should_flip_high_1bit = (a ^ ~b) & 0x80808080;
+
+    return result_low_7bits ^ should_flip_high_1bit;
 }
 
 static __device__ __forceinline__ unsigned int __vcmpeq4(unsigned int a, unsigned int b) {
@@ -293,13 +301,13 @@ static __device__ __forceinline__ unsigned int __vcmpeq4(unsigned int a, unsigne
 }
 
 static __device__ __forceinline__ unsigned int __vcmpne4(unsigned int a, unsigned int b) {
-    const uint8x4_t& va = reinterpret_cast<const uint8x4_t&>(a);
+    const unsigned int x = a ^ b;
-    const uint8x4_t& vb = reinterpret_cast<const uint8x4_t&>(b);
+
-    unsigned int c;
+    // any non-equal bit in a byte will set the high bit of that byte here
-    uint8x4_t& vc = reinterpret_cast<uint8x4_t&>(c);
+    // the addition will not overflow in the byte as op1 and op2 are both less than 0x80
-#pragma unroll
+    const unsigned int ne_low_7bits = ((x & 0x7f7f7f7f) + 0x7f7f7f7f) & 0x80808080;
-    for (int i = 0; i < 4; ++i) {
+    const unsigned int ne_high_1bit = x & 0x80808080;
-        vc[i] = va[i] == vb[i] ? 0x00 : 0xff;
+    const unsigned int ne_any_bit = ne_low_7bits | ne_high_1bit;
-    }
+
-    return c;
+    return (ne_any_bit >> 7) * 0xff;
 }

ggml/src/ggml-metal/ggml-metal-ops.cpp

@@ -3584,7 +3584,6 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
         auto cfg = use_sparse
                 ? ggml_metal_tuning::fa_vec_baseline_cfg((int) ne00, (int) ne20)
                 : ggml_metal_tuning::fa_vec_pick(
-                          props_dev->device_id,
                           props_dev->gpu_family,
                           (int) op->src[1]->type,
                           (int) ne00, (int) ne20,   // dk, dv (ne00 == dk for FA)

ggml/src/ggml-metal/ggml-metal-tuning.h

@@ -1,6 +1,5 @@
 #pragma once
 
-#include "ggml-metal-device.h"  // enum ggml_metal_device_id
 #include "ggml.h"
 
 #include <cstdint>
@@ -32,7 +31,7 @@ constexpr int8_t FA_VEC_DOMAIN_DECODE = 0;  // ne01 == 1
 constexpr int8_t FA_VEC_DOMAIN_BATCH  = 1;  // ne01 >= 2
 
 struct fa_vec_key_t {
-    int8_t  device_id;
+    int8_t  family;
     int8_t  dtype;
     int16_t dk;
     int16_t dv;
@@ -70,8 +69,7 @@ void         fa_vec_set_override(fa_vec_cfg_t cfg);
 void         fa_vec_clear_override();
 fa_vec_cfg_t fa_vec_baseline_cfg(int dk, int dv);
 
-// device_id selects a per-SKU row; on a miss, gpu_family (0 if unknown) maps to a representative
+// Keyed by Apple GPU family; an untuned family matches no row and gets the baseline.
-// SKU and the table is retried. No match -> baseline.
+fa_vec_cfg_t fa_vec_pick(int gpu_family, int dtype, int dk, int dv, int64_t ne11, int64_t ne01);
-fa_vec_cfg_t fa_vec_pick(enum ggml_metal_device_id device_id, int gpu_family, int dtype, int dk, int dv, int64_t ne11, int64_t ne01);
 
 }  // namespace ggml_metal_tuning

ggml/src/ggml-metal/kernels/mul_mv.metal

@@ -1043,6 +1043,7 @@ template [[host_name("kernel_mul_mv_f16_f16")]]   kernel mul_mv_t_t kernel_mul_m
 #if defined(GGML_METAL_HAS_BF16)
 template [[host_name("kernel_mul_mv_bf16_f32")]]  kernel mul_mv_t_t kernel_mul_mv_t_t<bfloat, float>;
 template [[host_name("kernel_mul_mv_bf16_bf16")]] kernel mul_mv_t_t kernel_mul_mv_t_t<bfloat, bfloat>;
+template [[host_name("kernel_mul_mv_f32_bf16")]]  kernel mul_mv_t_t kernel_mul_mv_t_t<float,  bfloat>;
 #endif
 
 template<typename T0, typename T04, typename T1, typename T14, short NR0, typename args_t>
@@ -1167,6 +1168,7 @@ template [[host_name("kernel_mul_mv_f16_f16_4")]]   kernel mul_mv_t_t_4 kernel_m
 #if defined(GGML_METAL_HAS_BF16)
 template [[host_name("kernel_mul_mv_bf16_f32_4")]]  kernel mul_mv_t_t_4 kernel_mul_mv_t_t_4<bfloat, bfloat4, float,  float4>;
 template [[host_name("kernel_mul_mv_bf16_bf16_4")]] kernel mul_mv_t_t_4 kernel_mul_mv_t_t_4<bfloat, bfloat4, bfloat, bfloat4>;
+template [[host_name("kernel_mul_mv_f32_bf16_4")]]  kernel mul_mv_t_t_4 kernel_mul_mv_t_t_4<float,  float4,  bfloat, bfloat4>;
 #endif
 
 template<typename T0, typename T1, typename args_t>
@@ -1232,6 +1234,7 @@ template [[host_name("kernel_mul_mv_f16_f16_short")]]  kernel mul_mv_t_t_short_t
 #if defined(GGML_METAL_HAS_BF16)
 template [[host_name("kernel_mul_mv_bf16_f32_short")]]  kernel mul_mv_t_t_short_t kernel_mul_mv_t_t_short<bfloat, float>;
 template [[host_name("kernel_mul_mv_bf16_bf16_short")]] kernel mul_mv_t_t_short_t kernel_mul_mv_t_t_short<bfloat, bfloat>;
+template [[host_name("kernel_mul_mv_f32_bf16_short")]]  kernel mul_mv_t_t_short_t kernel_mul_mv_t_t_short<float,  bfloat>;
 #endif
 
 template<int nr0, typename args_t>

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -1492,6 +1492,7 @@ static bool ggml_vk_matmul_int_shmem_support(const vk_device& device, const std:
         case GGML_TYPE_Q5_0:    block_a_size = std430_size({{16, 4}, {4, 4}, {fp_size,  fp_align}});          break; // qs[16/4] + qh + dm
         case GGML_TYPE_Q5_1:    block_a_size = std430_size({{16, 4}, {4, 4}, {fp2_size, fp2_align}});         break; // qs[16/4] + qh + dm(vec2)
         case GGML_TYPE_Q8_0:    block_a_size = std430_size({{32, 4}, {fp_size,  fp_align}});                  break; // qs[8] + dm
+        case GGML_TYPE_IQ4_XS:  block_a_size = std430_size({{32, 4}, {fp_size,  fp_align}});                  break; // qs[8] + d
         case GGML_TYPE_MXFP4:   block_a_size = std430_size({{32, 4}, {fp_size,  fp_align}});                  break; // qs[8] + d
         case GGML_TYPE_Q2_K:    block_a_size = std430_size({{ 8, 4}, {2, 2}, {fp2_size, fp2_align}});         break; // qs[2] + scales(u8vec2) + dm(vec2)
         case GGML_TYPE_Q3_K:    block_a_size = std430_size({{16, 4}, {fp2_size, fp2_align}});                 break; // qs[4] + d_scales(vec2)
@@ -2460,6 +2461,7 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
                 sg_create_mmq({GGML_TYPE_Q5_1, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_q5_1_q8_1", matmul_q5_1_q8_1_len, matmul_q5_1_q8_1_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q8_0, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_q8_0_q8_1", matmul_q8_0_q8_1_len, matmul_q8_0_q8_1_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_MXFP4, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_mxfp4_q8_1", matmul_mxfp4_q8_1_len, matmul_mxfp4_q8_1_data, sizeof(vk_mat_mat_push_constants), 3);
+                sg_create_mmq({GGML_TYPE_IQ4_XS, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_iq4_xs_q8_1", matmul_iq4_xs_q8_1_len, matmul_iq4_xs_q8_1_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q2_K, GGML_TYPE_Q8_1, false, false}, tc_mmq_int_k, "matmul_q2_k_q8_1", matmul_q2_k_q8_1_len, matmul_q2_k_q8_1_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q3_K, GGML_TYPE_Q8_1, false, false}, tc_mmq_int_k, "matmul_q3_k_q8_1", matmul_q3_k_q8_1_len, matmul_q3_k_q8_1_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q4_K, GGML_TYPE_Q8_1, false, false}, tc_mmq_int_k, "matmul_q4_k_q8_1", matmul_q4_k_q8_1_len, matmul_q4_k_q8_1_data, sizeof(vk_mat_mat_push_constants), 3);
@@ -2497,6 +2499,7 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
                     sg_create_mmq({GGML_TYPE_Q5_1, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_subgroup_q5_1_q8_1", matmul_id_subgroup_q5_1_q8_1_len, matmul_id_subgroup_q5_1_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count, mul_mat_subgroup_size);
                     sg_create_mmq({GGML_TYPE_Q8_0, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_subgroup_q8_0_q8_1", matmul_id_subgroup_q8_0_q8_1_len, matmul_id_subgroup_q8_0_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count, mul_mat_subgroup_size);
                     sg_create_mmq({GGML_TYPE_MXFP4, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_subgroup_mxfp4_q8_1", matmul_id_subgroup_mxfp4_q8_1_len, matmul_id_subgroup_mxfp4_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count, mul_mat_subgroup_size);
+                    sg_create_mmq({GGML_TYPE_IQ4_XS, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_subgroup_iq4_xs_q8_1", matmul_id_subgroup_iq4_xs_q8_1_len, matmul_id_subgroup_iq4_xs_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count, mul_mat_subgroup_size);
                     sg_create_mmq({GGML_TYPE_Q2_K, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int_k, "matmul_id_subgroup_q2_k_q8_1", matmul_id_subgroup_q2_k_q8_1_len, matmul_id_subgroup_q2_k_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count, mul_mat_subgroup_size_16);
                     sg_create_mmq({GGML_TYPE_Q3_K, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int_k, "matmul_id_subgroup_q3_k_q8_1", matmul_id_subgroup_q3_k_q8_1_len, matmul_id_subgroup_q3_k_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count, mul_mat_subgroup_size_16);
                     sg_create_mmq({GGML_TYPE_Q4_K, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int_k, "matmul_id_subgroup_q4_k_q8_1", matmul_id_subgroup_q4_k_q8_1_len, matmul_id_subgroup_q4_k_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count, mul_mat_subgroup_size_16);
@@ -2533,6 +2536,7 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
                     sg_create_mmq({GGML_TYPE_Q5_1, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_q5_1_q8_1", matmul_id_q5_1_q8_1_len, matmul_id_q5_1_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count);
                     sg_create_mmq({GGML_TYPE_Q8_0, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_q8_0_q8_1", matmul_id_q8_0_q8_1_len, matmul_id_q8_0_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count);
                     sg_create_mmq({GGML_TYPE_MXFP4, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_mxfp4_q8_1", matmul_id_mxfp4_q8_1_len, matmul_id_mxfp4_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count);
+                    sg_create_mmq({GGML_TYPE_IQ4_XS, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int, "matmul_id_iq4_xs_q8_1", matmul_id_iq4_xs_q8_1_len, matmul_id_iq4_xs_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count);
                     sg_create_mmq({GGML_TYPE_Q2_K, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int_k, "matmul_id_q2_k_q8_1", matmul_id_q2_k_q8_1_len, matmul_id_q2_k_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count);
                     sg_create_mmq({GGML_TYPE_Q3_K, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int_k, "matmul_id_q3_k_q8_1", matmul_id_q3_k_q8_1_len, matmul_id_q3_k_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count);
                     sg_create_mmq({GGML_TYPE_Q4_K, GGML_TYPE_Q8_1, true, false}, tc_mmqid_int_k, "matmul_id_q4_k_q8_1", matmul_id_q4_k_q8_1_len, matmul_id_q4_k_q8_1_data, sizeof(vk_mat_mat_id_push_constants), mul_mat_id_param_count);
@@ -2570,6 +2574,7 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
                 sg_create_mmq({GGML_TYPE_Q5_0, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_q5_0_q8_1", matmul_q5_0_q8_1_fp32_len, matmul_q5_0_q8_1_fp32_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q5_1, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_q5_1_q8_1", matmul_q5_1_q8_1_fp32_len, matmul_q5_1_q8_1_fp32_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q8_0, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_q8_0_q8_1", matmul_q8_0_q8_1_fp32_len, matmul_q8_0_q8_1_fp32_data, sizeof(vk_mat_mat_push_constants), 3);
+                sg_create_mmq({GGML_TYPE_IQ4_XS, GGML_TYPE_Q8_1, false, false}, tc_mmq_int, "matmul_iq4_xs_q8_1", matmul_iq4_xs_q8_1_fp32_len, matmul_iq4_xs_q8_1_fp32_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q2_K, GGML_TYPE_Q8_1, false, false}, tc_mmq_int_k, "matmul_q2_k_q8_1", matmul_q2_k_q8_1_fp32_len, matmul_q2_k_q8_1_fp32_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q3_K, GGML_TYPE_Q8_1, false, false}, tc_mmq_int_k, "matmul_q3_k_q8_1", matmul_q3_k_q8_1_fp32_len, matmul_q3_k_q8_1_fp32_data, sizeof(vk_mat_mat_push_constants), 3);
                 sg_create_mmq({GGML_TYPE_Q4_K, GGML_TYPE_Q8_1, false, false}, tc_mmq_int_k, "matmul_q4_k_q8_1", matmul_q4_k_q8_1_fp32_len, matmul_q4_k_q8_1_fp32_data, sizeof(vk_mat_mat_push_constants), 3);
@@ -2822,6 +2827,7 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
 
                 ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_q8_1_f32[w][GGML_TYPE_IQ1_S][i], "mul_mat_vec_iq1_s_q8_1_f32", arr_dmmv_iq1_s_q8_1_f32_len[reduc], arr_dmmv_iq1_s_q8_1_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1*rm_iq_int(i), 1, 1}, {wg_size_subgroup_int, 1*rm_iq_int(i), i+1}, 1, true, use_subgroups, subgroup_size_int);
                 ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_q8_1_f32[w][GGML_TYPE_IQ1_M][i], "mul_mat_vec_iq1_m_q8_1_f32", arr_dmmv_iq1_m_q8_1_f32_len[reduc], arr_dmmv_iq1_m_q8_1_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1*rm_iq_int(i), 1, 1}, {wg_size_subgroup_int, 1*rm_iq_int(i), i+1}, 1, true, use_subgroups, subgroup_size_int);
+                ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_q8_1_f32[w][GGML_TYPE_IQ4_XS][i], "mul_mat_vec_iq4_xs_q8_1_f32", arr_dmmv_iq4_xs_q8_1_f32_len[reduc], arr_dmmv_iq4_xs_q8_1_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_int_n(1*rm_stdq_int, i), 1, 1}, {wg_size_subgroup_int, rm_int_n(1*rm_stdq_int, i), i+1}, 1, true, use_subgroups, subgroup_size_int);
 
             }
 #endif // GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT
@@ -2878,6 +2884,7 @@ void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
 
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_q8_1_f32[w][GGML_TYPE_IQ1_S], "mul_mat_vec_id_iq1_s_q8_1_f32", arr_dmmv_id_iq1_s_q8_1_f32_len[reduc], arr_dmmv_id_iq1_s_q8_1_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {1*rm_iq_int(0), 1, 1}, {wg_size_subgroup_int, 1*rm_iq_int(0)}, 1, true, use_subgroups, subgroup_size_int);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_q8_1_f32[w][GGML_TYPE_IQ1_M], "mul_mat_vec_id_iq1_m_q8_1_f32", arr_dmmv_id_iq1_m_q8_1_f32_len[reduc], arr_dmmv_id_iq1_m_q8_1_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {1*rm_iq_int(0), 1, 1}, {wg_size_subgroup_int, 1*rm_iq_int(0)}, 1, true, use_subgroups, subgroup_size_int);
+            ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_q8_1_f32[w][GGML_TYPE_IQ4_XS], "mul_mat_vec_id_iq4_xs_q8_1_f32", arr_dmmv_id_iq4_xs_q8_1_f32_len[reduc], arr_dmmv_id_iq4_xs_q8_1_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_id(1*rm_stdq_int), 1, 1}, {wg_size_subgroup_int, rm_id(1*rm_stdq_int)}, 1, true, use_subgroups, subgroup_size_int);
         }
 #endif // GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT
     }
@@ -5341,6 +5348,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec(ggml_backend_vk_context *
             case GGML_TYPE_Q6_K:
             case GGML_TYPE_IQ1_S:
             case GGML_TYPE_IQ1_M:
+            case GGML_TYPE_IQ4_XS:
                 break;
             default:
                 return nullptr;
@@ -5428,6 +5436,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec_id(ggml_backend_vk_context
             case GGML_TYPE_Q6_K:
             case GGML_TYPE_IQ1_S:
             case GGML_TYPE_IQ1_M:
+            case GGML_TYPE_IQ4_XS:
                 break;
             default:
                 return nullptr;
@@ -6411,6 +6420,7 @@ static bool ggml_vk_should_use_mmvq(const vk_device& device, uint32_t m, uint32_
         // From tests on A770 Linux, may need more tuning
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q5_1:
+        case GGML_TYPE_IQ4_XS:
             return false;
         default:
             return true;

ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vecq.comp

@@ -5,6 +5,7 @@
 
 #define MMQ
 #define NEEDS_IQ1S_GRID_GPU
+#define KVALUES_IQ4NL_I8
 #define B_TYPE block_q8_1_x4
 
 #include "mul_mat_vec_base.glsl"
@@ -15,7 +16,7 @@ layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
 #define K_PER_ITER 16
 #elif defined(DATA_A_QUANT_LEGACY) || defined(DATA_A_MXFP4)
 #define K_PER_ITER 8
-#elif defined(DATA_A_IQ1_S) || defined(DATA_A_IQ1_M)
+#elif defined(DATA_A_IQ1_S) || defined(DATA_A_IQ1_M) || defined(DATA_A_IQ4_XS)
 #define K_PER_ITER 32
 #else
 #error unimplemented

ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vecq_funcs.glsl

@@ -448,6 +448,28 @@ FLOAT_TYPE mmvq_dot_product(const uint ib_a, const uint iqs) {
 }
 #endif
 
+#if defined(DATA_A_IQ4_XS)
+FLOAT_TYPE mmvq_dot_product(const uint ib_a, const uint iqs) {
+    const uint ib = ib_a / 8;
+    const uint ib32 = ib_a % 8;
+
+    int32_t q_sum = 0;
+    [[unroll]] for (uint j = 0; j < 4; ++j) {
+        const uint32_t vui = data_a_packed32[ib].qs[4 * ib32 + j];
+        const i32vec2 qs_a = iq4nl_to_i8x8(vui);
+
+        q_sum += dotPacked4x8EXT(qs_a.x, cache_b_qs[j]);
+        q_sum += dotPacked4x8EXT(qs_a.y, cache_b_qs[j + 4]);
+    }
+
+    const uint sl = (data_a_packed32[ib].scales_l >> (4 * ib32)) & 0xF;
+    const uint sh = (data_a_packed32[ib].scales_h >> (2 * ib32)) & 3;
+    const float d = float(data_a[ib].d) * float(int(sl | (sh << 4)) - 32);
+
+    return FLOAT_TYPE(float(cache_b_ds.x) * d * float(q_sum));
+}
+#endif
+
 #if defined(DATA_A_IQ1_S)
 void repack8(uint ib, uint iqs, out i32vec4 out0, out i32vec4 out1) {
     const uint ib32 = iqs / 32;

ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl

@@ -261,22 +261,28 @@ void load_a_to_shmem(const uint pos_a, const uint row, const uint col, const uin
 
 #elif defined(DATA_A_IQ4_XS)
     const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row;
-    const uint k_pair = row * LOAD_VEC_A / 2;
+    const uint k_pair = row * LOAD_VEC_A / 4;
 
-    const uint ib = idx / 64;
+    const uint ib = idx / 32;
-    const uint ib32 = (idx % 64) / 8;
+    const uint ib32 = (idx % 32) / 4;
     const uint iq = 4 * ib32 + (idx % 4);
 
     const uint sl = (data_a[ib].scales_l[ib32/2] >> (4 * (ib32 & 1))) & 0xF;
     const uint sh = ((data_a[ib].scales_h) >> (2 * ib32)) & 3;
-    const uint qshift = idx & 4;
-    u8vec4 qs = unpack8((uint(data_a_packed32[ib].qs[iq]) >> qshift) & 0x0F0F0F0F);
-
     const float d = float(data_a[ib].d);
-    const vec4 v = d * float(int(sl | (sh << 4)) - 32) * vec4(kvalues_iq4nl[qs.x], kvalues_iq4nl[qs.y], kvalues_iq4nl[qs.z], kvalues_iq4nl[qs.w]);
+    const float dl = d * float(int(sl | (sh << 4)) - 32);
+    const uint vui = uint(data_a_packed32[ib].qs[iq]);
+
+    const u8vec4 qs0 = unpack8( vui       & 0x0F0F0F0F);
+    const u8vec4 qs1 = unpack8((vui >> 4) & 0x0F0F0F0F);
+    const vec4 v0 = dl * vec4(kvalues_iq4nl[qs0.x], kvalues_iq4nl[qs0.y], kvalues_iq4nl[qs0.z], kvalues_iq4nl[qs0.w]);
+    const vec4 v1 = dl * vec4(kvalues_iq4nl[qs1.x], kvalues_iq4nl[qs1.y], kvalues_iq4nl[qs1.z], kvalues_iq4nl[qs1.w]);
+
+    store_a(col, k_pair,     FLOAT_TYPEV2(v0.xy));
+    store_a(col, k_pair + 1, FLOAT_TYPEV2(v0.zw));
+    store_a(col, k_pair + 8, FLOAT_TYPEV2(v1.xy));
+    store_a(col, k_pair + 9, FLOAT_TYPEV2(v1.zw));
 
-    store_a(col, k_pair, FLOAT_TYPEV2(v.xy));
-    store_a(col, k_pair + 1, FLOAT_TYPEV2(v.zw));
 #elif defined(DATA_A_IQ4_NL)
     const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row;
     const uint k_pair = row * LOAD_VEC_A / 4;

ggml/src/ggml-vulkan/vulkan-shaders/mul_mmq.comp

@@ -6,6 +6,8 @@
 
 #extension GL_EXT_integer_dot_product : require
 
+#define KVALUES_IQ4NL_I8
+
 #ifdef FLOAT16
 #extension GL_EXT_shader_explicit_arithmetic_types_float16 : require
 #endif

ggml/src/ggml-vulkan/vulkan-shaders/mul_mmq_funcs.glsl

@@ -217,6 +217,41 @@ ACC_TYPE mmq_dot_product(const uint ib_a) {
 }
 #endif
 
+#if defined(DATA_A_IQ4_XS)
+void block_a_to_shmem(const uint buf_ib, const uint ib, const uint iqs) {
+    const uint ib_k = ib / 8;
+    const uint ib32 = ib % 8;
+    const uint32_t vui = data_a_packed32[ib_k].qs[4 * ib32 + iqs];
+    const i32vec2 qs = iq4nl_to_i8x8(vui);
+
+    buf_a[buf_ib].qs[iqs    ] = qs.x;
+    buf_a[buf_ib].qs[iqs + 4] = qs.y;
+
+    if (iqs == 0) {
+        const uint sl = (data_a_packed32[ib_k].scales_l >> (4 * ib32)) & 0xF;
+        const uint sh = (data_a_packed32[ib_k].scales_h >> (2 * ib32)) & 3;
+        buf_a[buf_ib].d = FLOAT_TYPE(float(data_a[ib_k].d) * float(int(sl | (sh << 4)) - 32));
+    }
+}
+
+void block_a_to_registers(const uint reg_ib, const uint buf_ib) {
+    cache_a[reg_ib].d = buf_a[buf_ib].d;
+
+    [[unroll]] for (uint iqs = 0; iqs < 8; iqs++) {
+        cache_a[reg_ib].qs[iqs] = buf_a[buf_ib].qs[iqs];
+    }
+}
+
+ACC_TYPE mmq_dot_product(const uint ib_a) {
+    int32_t q_sum = 0;
+    [[unroll]] for (uint iqs = 0; iqs < 8; iqs++) {
+        q_sum += dotPacked4x8EXT(cache_a[ib_a].qs[iqs], cache_b.qs[iqs]);
+    }
+
+    return ACC_TYPE(float(cache_a[ib_a].d) * float(cache_b.ds.x) * float(q_sum));
+}
+#endif
+
 // For k-quants, ib and iqs still assume 32-wide blocks, but k-quants are 256-wide
 // iqs still refers to a 32-bit integer, meaning 0..7 for 32-wide quants
 #if defined(DATA_A_Q2_K)

ggml/src/ggml-vulkan/vulkan-shaders/mul_mmq_shmem_types.glsl

@@ -53,6 +53,12 @@ struct block_a_cache {
     int32_t qs[8];
     FLOAT_TYPE dm;
 };
+#elif defined(DATA_A_IQ4_XS)
+#define QUANT_R_MMQ 2
+struct block_a_cache {
+    int32_t qs[8];
+    FLOAT_TYPE d;
+};
 #elif defined(DATA_A_MXFP4)
 #define QUANT_R_MMQ 2
 struct block_a_cache {

ggml/src/ggml-vulkan/vulkan-shaders/types.glsl

@@ -1875,7 +1875,11 @@ const int8_t kvalues_iq4nl_const[16] = {
     int8_t(1), int8_t(13), int8_t(25), int8_t(38), int8_t(53), int8_t(69), int8_t(89), int8_t(113)
 };
 
+#ifdef KVALUES_IQ4NL_I8
+shared int8_t kvalues_iq4nl[16];
+#else
 shared FLOAT_TYPE kvalues_iq4nl[16];
+#endif
 
 #if defined(DATA_A_IQ4_NL) || defined(DATA_A_IQ4_XS)
 #define NEEDS_INIT_IQ_SHMEM
@@ -1883,10 +1887,25 @@ void init_iq_shmem(uvec3 wgsize)
 {
     // copy the table into shared memory and sync
     for (uint i = gl_LocalInvocationIndex.x; i < kvalues_iq4nl.length(); i += wgsize.x) {
+#ifdef KVALUES_IQ4NL_I8
+        kvalues_iq4nl[i] = kvalues_iq4nl_const[i];
+#else
         kvalues_iq4nl[i] = FLOAT_TYPE(kvalues_iq4nl_const[i]);
+#endif
     }
     barrier();
 }
+
+#ifdef KVALUES_IQ4NL_I8
+i32vec2 iq4nl_to_i8x8(uint32_t vui) {
+    const u8vec4 i0 = unpack8( vui       & 0x0F0F0F0F);
+    const u8vec4 i1 = unpack8((vui >> 4) & 0x0F0F0F0F);
+
+    return i32vec2(
+        pack32(i8vec4(kvalues_iq4nl[i0.x], kvalues_iq4nl[i0.y], kvalues_iq4nl[i0.z], kvalues_iq4nl[i0.w])),
+        pack32(i8vec4(kvalues_iq4nl[i1.x], kvalues_iq4nl[i1.y], kvalues_iq4nl[i1.z], kvalues_iq4nl[i1.w])));
+}
+#endif
 #endif
 #endif
 

ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp

@@ -265,7 +265,7 @@ bool is_lut_quant(const std::string& type_name) {
 }
 
 std::string lut_load_vec_a(const std::string& type_name) {
-    if (type_name == "iq1_s" || type_name == "iq1_m" || type_name == "iq2_xxs" || type_name == "iq2_xs" || type_name == "iq2_s") {
+    if (type_name == "iq1_s" || type_name == "iq1_m" || type_name == "iq2_xxs" || type_name == "iq2_xs" || type_name == "iq2_s" || type_name == "iq4_xs") {
         return "8";
     }
     return "4";
@@ -650,7 +650,7 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c
         };
 
 #if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT)
-        if (!f16acc && !coopmat && !coopmat2 && !dot2 && (is_legacy_quant(tname) || is_k_quant(tname) || tname == "mxfp4" || tname == "iq3_s")) {
+        if (!f16acc && !coopmat && !coopmat2 && !dot2 && (is_legacy_quant(tname) || is_k_quant(tname) || tname == "mxfp4" || tname == "iq3_s" || tname == "iq4_xs")) {
             string_to_spv(shader_name + "_" + tname + "_q8_1", "mul_mmq.comp", merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"D_TYPE", "float"},}), fp16, coopmat, coopmat2, f16acc);
         }
 #endif
@@ -833,7 +833,7 @@ void process_shaders() {
 
         // mul mat vec with integer dot product
 #if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT)
-        if (is_legacy_quant(tname) || tname == "mxfp4" || is_k_quant(tname) || tname == "iq1_s" || tname == "iq1_m") {
+        if (is_legacy_quant(tname) || tname == "mxfp4" || is_k_quant(tname) || tname == "iq1_s" || tname == "iq1_m" || tname == "iq4_xs") {
             string_to_spv("mul_mat_vec_" + tname + "_q8_1_f32", "mul_mat_vecq.comp", merge_maps(base_dict, {{data_a_key, "1"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}, {"FLOAT_TYPEV2", "vec2"}, {"ACC_TYPE", "float"}}));
             string_to_spv("mul_mat_vec_" + tname + "_q8_1_f32_subgroup", "mul_mat_vecq.comp", merge_maps(base_dict, {{data_a_key, "1"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}, {"FLOAT_TYPEV2", "vec2"}, {"ACC_TYPE", "float"}, {"USE_SUBGROUP_ADD", "1"}}));
             string_to_spv("mul_mat_vec_" + tname + "_q8_1_f32_subgroup_no_shmem", "mul_mat_vecq.comp", merge_maps(base_dict, {{data_a_key, "1"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}, {"FLOAT_TYPEV2", "vec2"}, {"ACC_TYPE", "float"}, {"USE_SUBGROUP_ADD_NO_SHMEM", "1"}}));
@@ -1357,7 +1357,7 @@ void write_output_files() {
 
     for (const std::string& btype : btypes) {
     for (const auto& tname : type_names) {
-        if (btype == "q8_1" && !is_legacy_quant(tname) && tname != "mxfp4" && !is_k_quant(tname) && tname != "iq1_s" && tname != "iq1_m") {
+        if (btype == "q8_1" && !is_legacy_quant(tname) && tname != "mxfp4" && !is_k_quant(tname) && tname != "iq1_s" && tname != "iq1_m" && tname != "iq4_xs") {
             continue;
         }
         hdr << "extern const void * arr_dmmv_"   << tname << "_" << btype << "_f32_data[3];\n";

gguf-py/gguf/constants.py

@@ -650,6 +650,7 @@ class VISION_PROJECTOR_TYPE(IntEnum):
     GEMMA3N   = auto()
     GEMMA3    = auto()
     QWEN3VL   = auto()
+    LING3VL   = auto()
     STEP3VL   = auto()
     COGVLM    = auto()
 
@@ -1407,6 +1408,7 @@ VISION_PROJECTOR_TYPE_NAMES: dict[VISION_PROJECTOR_TYPE, str] = {
     VISION_PROJECTOR_TYPE.MERGER:    "qwen2vl_merger",
     VISION_PROJECTOR_TYPE.GEMMA3:    "gemma3",
     VISION_PROJECTOR_TYPE.QWEN3VL:   "qwen3vl_merger",
+    VISION_PROJECTOR_TYPE.LING3VL:   "ling3vl",
     VISION_PROJECTOR_TYPE.STEP3VL:   "step3vl",
 }
 
@@ -5207,6 +5209,10 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
         MODEL_TENSOR.D2T,
     ],
     MODEL_ARCH.DFLASH: [
+        MODEL_TENSOR.ATTN_POST_NORM,
+        MODEL_TENSOR.FFN_POST_NORM,
+        MODEL_TENSOR.LAYER_OUT_SCALE,
+        MODEL_TENSOR.ROPE_FREQS,
         MODEL_TENSOR.TOKEN_EMBD,
         MODEL_TENSOR.OUTPUT,
         MODEL_TENSOR.OUTPUT_NORM,
@@ -5825,6 +5831,7 @@ class VisionProjectorType:
     QWEN25VL = "qwen2.5vl_merger"
     EXAONE4_5 = "exaone4_5"
     QWEN3VL = "qwen3vl_merger"
+    LING3VL = "ling3vl"
     STEP3VL = "step3vl"
     ULTRAVOX = "ultravox"
     INTERNVL = "internvl"

src/llama-model.cpp

@@ -3111,7 +3111,6 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
         case LLM_ARCH_GRANITE_SWA:
         case LLM_ARCH_CHAMELEON:
         case LLM_ARCH_BAILINGMOE:
-        case LLM_ARCH_BAILINGMOE3:
         case LLM_ARCH_NEO_BERT:
         case LLM_ARCH_SMOLLM3:
         case LLM_ARCH_ARCEE:
@@ -3126,6 +3125,10 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
         case LLM_ARCH_DOTS3NOTE:
         case LLM_ARCH_NANBEIGE:
         case LLM_ARCH_POCKETTTS:
+            return LLAMA_ROPE_TYPE_NORM;
+        case LLM_ARCH_BAILINGMOE3:
+            // VL files carry mrope sections; text-only files keep NORM rope
+            return model->hparams.use_mrope() ? LLAMA_ROPE_TYPE_MROPE : LLAMA_ROPE_TYPE_NORM;
         // HY_V4 rotates consecutive pairs, matching the reference implementation
         case LLM_ARCH_HY_V4:
             return LLAMA_ROPE_TYPE_NORM;

src/llama-sampler.cpp

@@ -644,7 +644,7 @@ static bool llama_sampler_backend_support(
         return true;
     }
 
-    auto probe = llama_sampler_backend_probe_graph(smpl, 1024*1024, GGML_DEFAULT_GRAPH_SIZE, true);
+    auto probe = llama_sampler_backend_probe_graph(smpl, 128*1024, GGML_DEFAULT_GRAPH_SIZE, true);
 
     for (int i = 0; i < ggml_graph_n_nodes(probe.gf); i++) {
         struct ggml_tensor * op = ggml_graph_node(probe.gf, i);
@@ -764,7 +764,7 @@ static bool llama_sampler_chain_backend_init(
         res = res && cur_prefix;
     }
 
-    auto probe = llama_sampler_backend_probe_graph(smpl, 1024*1024, GGML_DEFAULT_GRAPH_SIZE, false);
+    auto probe = llama_sampler_backend_probe_graph(smpl, 128*1024, GGML_DEFAULT_GRAPH_SIZE, false);
     chain->n_nodes = llama_sampler_backend_probe_n_nodes(probe);
 
     return res;

src/models/bailingmoe3.cpp

@@ -15,6 +15,7 @@ void llama_model_bailingmoe3::load_arch_hparams(llama_model_loader & ml) {
         hparams.kda_safe_gate = true;
     }
     ml.get_key(LLM_KV_KDA_GATE_LOWER_BOUND,             hparams.kda_gate_lower_bound);
+    ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS, hparams.rope_sections, 4, false);
     ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
     ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
     ml.get_key(LLM_KV_EXPERT_SHARED_COUNT,              hparams.n_expert_shared);
@@ -233,6 +234,10 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph
     const int64_t d_conv = hparams.ssm_d_conv;
     const int64_t n_seqs = ubatch.n_seqs;
     const int64_t n_seq_tokens = ubatch.n_seq_tokens;
+
+    const bool use_mrope = hparams.use_mrope();
+    int sections[4];
+    std::copy(std::begin(hparams.rope_sections), std::begin(hparams.rope_sections) + 4, sections);
     const int64_t qk_head_dim = hparams.n_embd_head_k_mla();
     const int64_t v_head_dim = hparams.n_embd_head_v_mla();
     const int64_t qk_rope_head_dim = hparams.n_rot();
@@ -326,10 +331,17 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph
                     ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim),
                     ggml_row_size(kv_all->type, kv_lora_rank));
 
-            q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+            if (use_mrope) {
-                    ext_factor, attn_factor, beta_fast, beta_slow);
+                q_pe = ggml_rope_multi(ctx0, q_pe, inp_pos, nullptr, n_rot, sections, rope_type,
-            k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+                        n_ctx_orig, freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow);
-                    ext_factor, attn_factor, beta_fast, beta_slow);
+                k_pe = ggml_rope_multi(ctx0, k_pe, inp_pos, nullptr, n_rot, sections, rope_type,
+                        n_ctx_orig, freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow);
+            } else {
+                q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+                        ext_factor, attn_factor, beta_fast, beta_slow);
+                k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+                        ext_factor, attn_factor, beta_fast, beta_slow);
+            }
             kv = build_norm(kv, layer.attn_kv_a_norm, nullptr, LLM_NORM_RMS, il);
 
             q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3);
@@ -482,10 +494,21 @@ llama_model_bailingmoe3::graph_mtp::graph_mtp(const llama_model & model, const l
             ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim),
             ggml_row_size(kv_all->type, kv_lora_rank));
 
-    q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+    const bool use_mrope = hparams.use_mrope();
-            ext_factor, attn_factor, beta_fast, beta_slow);
+    int sections[4];
-    k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+    std::copy(std::begin(hparams.rope_sections), std::begin(hparams.rope_sections) + 4, sections);
-            ext_factor, attn_factor, beta_fast, beta_slow);
+
+    if (use_mrope) {
+        q_pe = ggml_rope_multi(ctx0, q_pe, inp_pos, nullptr, n_rot, sections, rope_type,
+                n_ctx_orig, freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow);
+        k_pe = ggml_rope_multi(ctx0, k_pe, inp_pos, nullptr, n_rot, sections, rope_type,
+                n_ctx_orig, freq_base, freq_scale, ext_factor, attn_factor, beta_fast, beta_slow);
+    } else {
+        q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+                ext_factor, attn_factor, beta_fast, beta_slow);
+        k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
+                ext_factor, attn_factor, beta_fast, beta_slow);
+    }
     kv = build_norm(kv, layer.attn_kv_a_norm, nullptr, LLM_NORM_RMS, il);
 
     q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3);

src/models/dflash.cpp

@@ -6,6 +6,19 @@
 
 void llama_model_dflash::load_arch_hparams(llama_model_loader & ml) {
 
+    ml.get_key(LLM_KV_EMBEDDING_SCALE, hparams.f_embedding_scale, false);
+    ml.get_key(LLM_KV_ATTENTION_SCALE, hparams.f_attention_scale, false);
+
+    hparams.llm_ffn_op = LLM_FFN_SILU;
+    std::string hidden_act;
+    if (ml.get_key(LLM_KV_HIDDEN_ACT, hidden_act, false)) {
+        if (hidden_act == "gelu" || hidden_act == "gelu_pytorch_tanh") {
+            hparams.llm_ffn_op = LLM_FFN_GELU;
+        } else if (hidden_act != "silu") {
+            throw std::runtime_error("unsupported DFlash hidden activation: " + hidden_act);
+        }
+    }
+
     ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
     ml.get_key(LLM_KV_LOGIT_SCALE,                 hparams.f_logit_scale, false);
     hparams.f_final_logit_softcapping = 0.0f;
@@ -108,9 +121,6 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
     }
 
     // DSpark = DFlash + a semi-autoregressive Markov head and Confidence head
-    //
-    // TODO: only Qwen3-style backbones are supported for now; other backbones (e.g. Gemma4)
-    //       need their own conversion path and graph tweaks
     const struct ggml_tensor * markov_meta = ml->get_tensor_meta("markov_w1.weight");
     if (markov_meta) {
         const int64_t dspark_markov_rank = markov_meta->ne[0];
@@ -156,6 +166,9 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
     // optional: reduced-vocab drafts ship their own lm head, full-vocab drafts can share the target's via ctx_other
     // a draft with its own embeddings + head references no target tensors and can run on devices the target does not use (e.g. -devd with a tensor-split target)
     output   = create_tensor(tn(LLM_TENSOR_OUTPUT,     "weight"), { n_embd, n_vocab_draft }, TENSOR_NOT_REQUIRED);
+    if (output == nullptr && tok_embd != nullptr) {
+        output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab_draft }, TENSOR_DUPLICATED);
+    }
 
     if (hparams.dsv4_hc_mult > 0) {
         const int64_t q_lora_rank     = hparams.n_lora_q;
@@ -214,12 +227,17 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
 
         layer.wq = create_tensor(tn(LLM_TENSOR_ATTN_Q,   "weight", i), { n_embd, n_embd_head_k * n_head }, 0);
         layer.wk = create_tensor(tn(LLM_TENSOR_ATTN_K,   "weight", i), { n_embd, n_embd_k_gqa }, 0);
-        layer.wv = create_tensor(tn(LLM_TENSOR_ATTN_V,   "weight", i), { n_embd, n_embd_v_gqa }, 0);
+        layer.wv = create_tensor(tn(LLM_TENSOR_ATTN_V,   "weight", i), { n_embd, n_embd_v_gqa }, TENSOR_NOT_REQUIRED);
         layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), { n_embd_head_k * n_head, n_embd }, 0);
 
         layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), { n_embd_head_k }, 0);
         layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), { n_embd_head_k }, 0);
 
+        layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", i), { n_embd }, TENSOR_NOT_REQUIRED);
+        layer.ffn_post_norm  = create_tensor(tn(LLM_TENSOR_FFN_POST_NORM,  "weight", i), { n_embd }, TENSOR_NOT_REQUIRED);
+        layer.out_scale      = create_tensor(tn(LLM_TENSOR_LAYER_OUT_SCALE, "weight", i), { 1 }, TENSOR_NOT_REQUIRED);
+        layer.rope_freqs     = create_tensor(tn(LLM_TENSOR_ROPE_FREQS, "weight", i), { n_embd_head_k/2 }, TENSOR_NOT_REQUIRED | (i > 0 ? TENSOR_DUPLICATED : 0));
+
         // optional per-head attention sinks (e.g. Nemotron DSpark)
         layer.attn_sinks = create_tensor(tn(LLM_TENSOR_ATTN_SINKS, "weight", i), { n_head }, TENSOR_NOT_REQUIRED);
 
@@ -571,7 +589,7 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
         inp_attn = build_attn_inp_kv();
     }
 
-    const float kq_scale = 1.0f/sqrtf(float(n_embd_head));
+    const float kq_scale = hparams.f_attention_scale != 0.0f ? hparams.f_attention_scale : 1.0f/sqrtf(float(n_embd_head));
 
     // drafts for M-RoPE targets use degenerate sections (temporal dim only)
     int sections[4];
@@ -582,7 +600,7 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
             ? ggml_rope_multi(ctx0, cur, pos, nullptr,
                     n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale,
                     ext_factor, attn_factor, beta_fast, beta_slow)
-            : ggml_rope_ext(ctx0, cur, pos, nullptr,
+            : ggml_rope_ext(ctx0, cur, pos, model.layers[0].rope_freqs,
                     n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
                     ext_factor, attn_factor, beta_fast, beta_slow);
     };
@@ -608,12 +626,16 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
             const auto & layer = model.layers[il];
 
             ggml_tensor * Kcur = build_lora_mm(layer.wk, inp_g, layer.wk_s);
-            ggml_tensor * Vcur = build_lora_mm(layer.wv, inp_g, layer.wv_s);
+            const bool shared_kv = layer.wv == nullptr;
+            ggml_tensor * Vcur = shared_kv ? Kcur : build_lora_mm(layer.wv, inp_g, layer.wv_s);
 
             Kcur = ggml_reshape_3d(ctx0, Kcur, n_embd_head, n_head_kv, n_tokens);
             Vcur = ggml_reshape_3d(ctx0, Vcur, n_embd_head, n_head_kv, n_tokens);
 
             Kcur = build_norm(Kcur, layer.attn_k_norm, NULL, LLM_NORM_RMS, il);
+            if (shared_kv) {
+                Vcur = ggml_rms_norm(ctx0, Vcur, hparams.f_norm_rms_eps);
+            }
             Kcur = build_rope(Kcur, inp_pos);
             cb(Kcur, "Kcur_injected", il);
             cb(Vcur, "Vcur_injected", il);
@@ -673,6 +695,9 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
     ggml_tensor * inp_tokens = inp->tokens;
 
     ggml_tensor * inpL = ggml_get_rows(ctx0, tok_embd, inp->tokens);
+    if (hparams.f_embedding_scale != 0.0f) {
+        inpL = ggml_scale(ctx0, inpL, hparams.f_embedding_scale);
+    }
     cb(inpL, "inp_noise_embd", -1);
 
     res->add_input(std::move(inp));
@@ -692,7 +717,8 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
 
         ggml_tensor * Qcur = build_lora_mm(layer.wq, noise_norm, layer.wq_s);
         ggml_tensor * Kcur = build_lora_mm(layer.wk, noise_norm, layer.wk_s);
-        ggml_tensor * Vcur = build_lora_mm(layer.wv, noise_norm, layer.wv_s);
+        const bool shared_kv = layer.wv == nullptr;
+        ggml_tensor * Vcur = shared_kv ? Kcur : build_lora_mm(layer.wv, noise_norm, layer.wv_s);
 
         Qcur = ggml_reshape_3d(ctx0, Qcur, n_embd_head, n_head,    n_tokens);
         Kcur = ggml_reshape_3d(ctx0, Kcur, n_embd_head, n_head_kv, n_tokens);
@@ -700,6 +726,9 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
 
         Qcur = build_norm(Qcur, layer.attn_q_norm, NULL, LLM_NORM_RMS, il);
         Kcur = build_norm(Kcur, layer.attn_k_norm, NULL, LLM_NORM_RMS, il);
+        if (shared_kv) {
+            Vcur = ggml_rms_norm(ctx0, Vcur, hparams.f_norm_rms_eps);
+        }
 
         Qcur = build_rope(Qcur, inp_pos);
         Kcur = build_rope(Kcur, inp_pos);
@@ -717,6 +746,11 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
             cb(cur, "attn_conv_out", il);
         }
 
+        if (layer.attn_post_norm) {
+            cur = build_norm(cur, layer.attn_post_norm, NULL, LLM_NORM_RMS, il);
+            cb(cur, "attn_post_norm", il);
+        }
+
         ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpL);
         cb(ffn_inp, "ffn_inp", il);
 
@@ -735,7 +769,7 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
                 layer.ffn_gate, NULL, layer.ffn_gate_s,
                 layer.ffn_down, NULL, layer.ffn_down_s,
                 NULL,
-                LLM_FFN_SILU, LLM_FFN_PAR, il);
+                hparams.llm_ffn_op, LLM_FFN_PAR, il);
         cb(cur, "ffn_out", il);
 
         if (ffn_dynamic) {
@@ -743,7 +777,15 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
             cb(cur, "ffn_conv_out", il);
         }
 
+        if (layer.ffn_post_norm) {
+            cur = build_norm(cur, layer.ffn_post_norm, NULL, LLM_NORM_RMS, il);
+            cb(cur, "ffn_post_norm", il);
+        }
+
         cur = ggml_add(ctx0, cur, ffn_inp);
+        if (layer.out_scale) {
+            cur = ggml_mul(ctx0, cur, layer.out_scale);
+        }
         cb(cur, "l_out", il);
 
         inpL = cur;

tools/mtmd/clip-impl.h

@@ -450,6 +450,7 @@ enum projector_type {
     PROJECTOR_TYPE_GLM_EDGE,
     PROJECTOR_TYPE_QWEN2VL,
     PROJECTOR_TYPE_QWEN3VL,
+    PROJECTOR_TYPE_LING3VL,
     PROJECTOR_TYPE_STEP3VL,
     PROJECTOR_TYPE_GEMMA3,
     PROJECTOR_TYPE_GEMMA3NV,
@@ -516,6 +517,7 @@ static std::map<projector_type, std::string> PROJECTOR_TYPE_NAMES = {
     { PROJECTOR_TYPE_QWEN2VL,           "qwen2vl_merger"},
     { PROJECTOR_TYPE_QWEN25VL,          "qwen2.5vl_merger"},
     { PROJECTOR_TYPE_QWEN3VL,           "qwen3vl_merger"},
+    { PROJECTOR_TYPE_LING3VL,           "ling3vl"},
     { PROJECTOR_TYPE_STEP3VL,           "step3vl"},
     { PROJECTOR_TYPE_GEMMA3,            "gemma3"},
     { PROJECTOR_TYPE_GEMMA3NV,          "gemma3nv"},

tools/mtmd/clip.cpp

@@ -1035,6 +1035,10 @@ static std::unique_ptr<clip_graph> clip_get_graph_builder(clip_ctx * ctx, const
             {
                 builder = std::make_unique<clip_graph_qwen3vl>(ctx, img);
             } break;
+        case PROJECTOR_TYPE_LING3VL:
+            {
+                builder = std::make_unique<clip_graph_ling3vl>(ctx, img);
+            } break;
         case PROJECTOR_TYPE_EXAONE4_5:
             {
                 builder = std::make_unique<clip_graph_exaone4_5>(ctx, img);
@@ -1743,6 +1747,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_QWEN2VL:
                 case PROJECTOR_TYPE_QWEN25VL:
                 case PROJECTOR_TYPE_QWEN3VL:
+                case PROJECTOR_TYPE_LING3VL:
                     {
                         hparams.n_merge = 2; // default value for Qwen 2 and 2.5
                         hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
@@ -2580,6 +2585,15 @@ struct clip_model_loader {
                     model.mm_1_w = get_tensor(string_format(TN_LLAVA_PROJ, 2, "weight"));
                     model.mm_1_b = get_tensor(string_format(TN_LLAVA_PROJ, 2, "bias"));
                 } break;
+            case PROJECTOR_TYPE_LING3VL:
+                {
+                    model.mm_input_norm_w = get_tensor(TN_MM_INP_NORM);        // merger.norm
+                    model.mm_input_norm_b = get_tensor(TN_MM_INP_NORM_B);     // merger.norm
+                    model.mm_0_w = get_tensor(string_format(TN_LLAVA_PROJ, 0, "weight"));  // linear_proj.0
+                    model.mm_0_b = get_tensor(string_format(TN_LLAVA_PROJ, 0, "bias"));
+                    model.mm_1_w = get_tensor(string_format(TN_LLAVA_PROJ, 2, "weight"));  // linear_proj.2
+                    model.mm_1_b = get_tensor(string_format(TN_LLAVA_PROJ, 2, "bias"));
+                } break;
             case PROJECTOR_TYPE_MIMOVL:
                 {
                     model.mm_0_w = get_tensor(string_format(TN_LLAVA_PROJ, 0, "weight"));
@@ -4141,6 +4155,7 @@ int clip_n_output_tokens_x(const clip_ctx * ctx, const clip_image_f32 * img) {
         case PROJECTOR_TYPE_QWEN2VL:
         case PROJECTOR_TYPE_QWEN25VL:
         case PROJECTOR_TYPE_QWEN3VL:
+        case PROJECTOR_TYPE_LING3VL:
         case PROJECTOR_TYPE_EXAONE4_5:
         case PROJECTOR_TYPE_MIMOVL:
         case PROJECTOR_TYPE_GLM4V:
@@ -4167,6 +4182,7 @@ int clip_n_output_tokens_y(const clip_ctx * ctx, const clip_image_f32 * img) {
         case PROJECTOR_TYPE_QWEN2VL:
         case PROJECTOR_TYPE_QWEN25VL:
         case PROJECTOR_TYPE_QWEN3VL:
+        case PROJECTOR_TYPE_LING3VL:
         case PROJECTOR_TYPE_EXAONE4_5:
         case PROJECTOR_TYPE_MIMOVL:
         case PROJECTOR_TYPE_GLM4V:
@@ -4247,6 +4263,7 @@ int clip_n_output_tokens(const clip_ctx * ctx, const clip_image_f32 * img) {
         case PROJECTOR_TYPE_QWEN2VL:
         case PROJECTOR_TYPE_QWEN25VL:
         case PROJECTOR_TYPE_QWEN3VL:
+        case PROJECTOR_TYPE_LING3VL:
         case PROJECTOR_TYPE_EXAONE4_5:
         case PROJECTOR_TYPE_MIMOVL:
         case PROJECTOR_TYPE_MINIMAX_M3:
@@ -4887,6 +4904,7 @@ bool clip_encode(struct clip_ctx * ctx, struct clip_encode_params * params) {
             } break;
         case PROJECTOR_TYPE_QWEN2VL:
         case PROJECTOR_TYPE_QWEN3VL:
+        case PROJECTOR_TYPE_LING3VL:
         case PROJECTOR_TYPE_GLM4V:
             {
                 const int merge_ratio = hparams.n_merge;
@@ -6051,6 +6069,8 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) {
         case PROJECTOR_TYPE_QWEN3VL:
             // main path + deepstack paths
             return ctx->model.mm_1_b->ne[0] * (1 + ctx->model.n_deepstack_layers);
+        case PROJECTOR_TYPE_LING3VL:
+            return ctx->model.mm_1_b->ne[0];
         case PROJECTOR_TYPE_MIMOVL:
             return ctx->model.mm_1_w->ne[1];
         case PROJECTOR_TYPE_STEP3VL:
@@ -6144,6 +6164,7 @@ int clip_model_n_temporal_merge(const struct clip_ctx * ctx) {
         case PROJECTOR_TYPE_QWEN2VL:
         case PROJECTOR_TYPE_QWEN25VL:
         case PROJECTOR_TYPE_QWEN3VL:
+        case PROJECTOR_TYPE_LING3VL:
             return 2;
         default:
             return 1;

tools/mtmd/models/ling3vl.cpp

@@ -0,0 +1,86 @@
+#include "models.h"
+
+ggml_cgraph * clip_graph_ling3vl::build() {
+    // same vision tower as qwen3vl, but the merger is norm-only (no fc1/fc2) and
+    // the projector MLP lives at the top level (mm.0 / mm.2)
+    GGML_ASSERT(model.class_embedding == nullptr);
+    GGML_ASSERT(model.mm_input_norm_w != nullptr); // merger norm (pre spatial merge)
+
+    const int batch_size = 1;
+    const int n_pos      = n_patches;
+
+    norm_type norm_t = NORM_TYPE_NORMAL;
+
+    // vision M-RoPE, same layout as qwen3vl: [row, col, row, col] quarters
+    int mrope_sections[4] = {d_head/4, d_head/4, d_head/4, d_head/4};
+
+    ggml_tensor * inp = build_inp_with_temporal_merge();
+
+    // spatial merge
+    {
+        inp = ggml_permute(ctx0, inp, 1, 2, 0, 3);  // [w, h, c, b] -> [c, w, h, b]
+        inp = ggml_cont_4d(
+            ctx0, inp,
+            n_embd * 2, n_patches_x / 2, n_patches_y, batch_size);
+        inp = ggml_reshape_4d(
+            ctx0, inp,
+            n_embd * 2, n_patches_x / 2, 2, batch_size * (n_patches_y / 2));
+        inp = ggml_permute(ctx0, inp, 0, 2, 1, 3);
+        inp = ggml_cont_3d(
+            ctx0, inp,
+            n_embd, n_patches_x * n_patches_y, batch_size);
+    }
+
+    // add patch bias
+    if (model.patch_bias != nullptr) {
+        inp = ggml_add(ctx0, inp, model.patch_bias);
+        cb(inp, "patch_bias", -1);
+    }
+
+    // calculate absolute position embedding and apply
+    ggml_tensor * learned_pos_embd = resize_position_embeddings(GGML_SCALE_MODE_BILINEAR | GGML_SCALE_FLAG_ALIGN_CORNERS);
+    learned_pos_embd = ggml_cont_4d(
+        ctx0, learned_pos_embd,
+        n_embd * 2, n_patches_x / 2, n_patches_y, batch_size);
+    learned_pos_embd = ggml_reshape_4d(
+        ctx0, learned_pos_embd,
+        n_embd * 2, n_patches_x / 2, 2, batch_size * (n_patches_y / 2));
+    learned_pos_embd = ggml_permute(ctx0, learned_pos_embd, 0, 2, 1, 3);
+    learned_pos_embd = ggml_cont_3d(
+        ctx0, learned_pos_embd,
+        n_embd, n_patches_x * n_patches_y, batch_size);
+
+    const int num_position_ids = n_pos * 4; // m-rope requires 4 dim per position
+    ggml_tensor * positions = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, num_position_ids);
+    ggml_set_name(positions, "positions");
+    ggml_set_input(positions);
+
+    ggml_tensor * inpL = build_vit(
+        inp, n_pos, norm_t, hparams.ffn_op, learned_pos_embd,
+        [&](ggml_tensor * c, const clip_layer &) {
+            return ggml_rope_multi(
+                ctx0, c, positions, nullptr,
+                d_head/2, mrope_sections, GGML_ROPE_TYPE_VISION, 32768, 10000, 1, 0, 1, 32, 1);
+        });
+
+    // multimodal projection (linear_proj MLP over the merged patches)
+    ggml_tensor * embeddings = inpL;
+
+    // per-patch merger norm, applied post-blocks before the 2x2 merge
+    // (merger.norm, LayerNorm over n_embd)
+    embeddings = build_norm(embeddings, model.mm_input_norm_w, model.mm_input_norm_b, norm_t, eps, -1);
+    cb(embeddings, "merger_norm", -1);
+
+    embeddings = ggml_reshape_3d(ctx0, embeddings, n_embd * 4, n_pos / 4, batch_size);
+
+    embeddings = build_ffn(embeddings,
+        model.mm_0_w, model.mm_0_b,
+        nullptr, nullptr,
+        model.mm_1_w, model.mm_1_b,
+        ffn_op_type::FFN_GELU, -1);
+
+    // build the graph
+    ggml_build_forward_expand(gf, embeddings);
+
+    return gf;
+}

tools/mtmd/models/models.h

@@ -50,6 +50,11 @@ struct clip_graph_qwen3vl : clip_graph_qwen2vl {
     ggml_cgraph * build() override;
 };
 
+struct clip_graph_ling3vl : clip_graph_qwen3vl {
+    clip_graph_ling3vl(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph_qwen3vl(ctx, img) {}
+    ggml_cgraph * build() override;
+};
+
 struct clip_graph_minimax_m3 : clip_graph {
     clip_graph_minimax_m3(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
     ggml_cgraph * build() override;

tools/mtmd/mtmd.cpp

@@ -694,6 +694,7 @@ struct mtmd_context {
             case PROJECTOR_TYPE_QWEN2VL:
             case PROJECTOR_TYPE_QWEN25VL:
             case PROJECTOR_TYPE_QWEN3VL:
+            case PROJECTOR_TYPE_LING3VL:
             case PROJECTOR_TYPE_MIMOVL:
                 {
                     // <|vision_start|> ... (image embeddings) ... <|vision_end|>

tools/server/server-chat.cpp

@@ -203,10 +203,24 @@ json server_chat_convert_responses_to_chatcmpl(const json & response_body) {
                 } else {
                     json chatcmpl_outputs = item.at("output");
                     for (json & chatcmpl_output : chatcmpl_outputs) {
-                        if (!chatcmpl_output.contains("type") || chatcmpl_output.at("type") != "input_text") {
+                        if (!chatcmpl_output.contains("type")) {
-                            throw std::invalid_argument("Output of tool call should be 'Input text'");
+                            throw std::invalid_argument("Output of tool call missing 'type' field");
+                        }
+                        const auto type = chatcmpl_output.at("type");
+                        if (type != "input_text" && type != "input_image") {
+                            throw std::invalid_argument("Output of tool call should be 'Input text' or 'Input image'");
+                        }
+                        if (type == "input_text") {
+                            chatcmpl_output["type"] = "text";
+                        } else if (type == "input_image") {
+                            if (!chatcmpl_output.contains("image_url")) {
+                                throw std::invalid_argument("'image_url' is required");
+                            }
+                            chatcmpl_output["type"] = "image_url";
+                            chatcmpl_output["image_url"] = json {
+                                {"url", chatcmpl_output.at("image_url")}
+                            };
                         }
-                        chatcmpl_output["type"] = "text";
                     }
                     chatcmpl_messages.push_back(json {
                         {"content",      chatcmpl_outputs},

tools/server/server-common.cpp

@@ -1263,12 +1263,13 @@ json oaicompat_chat_params_parse(
                 p["text"] = get_media_marker();
                 p.erase("input_audio");
 
-            } else if (type == "input_video") {
+            } else if (type == "input_video" || type == "video_url") {
                 if (!opt.allow_video) {
                     throw std::runtime_error("video input is not supported - hint: if this is unexpected, you may need to provide the mmproj");
                 }
 
-                json input_video = json_value(p, "input_video", json::object());
+                // accept the OpenAI-style "video_url" key as an alias of "input_video"
+                json input_video = json_value(p, type, json::object());
                 std::string url  = json_value(input_video, "data",
                                         json_value(input_video, "url", std::string()));
                 handle_media(out_files, url, opt.media_path);
@@ -1276,6 +1277,7 @@ json oaicompat_chat_params_parse(
                 p["type"] = "media_marker";
                 p["text"] = get_media_marker();
                 p.erase("input_video");
+                p.erase("video_url");
 
             } else if (type != "text") {
                 throw std::invalid_argument("unsupported content[].type");

tools/server/server-context.cpp

@@ -2060,7 +2060,7 @@ private:
             res->is_begin = true;
         } else {
             res->content = tkn.text_to_send;
-            res->tokens  = { tkn.tok };
+            res->tokens.assign(1, tkn.tok);
         }
 
         res->n_decoded             = slot.stats.n_gen;
@@ -4944,7 +4944,7 @@ void server_routes::init_routes() {
     };
 
     this->post_chat_completions_tok = [this](const server_http_req & req) {
-        return handle_count_tokens(ctx_server.vocab, ctx_server.mctx, ctx_server.init_opt, req, TASK_RESPONSE_TYPE_OAI_CHAT);
+        return handle_count_tokens(req, TASK_RESPONSE_TYPE_OAI_CHAT);
     };
 
     this->post_control = [this](const server_http_req & req) {
@@ -5003,7 +5003,7 @@ void server_routes::init_routes() {
     };
 
     this->post_responses_tok_oai = [this](const server_http_req & req) {
-        return handle_count_tokens(ctx_server.vocab, ctx_server.mctx, ctx_server.init_opt, req, TASK_RESPONSE_TYPE_OAI_RESP);
+        return handle_count_tokens(req, TASK_RESPONSE_TYPE_OAI_RESP);
     };
 
     this->post_transcriptions_oai = [this](const server_http_req & req) {
@@ -5053,7 +5053,7 @@ void server_routes::init_routes() {
     };
 
     this->post_anthropic_count_tokens = [this](const server_http_req & req) {
-        return handle_count_tokens(ctx_server.vocab, ctx_server.mctx, ctx_server.init_opt, req, TASK_RESPONSE_TYPE_ANTHROPIC);
+        return handle_count_tokens(req, TASK_RESPONSE_TYPE_ANTHROPIC);
     };
 
     // same with handle_chat_completions, but without inference part
@@ -5485,7 +5485,7 @@ std::unique_ptr<server_res_generator> server_routes::handle_embeddings_impl(cons
     return res;
 }
 
-std::unique_ptr<server_res_generator> server_routes::handle_count_tokens(const llama_vocab * vocab, mtmd_context * mctx, const mtmd_helper_init_opt & init_opt, const server_http_req & req, task_response_type res_type) {
+std::unique_ptr<server_res_generator> server_routes::handle_count_tokens(const server_http_req & req, task_response_type res_type) {
     auto res = create_response();
     std::vector<raw_buffer> files;
     json body = json::parse(req.body);
@@ -5519,13 +5519,13 @@ std::unique_ptr<server_res_generator> server_routes::handle_count_tokens(const l
 
     // TODO @ngxson : refactor this code block, move this to server-common and reuse it in other places
     size_t n_tokens;
-    if (mctx != nullptr) {
+    if (ctx_server.mctx != nullptr) {
         if (!prompt.is_string()) {
             throw std::runtime_error("for mtmd, input prompt must be a string.");
         }
-        n_tokens = process_mtmd_prompt(mctx, prompt.get<std::string>(), files, init_opt, true).size();
+        n_tokens = process_mtmd_prompt(ctx_server.mctx, prompt.get<std::string>(), files, ctx_server.init_opt, true).size();
     } else {
-        n_tokens = tokenize_mixed(vocab, prompt, true, true).size();
+        n_tokens = tokenize_mixed(ctx_server.vocab, prompt, true, true).size();
     }
 
     json response = {{"input_tokens", static_cast<int64_t>(n_tokens)}};

tools/server/server-context.h

@@ -169,7 +169,7 @@ private:
     std::unique_ptr<server_res_generator> handle_slots_restore(const server_http_req & req, int id_slot);
     std::unique_ptr<server_res_generator> handle_slots_erase(const server_http_req &, int id_slot);
     std::unique_ptr<server_res_generator> handle_embeddings_impl(const server_http_req & req, task_response_type res_type);
-    std::unique_ptr<server_res_generator> handle_count_tokens(const llama_vocab * vocab, mtmd_context * mctx, const mtmd_helper_init_opt & init_opt, const server_http_req & req, task_response_type res_type);
+    std::unique_ptr<server_res_generator> handle_count_tokens(const server_http_req & req, task_response_type res_type);
 
     // using unique_ptr to allow late initialization of const
     std::unique_ptr<const server_context_meta> meta;

tools/server/server-models.cpp

@@ -297,7 +297,7 @@ struct server_lru_sched {
             return;
         }
         queue.push_back({ model_id, 1, false });
-        SRV_INF("models_max reached, request for name=%s queued at position %zu\n",
+        SRV_INF("request for name=%s queued at position %zu\n",
                 model_id.c_str(), queue.size());
     }
 
@@ -584,8 +584,12 @@ server_models::server_models(
               base_preset(ctx_preset.load_from_args(argc, argv)),
               sched(std::make_unique<server_lru_sched>(*this)),
               monitor(std::make_unique<server_monitor>(*this)) {
-    // clean up base preset
+    // propagate base params to child
     unset_reserved_args(base_preset, true);
+
+    // do not propagate these options, but allow preset to explicitly set them
+    base_preset.unset_option("LLAMA_ARG_LOG_FILE");
+
     // set binary path
     try {
         bin_path = get_server_exec_path().string();
@@ -733,21 +737,25 @@ void server_models::load_models() {
     std::set<std::string> hidden_models;
     {
         std::set<std::string> preset_paths;
+        auto add_hf_path = [&preset_paths](const common_preset & preset, const char * repo_key, const char * file_key) {
+            std::string hf_repo;
+            if (!preset.get_option(repo_key, hf_repo) || hf_repo.empty()) {
+                return;
+            }
+            std::string hf_file;
+            preset.get_option(file_key, hf_file);
+            std::string path = common_download_resolve_path(hf_repo, hf_file);
+            if (!path.empty()) {
+                preset_paths.insert(path);
+            }
+        };
         for (const auto & [name, preset] : custom_presets) {
             std::string val;
             if (!preset.get_option(COMMON_ARG_PRESET_DEDUP_CACHE_MODELS, val) || !common_arg_utils::is_truthy(val)) {
                 continue;
             }
-            std::string hf_repo;
+            add_hf_path(preset, "LLAMA_ARG_HF_REPO", "LLAMA_ARG_HF_FILE");
-            if (!preset.get_option("LLAMA_ARG_HF_REPO", hf_repo) || hf_repo.empty()) {
+            add_hf_path(preset, "LLAMA_ARG_SPEC_DRAFT_HF_REPO", "LLAMA_ARG_SPEC_DRAFT_MODEL");
-                continue;
-            }
-            std::string hf_file;
-            preset.get_option("LLAMA_ARG_HF_FILE", hf_file);
-            std::string path = common_download_resolve_path(hf_repo, hf_file);
-            if (!path.empty()) {
-                preset_paths.insert(path);
-            }
         }
         if (!preset_paths.empty()) {
             for (const auto & [name, preset] : cached_models) {
@@ -1222,15 +1230,16 @@ void server_models::request_stop(const std::string & name, bool send_exit) {
 void server_models::on_child_exit(const std::string & name, const std::shared_ptr<server_subproc> & proc, server_child_mode mode, int exit_code) {
     {
         std::lock_guard<std::mutex> lk(mutex);
-        stopping_models.erase(name);
         auto it = mapping.find(name);
         if (it == mapping.end() || it->second.subproc != proc) {
+            stopping_models.erase(name);
             return; // entry erased, or a newer instance took the name
         }
     }
     if (mode == SERVER_CHILD_MODE_DOWNLOAD) {
         // instance will be cleaned up on next load_models() call
         std::lock_guard<std::mutex> lk(mutex);
+        stopping_models.erase(name);
         cv.notify_all();
     } else {
         update_status(name, {
@@ -1300,6 +1309,9 @@ void server_models::update_status(const std::string & name, const update_status_
         auto & meta = it->second.meta;
         meta.status      = args.status;
         meta.exit_code   = args.exit_code;
+        if (args.status == SERVER_MODEL_STATUS_UNLOADED) {
+            stopping_models.erase(name);
+        }
         if (!args.loaded_info.is_null()) {
             meta.loaded_info = args.loaded_info;
         }
@@ -1439,10 +1451,15 @@ bool server_models::ensure_model_ready(const std::string & name, const std::func
     if (!meta.has_value()) {
         throw std::runtime_error("model name=" + name + " is not found");
     }
-    if (meta->is_ready()) {
+    bool stopping;
+    {
+        std::lock_guard<std::mutex> lk(mutex);
+        stopping = stopping_models.count(name) > 0;
+    }
+    if (!stopping && meta->is_ready()) {
         return false; // ready for taking requests
     }
-    if (meta->status == SERVER_MODEL_STATUS_SLEEPING) {
+    if (!stopping && meta->status == SERVER_MODEL_STATUS_SLEEPING) {
         return false; // child is sleeping but still running; new request will wake it up
     }
 
@@ -1452,17 +1469,10 @@ bool server_models::ensure_model_ready(const std::string & name, const std::func
         std::unique_lock<std::mutex> lk(mutex);
         auto it = mapping.find(name);
         if (it != mapping.end() && it->second.meta.status == SERVER_MODEL_STATUS_UNLOADED) {
-            if (sched->has_capacity(lk) && sched->queue_empty(lk)) {
+            // the queue entry protects the model from eviction until its waiters leave
-                lk.unlock();
+            sched->join(lk, name);
-                SRV_INF("model name=%s is not loaded, loading...\n", name.c_str());
+            sched->tick(lk);
-                load(name);
+            queued = true;
-                did_load = true;
-            } else {
-                // also queue when a slot looks free but others wait already, else they starve
-                sched->join(lk, name);
-                sched->tick(lk);
-                queued = true;
-            }
         }
     }
 
@@ -1483,6 +1493,19 @@ bool server_models::ensure_model_ready(const std::string & name, const std::func
             if (it == mapping.end()) {
                 break; // removed by another code path, nothing to wait for
             }
+            if (stopping_models.count(name)) {
+                // a stopping instance takes no new request, the next instance serves it
+                if (!queued) {
+                    sched->join(lk, name);
+                    sched->tick(lk);
+                    queued = true;
+                }
+                if (should_stop && should_stop()) {
+                    throw std::runtime_error("request cancelled while waiting for model name=" + name);
+                }
+                cv.wait_for(lk, std::chrono::milliseconds(200));
+                continue;
+            }
             const server_model_status status = it->second.meta.status;
 
             if (status == SERVER_MODEL_STATUS_LOADED || status == SERVER_MODEL_STATUS_SLEEPING) {

tools/server/tests/unit/test_router.py

@@ -433,12 +433,14 @@ def test_router_dedup_cache_models():
     global server
 
     preset_path = os.path.join(TMP_DIR, "test_dedup.ini")
-    cache_id = "ggml-org/test-model-stories260K:F32"
+    main_cache_id = "ggml-org/test-model-stories260K:F32"
+    draft_cache_id = "ggml-org/test-model-stories260K-infill:F32"
 
     with open(preset_path, "w") as f:
         f.write(
             "[model-dedup]\n"
             "hf-repo = ggml-org/test-model-stories260K\n"
+            "spec-draft-hf = ggml-org/test-model-stories260K-infill\n"
             "dedup-cache-models = 1\n"
         )
 
@@ -448,12 +450,13 @@ def test_router_dedup_cache_models():
     try:
         ids = _get_model_ids(is_reload=False)
         assert "model-dedup" in ids
-        assert cache_id not in ids, "cache model should be hidden by dedup"
+        assert main_cache_id not in ids, "main cache model should be hidden by dedup"
+        assert draft_cache_id not in ids, "draft cache model should be hidden by dedup"
         # other cache models are unaffected
         assert "ggml-org/tinygemma3-GGUF:Q8_0" in ids
 
         # the hidden model is only hidden from the listing, it can still be used
-        res = server.make_request("POST", "/tokenize", data={"model": cache_id, "content": "hello"})
+        res = server.make_request("POST", "/tokenize", data={"model": main_cache_id, "content": "hello"})
         assert res.status_code == 200
 
         # disabling the flag brings the cache entry back on reload
@@ -461,9 +464,11 @@ def test_router_dedup_cache_models():
             f.write(
                 "[model-dedup]\n"
                 "hf-repo = ggml-org/test-model-stories260K\n"
+                "spec-draft-hf = ggml-org/test-model-stories260K-infill\n"
             )
         ids = _get_model_ids(is_reload=True)
-        assert cache_id in ids
+        assert main_cache_id in ids
+        assert draft_cache_id in ids
 
         # the flag also works from the global section
         with open(preset_path, "w") as f:
@@ -473,10 +478,12 @@ def test_router_dedup_cache_models():
                 "\n"
                 "[model-dedup]\n"
                 "hf-repo = ggml-org/test-model-stories260K\n"
+                "spec-draft-hf = ggml-org/test-model-stories260K-infill\n"
             )
         ids = _get_model_ids(is_reload=True)
         assert "model-dedup" in ids
-        assert cache_id not in ids, "cache model should be hidden by global dedup"
+        assert main_cache_id not in ids, "main cache model should be hidden by global dedup"
+        assert draft_cache_id not in ids, "draft cache model should be hidden by global dedup"
     finally:
         os.remove(preset_path)
 

tools/server/tests/unit/test_sleep.py

@@ -125,3 +125,43 @@ def test_server_sleep_metrics_buckets():
     assert res.status_code == 200
     assert is_sleeping(server) == False
     assert get_metric(fetch_metrics(server), "predicted_tokens_seconds") == 0
+
+
+def test_server_sleep_token_counting_wake():
+    global server
+    server.sleep_idle_seconds = 1
+    server.start()
+
+    wait_for_sleep(server)
+    assert is_sleeping(server)
+
+    res = server.make_request("POST", "/chat/completions/input_tokens", data={
+        "messages": [
+            {"role": "user", "content": "Hello world"}
+        ]
+    })
+    assert res.status_code == 200
+    assert res.body["input_tokens"] > 0
+    assert is_sleeping(server) == False
+
+    wait_for_sleep(server)
+    assert is_sleeping(server)
+
+    res = server.make_request("POST", "/v1/responses/input_tokens", data={
+        "input": "Hello world"
+    })
+    assert res.status_code == 200
+    assert res.body["input_tokens"] > 0
+    assert is_sleeping(server) == False
+
+    wait_for_sleep(server)
+    assert is_sleeping(server)
+
+    res = server.make_request("POST", "/v1/messages/count_tokens", data={
+        "messages": [
+            {"role": "user", "content": "Hello world"}
+        ]
+    })
+    assert res.status_code == 200
+    assert res.body["input_tokens"] > 0
+    assert is_sleeping(server) == False

tools/tuning/README.md

@@ -26,10 +26,13 @@ Sweep the grid (6 dtypes x 10 head sizes x 4 KV depths x 9 batch widths; a few h
 ./build/bin/ggml-metal-tuning fa-vec > fa_vec_rows.txt 2> fa_vec_sweep.log
 ```
 
-`fa_vec_rows.txt` holds nothing but table rows, ready to paste into `fa_vec_tuned_table`: the min-max-regret target, the aggregate benefit gate, the short-KV drop and the pointwise compression are already applied.
+`fa_vec_rows.txt` holds nothing but table rows: the min-max-regret target, the aggregate benefit gate, the short-KV drop and the pointwise compression are already applied.
+The rows carry the SKU token the runtime reported, but `fa_vec_tuned_table` is keyed by Apple GPU family, so that column has to be retagged before the rows compile.
+Your family number is on the `MTLGPUFamilyApple<N>` line the backend logs at init, near the top of `fa_vec_sweep.log`; `N` is the value, and the `MTLGPUFamilyCommon`/`MTLGPUFamilyMetal` lines beside it are not it.
+If your log is the only sweep for that family, its rows become the family's segment; where the family already has rows, post the log and let the two be compared before anything is replaced.
 A config represents a bucket only if it is no slower than the baseline config at every point that bucket covers, so a config that wins on average but loses at one batch width leaves its bucket at baseline.
 `fa_vec_sweep.log` holds the per-cell timings, bucket coverage, noise floor, any cooldown activity, and every config the no-harm rule refused together with the point that refused it.
-Post both: the log is what makes the rows reviewable.
+Post both, always: the rows now speak for every device in the family, so the log is what makes them reviewable.
 
 Long sweeps can be split.
 `--dtype f16,q4_0` and `--dk 128,192` restrict the grid, and the emitted rows for one `(dtype, head size)` do not depend on the others.