ui : model download pipeline

Track HuggingFace downloads end to end: the server download/cancel endpoints, a status manager fed by the /models/sse download progress events, and a models-discover store holding the catalog and detail state for the discover view. Downloaded and in-flight entries are excluded from the loadable model list. Assisted-by: pi:GLM-5.3-Flash
ui : model memory-fit estimation
2026-09-07 16:37:57 +02:00 · 2026-09-07 16:08:16 +02:00 · 2026-09-07 16:08:16 +02:00 · 2026-09-07 16:08:16 +02:00 · 2026-09-07 16:08:16 +02:00 · 2026-09-07 16:08:15 +02:00
72 changed files with 4108 additions and 491 deletions
@@ -31,7 +31,7 @@
    ]
    && blas.meta.available,
  useCuda ? config.cudaSupport,
-  useMetalKit ? stdenv.isAarch64 && stdenv.isDarwin,
+  useMetalKit ? stdenv.hostPlatform.isAarch64 && stdenv.hostPlatform.isDarwin,
  # Increases the runtime closure size by ~700M
  useMpi ? false,
  useRocm ? config.rocmSupport,
@@ -92,7 +92,7 @@ let

  cudaBuildInputs = with cudaPackages; [
    cuda_cudart
-    cuda_cccl # <nv/target>
+    cccl # <nv/target>
    libcublas
  ];

@@ -166,7 +166,7 @@ effectiveStdenv.mkDerivation (finalAttrs: {
  # `xcrun` is used find the path of the Metal compiler, which is varible
  # and not on $PATH
  # see https://github.com/ggml-org/llama.cpp/pull/6118 for discussion
-  __noChroot = effectiveStdenv.isDarwin && useMetalKit && precompileMetalShaders;
+  __noChroot = effectiveStdenv.hostPlatform.isDarwin && useMetalKit && precompileMetalShaders;

  nativeBuildInputs =
    [
@@ -181,10 +181,10 @@ effectiveStdenv.mkDerivation (finalAttrs: {
      autoAddDriverRunpath
    ]
    ++ optionals (effectiveStdenv.hostPlatform.isGnu && enableStatic) [ glibc.static ]
-    ++ optionals (effectiveStdenv.isDarwin && useMetalKit && precompileMetalShaders) [ xcrunHost ];
+    ++ optionals (effectiveStdenv.hostPlatform.isDarwin && useMetalKit && precompileMetalShaders) [ xcrunHost ];

  buildInputs =
-    optionals effectiveStdenv.isDarwin darwinBuildInputs
+    optionals effectiveStdenv.hostPlatform.isDarwin darwinBuildInputs
    ++ optionals useCuda cudaBuildInputs
    ++ optionals useMpi [ mpi ]
    ++ optionals useRocm rocmBuildInputs
@@ -245,7 +245,7 @@ effectiveStdenv.mkDerivation (finalAttrs: {

    # Configurations that are known to result in build failures. Can be
    # overridden by importing Nixpkgs with `allowBroken = true`.
-    broken = (useMetalKit && !effectiveStdenv.isDarwin);
+    broken = (useMetalKit && !effectiveStdenv.hostPlatform.isDarwin);

    description = "Inference of LLaMA model in pure C/C++${descriptionSuffix}";
    homepage = "https://github.com/ggml-org/llama.cpp/";
@@ -8,6 +8,7 @@
 #include "json.h"

 #include <algorithm>
+#include <cctype>
 #include <filesystem>
 #include <fstream>
 #include <future>
@@ -534,15 +535,18 @@ static gguf_split_info get_gguf_split_info(const std::string & path) {
 }

 // Q4_0 -> 4, F16 -> 16, NVFP4 -> 4, Q8_K_M -> 8, etc
-static int extract_quant_bits(const std::string & filename) {
-    auto split = get_gguf_split_info(filename);
+static int quant_bits_from_tag(const std::string & tag) {
+    auto pos = tag.find_first_of("0123456789");

-    auto pos = split.tag.find_first_of("0123456789");
    if (pos == std::string::npos) {
        return 0;
    }

-    return std::stoi(split.tag.substr(pos));
+    return std::stoi(tag.substr(pos));
+}
+
+static int extract_quant_bits(const std::string & filename) {
+    return quant_bits_from_tag(get_gguf_split_info(filename).tag);
 }

 static hf_cache::hf_files get_split_files(const hf_cache::hf_files & files,
@@ -563,12 +567,127 @@ static hf_cache::hf_files get_split_files(const hf_cache::hf_files & files,
    return result;
 }

-// pick the best sibling GGUF whose filename contains `keyword` (e.g. "mmproj" / "mtp"),
+// sidecar filename tokens, as used in `<quant>-<sidecar>` download tags,
+// e.g. `Q4_0-mtp` for `mtp-Model-Q4_0.gguf`, `BF16-mmproj` for `mmproj-BF16.gguf`
+static const std::vector<std::string> sidecar_tokens = {
+    "mtp", "eagle3", "dflash", "dspark", "mmproj", "imatrix",
+};
+
+static bool iequals(const std::string & a, const std::string & b) {
+    return a.size() == b.size() &&
+        std::equal(a.begin(), a.end(), b.begin(), [](char x, char y) {
+            return std::tolower((unsigned char) x) == std::tolower((unsigned char) y);
+        });
+}
+
+// split a `<quant>-<sidecar>` tag into its parts, e.g. `Q4_0-mtp` -> {`Q4_0`, `mtp`};
+// a bare sidecar tag (`mtp`) yields an empty quant; no sidecar yields an empty token
+static std::pair<std::string, std::string> split_sidecar_tag(const std::string & tag) {
+    for (const auto & t : sidecar_tokens) {
+        if (tag.size() > t.size() + 1 && iequals(tag.substr(tag.size() - t.size() - 1), "-" + t)) {
+            return { tag.substr(0, tag.size() - t.size() - 1), t };
+        }
+        if (iequals(tag, t)) {
+            return { "", t };
+        }
+    }
+    return { tag, "" };
+}
+
+// filename with directory and extension removed, e.g. `sub/mtp-Model-Q4_0.gguf` -> `mtp-Model-Q4_0`
+static std::string stem_of(const std::string & path) {
+    std::string base = path;
+    if (auto pos = base.rfind('/'); pos != std::string::npos) {
+        base = base.substr(pos + 1);
+    }
+    string_remove_suffix(base, ".gguf");
+    return base;
+}
+
+// a sidecar file carries its token as a name segment in any position and case,
+// e.g. `mmproj-Model-F16.gguf`, `Model-mtp-Q4_0.gguf`, `Model-Q4_0-mtp.gguf`
+// or the short `mmproj-F16.gguf`; the token must be a whole segment, so an
+// unrelated name that merely contains it (`smtp-Model.gguf`) is a plain model
+static std::string sidecar_token_of(const std::string & path) {
+    std::string base = stem_of(path);
+
+    for (char & c : base) {
+        c = (char) std::tolower((unsigned char) c);
+    }
+    if (base.empty()) {
+        return {};
+    }
+
+    for (const auto & t : sidecar_tokens) {
+        if (base == t) {
+            return t; // the sidecar file itself, e.g. `imatrix.gguf`
+        }
+        if (base.rfind(t + "-", 0) == 0) {
+            return t; // `mtp-Model-Q4_0.gguf`
+        }
+        if (base.find("-" + t + "-") != std::string::npos) {
+            return t; // `Model-mtp-Q4_0.gguf`
+        }
+        // `Model-Q4_0-mtp.gguf`, optionally with a `-draft` tail
+        if (string_ends_with(base, "-" + t) || string_ends_with(base, "-" + t + "-draft")) {
+            return t;
+        }
+    }
+    return {};
+}
+
+// name with the sidecar token segment removed, lowercased for tag parsing,
+// e.g. `Model-MTP-Q4_0` -> `model-q4_0`
+static std::string strip_sidecar_token(const std::string & base, const std::string & token) {
+    std::string lower = base;
+
+    for (char & c : lower) {
+        c = (char) std::tolower((unsigned char) c);
+    }
+
+    if (lower == token) {
+        return {};
+    }
+    if (lower.rfind(token + "-", 0) == 0) {
+        return lower.substr(token.size() + 1);
+    }
+    const std::string seg = "-" + token + "-";
+    if (auto pos = lower.find(seg); pos != std::string::npos) {
+        return lower.substr(0, pos) + "-" + lower.substr(pos + seg.size());
+    }
+    if (string_ends_with(lower, "-" + token + "-draft")) {
+        return lower.substr(0, lower.size() - token.size() - 7);
+    }
+    if (string_ends_with(lower, "-" + token)) {
+        return lower.substr(0, lower.size() - token.size() - 1);
+    }
+    return lower;
+}
+
+// the quant a sidecar file belongs to, from its name with the token stripped;
+// a short-form name (`mmproj-F16.gguf`) leaves the bare quant, which has no
+// `-` separator for the tag regex, so it becomes the tag directly
+static std::string sidecar_quant(const std::string & path, const std::string & token) {
+    std::string name = strip_sidecar_token(stem_of(path), token);
+    std::string tag  = get_gguf_split_info(name + ".gguf").tag;
+
+    if (tag.empty() && name.find('-') == std::string::npos) {
+        for (char & c : name) {
+            c = (char) std::toupper((unsigned char) c);
+        }
+        tag = name;
+    }
+
+    return tag;
+}
+
+// pick the best sibling GGUF carrying the sidecar `token` (e.g. "mmproj" / "mtp"),
 // preferring deeper shared directory prefix with the model, then exact `tag` match,
 // then closest quantization to the tag when given, or to the model otherwise
+// an empty `model` skips the directory constraint: the sidecar is matched by tag alone
 static hf_cache::hf_file find_best_sibling(const hf_cache::hf_files & files,
                                           const std::string        & model,
-                                           const std::string        & keyword,
+                                           const std::string        & token,
                                           const std::string        & tag = "") {
    hf_cache::hf_file best;
    size_t best_depth = 0;
@@ -589,32 +708,32 @@ static hf_cache::hf_file find_best_sibling(const hf_cache::hf_files & files,
        model_bits = extract_quant_bits(model);
    }
    auto model_parts = string_split<std::string>(model, '/');
-    auto model_dir = model_parts.end() - 1;

    for (const auto & f : files) {
-        if (!string_ends_with(f.path, ".gguf") ||
-            f.path.find(keyword) == std::string::npos) {
+        if (sidecar_token_of(f.path) != token) {
            continue;
        }

        auto sib_parts = string_split<std::string>(f.path, '/');
        auto sib_dir = sib_parts.end() - 1;

-        auto [_, dir] = std::mismatch(model_parts.begin(), model_dir,
-                                      sib_parts.begin(), sib_dir);
-        if (dir != sib_dir) {
-            continue;
+        size_t depth = 0;
+        if (!model.empty()) {
+            auto model_dir = model_parts.end() - 1;
+            auto [_, dir] = std::mismatch(model_parts.begin(), model_dir,
+                                          sib_parts.begin(), sib_dir);
+            if (dir != sib_dir) {
+                continue;
+            }
+            depth = dir - sib_parts.begin();
        }

-        size_t depth = dir - sib_parts.begin();
-        auto bits = extract_quant_bits(f.path);
-        auto diff = std::abs(bits - model_bits);
-
-        std::string path_upper = f.path;
-        for (char & c : path_upper) {
-            c = (char) std::toupper((unsigned char) c);
-        }
-        bool exact = !tag_upper.empty() && path_upper.find("-" + tag_upper + ".") != std::string::npos;
+        // rank by the quant the sidecar belongs to, with the token segment
+        // stripped from its name
+        auto tag   = sidecar_quant(f.path, token);
+        auto bits  = quant_bits_from_tag(tag);
+        auto diff  = std::abs(bits - model_bits);
+        bool exact = !tag_upper.empty() && tag == tag_upper;

        if (!found || depth > best_depth ||
            (depth == best_depth && exact && !best_exact) ||
@@ -637,43 +756,31 @@ static hf_cache::hf_file find_best_mmproj(const hf_cache::hf_files & files,
 static hf_cache::hf_file find_best_mtp(const hf_cache::hf_files & files,
                                       const std::string        & model,
                                       const std::string        & tag = "") {
-    return find_best_sibling(files, model, "mtp-", tag);
+    return find_best_sibling(files, model, "mtp", tag);
 }

 static hf_cache::hf_file find_best_eagle3(const hf_cache::hf_files & files,
                                          const std::string        & model,
                                          const std::string        & tag = "") {
-    return find_best_sibling(files, model, "eagle3-", tag);
+    return find_best_sibling(files, model, "eagle3", tag);
 }

 static hf_cache::hf_file find_best_dflash(const hf_cache::hf_files & files,
                                          const std::string        & model,
                                          const std::string        & tag = "") {
-    return find_best_sibling(files, model, "dflash-", tag);
+    return find_best_sibling(files, model, "dflash", tag);
 }

 static hf_cache::hf_file find_best_dspark(const hf_cache::hf_files & files,
                                          const std::string        & model,
                                          const std::string        & tag = "") {
-    return find_best_sibling(files, model, "dspark-", tag);
+    return find_best_sibling(files, model, "dspark", tag);
 }

+// a plain model file: a GGUF whose name carries no sidecar token segment,
+// so `smtp-Model.gguf` counts and every sidecar form does not
 static bool gguf_filename_is_model(const std::string & filepath) {
-    if (!string_ends_with(filepath, ".gguf")) {
-        return false;
-    }
-
-    std::string filename = filepath;
-    if (auto pos = filename.rfind('/'); pos != std::string::npos) {
-        filename = filename.substr(pos + 1);
-    }
-
-    return filename.find("mmproj")  == std::string::npos &&
-           filename.find("imatrix") == std::string::npos &&
-           filename.find("mtp-")    == std::string::npos &&
-           filename.find("eagle3-") == std::string::npos &&
-           filename.find("dflash-") == std::string::npos &&
-           filename.find("dspark-") == std::string::npos;
+    return string_ends_with(filepath, ".gguf") && sidecar_token_of(filepath).empty();
 }

 static hf_cache::hf_file find_best_model(const hf_cache::hf_files & files,
@@ -765,8 +872,27 @@ common_download_hf_plan common_download_get_hf_plan(const common_params_model &
        }
    } else {
        primary = find_best_model(all, tag);
+
+        // a `<quant>-<sidecar>` tag (e.g. `Q4_0-mtp`) requests that sidecar alone;
+        // every token-bearing file is a sidecar (find_best_model skips them),
+        // so the sidecar resolves here whenever the tag carries one
+        auto [base_tag, sidecar] = split_sidecar_tag(tag);
+
+        if (primary.path.empty() && !sidecar.empty()) {
+            auto found = find_best_sibling(all, "", sidecar, base_tag);
+
+            if (!found.path.empty()) {
+                if      (sidecar == "mtp")    plan.mtp    = found;
+                else if (sidecar == "eagle3") plan.eagle3 = found;
+                else if (sidecar == "dflash") plan.dflash = found;
+                else if (sidecar == "dspark") plan.dspark = found;
+                else                          plan.mmproj = found;
+            }
+        }
+
        // a requested sidecar can resolve on its own, without a full model of the same tag
-        if (primary.path.empty() && !opts.download_mtp && !opts.download_dflash && !opts.download_eagle3 && !opts.download_dspark) {
+        if (primary.path.empty() && sidecar.empty() &&
+            !opts.download_mtp && !opts.download_dflash && !opts.download_eagle3 && !opts.download_dspark) {
            LOG_ERR("%s: no GGUF files found in repository %s\n", __func__, repo.c_str());
            list_available_gguf_files(all);
            return plan;
@@ -794,7 +920,7 @@ common_download_hf_plan common_download_get_hf_plan(const common_params_model &
        plan.dspark = find_best_dspark(all, primary.path, tag);
    }

-    if (primary.path.empty() &&
+    if (primary.path.empty() && plan.mmproj.local_path.empty() &&
        plan.mtp.local_path.empty() && plan.dflash.local_path.empty() && plan.eagle3.local_path.empty() && plan.dspark.local_path.empty()) {
        LOG_ERR("%s: no GGUF files found in repository %s\n", __func__, repo.c_str());
        list_available_gguf_files(all);
@@ -968,17 +1094,34 @@ std::vector<common_cached_model_info> common_list_cached_models() {
    auto files = hf_cache::get_cached_files();

    for (const auto & f : files) {
-        auto split = get_gguf_split_info(f.path);
-        if (split.index != 1 || split.tag.empty() ||
-            split.prefix.find("mmproj")  != std::string::npos ||
-            split.prefix.find("mtp-")    != std::string::npos ||
-            split.prefix.find("eagle3-") != std::string::npos ||
-            split.prefix.find("dflash-") != std::string::npos ||
-            split.prefix.find("dspark-") != std::string::npos) {
-            continue;
+        // a sidecar file is listed under its own `<quant>-<sidecar>` tag, so a
+        // cached `mtp-Model-Q4_0.gguf`, `Model-mtp-Q4_0.gguf`, `Model-Q4_0-mtp.gguf`
+        // or short `mmproj-F16.gguf` shows up as `<repo>:Q4_0-mtp` / `<repo>:F16-mmproj`;
+        // files whose name carries no token stay loadable models
+        auto token = sidecar_token_of(f.path);
+
+        std::string tag;
+        if (token.empty()) {
+            auto split = get_gguf_split_info(f.path);
+
+            if (split.index != 1 || split.tag.empty()) {
+                continue;
+            }
+
+            tag = split.tag;
+        } else {
+            tag = sidecar_quant(f.path, token);
+
+            // a bare sidecar file has no tag to request it by, stay hidden
+            if (tag.empty()) {
+                continue;
+            }
+
+            tag += "-" + token;
        }
-        if (seen.insert(f.repo_id + ":" + split.tag).second) {
-            result.push_back({f.repo_id, split.tag});
+
+        if (seen.insert(f.repo_id + ":" + tag).second) {
+            result.push_back({f.repo_id, tag});
        }
    }

@@ -1014,7 +1157,15 @@ bool common_download_remove(const std::string & hf_repo_with_tag) {
        return hf_cache::remove_cached_repo(repo_id);
    }

-    std::string tag_upper = tag;
+    // a `<quant>-<sidecar>` tag (`Q4_0-mtp`) targets that sidecar alone; a bare
+    // sidecar tag (`mtp`) is ambiguous across quants and is rejected
+    auto [base_tag, sidecar] = split_sidecar_tag(tag);
+    if (!sidecar.empty() && base_tag.empty()) {
+        LOG_ERR("%s: bare sidecar tag '%s': use `<quant>-<sidecar>`\n", __func__, tag.c_str());
+        return false;
+    }
+
+    std::string tag_upper = sidecar.empty() ? tag : base_tag;
    for (char & c : tag_upper) {
        c = (char) std::toupper((unsigned char) c);
    }
@@ -1024,13 +1175,25 @@ bool common_download_remove(const std::string & hf_repo_with_tag) {
        return false;
    }

-    // collect snapshot entries whose tag matches
+    // collect the snapshot entries the tag selects; sidecar files keep their
+    // own tags, so a plain quant tag never removes them
    std::vector<fs::path> to_remove;
    for (const auto & f : files) {
-        auto split = get_gguf_split_info(f.path);
-        if (split.tag == tag_upper) {
-            to_remove.emplace_back(f.local_path);
+        auto token = sidecar_token_of(f.path);
+
+        if (sidecar.empty()) {
+            if (!token.empty()) {
+                continue;
+            }
+            if (get_gguf_split_info(f.path).tag != tag_upper) {
+                continue;
+            }
+        } else {
+            if (token != sidecar || sidecar_quant(f.path, sidecar) != tag_upper) {
+                continue;
+            }
        }
+        to_remove.emplace_back(f.local_path);
    }

    if (to_remove.empty()) {
@@ -9,20 +9,6 @@ from .base import ModelBase, gguf, logger
 from .deepseek import DeepseekV2Model


-def split_kv_b_proj(weight: torch.Tensor, n_head: int, qk_nope: int, v_head_dim: int):
-    """Split kv_b_proj into k_b (transposed) and v_b, matching DeepSeek MLA absorption.
-
-    weight: [n_head*(qk_nope+v_head_dim), kv_lora_rank].
-    Returns (k_b, v_b): k_b [n_head, kv_lora_rank, qk_nope], v_b [n_head, v_head_dim, kv_lora_rank].
-    """
-    kv_lora = weight.shape[-1]
-    assert weight.shape[0] == n_head * (qk_nope + v_head_dim)
-    kv_b = weight.view(n_head, qk_nope + v_head_dim, kv_lora)
-    k_b, v_b = torch.split(kv_b, [qk_nope, v_head_dim], dim=1)
-    k_b = k_b.transpose(1, 2).contiguous()  # [n_head, kv_lora, qk_nope]
-    return k_b, v_b.contiguous()
-
-
 def split_gate_up(weight: torch.Tensor, moe_intermediate_size: int):
    """Split a fused stacked gate_up expert tensor into (gate, up).

@@ -36,6 +22,7 @@ def split_gate_up(weight: torch.Tensor, moe_intermediate_size: int):


@ModelBase.register("HYV4ForCausalLM")
+@ModelBase.example("tencent/Hy4-preview")
 class HYV4Model(DeepseekV2Model):
    """HY_V4: DeepSeek-V3 style MLA + MoE with iHC, a gated MLA output and a learnable sink.

@@ -54,6 +41,8 @@ class HYV4Model(DeepseekV2Model):

    model_arch = gguf.MODEL_ARCH.HY_V4

+    merge_expert = False
+
    # tensors a "full" indexer layer must carry
    INDEXER_SUFFIXES = frozenset({
        "self_attn.indexer.wq_b.weight",
@@ -186,6 +175,10 @@ class HYV4Model(DeepseekV2Model):
            )

    def prepare_tensors(self):
+        # Hy4-preview for some reason has num_key_value_heads equal to 8, so override it here
+        # without this conversion/deepseek.py fails on assert
+        self.hparams["num_key_value_heads"] = self.hparams["num_attention_heads"]
+
        # validate before the base materializes tensors, so a mismatch fails early
        is_full = self.indexer_is_full()
        if is_full is not None:
@@ -227,85 +220,25 @@ class HYV4Model(DeepseekV2Model):

    def modify_tensors(self, data_torch: torch.Tensor, name: str, bid: int | None) -> Iterable[tuple[str, torch.Tensor]]:
        hparams = self.hparams
-        n_head = hparams["num_attention_heads"]
-        qk_nope = hparams["qk_nope_head_dim"]
-        v_head_dim = hparams["v_head_dim"]
        moe_inter = hparams["moe_intermediate_size"]

        tn = self.format_tensor_name

-        # ---- global (non per-layer) ----
-        if name == "model.embed_tokens.weight":
-            return [(tn(gguf.MODEL_TENSOR.TOKEN_EMBD), data_torch)]
-        if name == "model.norm.weight":
-            return [(tn(gguf.MODEL_TENSOR.OUTPUT_NORM), data_torch)]
-        if name == "lm_head.weight":
-            return [(tn(gguf.MODEL_TENSOR.OUTPUT), data_torch)]
-        if name == "model.hc_head.hc_head_fn":
-            return [(tn(gguf.MODEL_TENSOR.HC_HEAD_FN), data_torch)]
-        if name == "model.hc_head.hc_head_base":
-            return [(tn(gguf.MODEL_TENSOR.HC_HEAD_BASE), data_torch)]
-        if name == "model.hc_head.hc_head_scale":
-            return [(tn(gguf.MODEL_TENSOR.HC_HEAD_SCALE), data_torch)]
-
-        assert bid is not None, f"expected a per-layer tensor, got {name!r}"
-
-        # ---- per-layer, keyed by suffix after 'model.layers.{bid}.' ----
-        suffix = name.split(f"model.layers.{bid}.", 1)[-1]
-
-        # note: q_b_proj and kv_a_proj_with_mqa are mapped straight through (no RoPE permute),
-        # the graph rotates consecutive pairs so the rows need no reordering
-        simple = {
-            "input_layernorm.weight":          (gguf.MODEL_TENSOR.ATTN_NORM, ".weight"),
-            "post_attention_layernorm.weight": (gguf.MODEL_TENSOR.FFN_NORM,  ".weight"),
-            "self_attn.q_a_proj.weight":       (gguf.MODEL_TENSOR.ATTN_Q_A, ".weight"),
-            "self_attn.q_a_layernorm.weight":  (gguf.MODEL_TENSOR.ATTN_Q_A_NORM, ".weight"),
-            "self_attn.q_b_proj.weight":       (gguf.MODEL_TENSOR.ATTN_Q_B, ".weight"),
-            "self_attn.kv_a_proj_with_mqa.weight": (gguf.MODEL_TENSOR.ATTN_KV_A_MQA, ".weight"),
-            "self_attn.kv_a_layernorm.weight": (gguf.MODEL_TENSOR.ATTN_KV_A_NORM, ".weight"),
-            "self_attn.o_proj.weight":         (gguf.MODEL_TENSOR.ATTN_OUT, ".weight"),
-            "self_attn.linear_gate.weight":    (gguf.MODEL_TENSOR.ATTN_GATE, ".weight"),
-            "self_attn.learnable_sink_param":  (gguf.MODEL_TENSOR.ATTN_SINKS, ".weight"),
-            "self_attn.indexer.wq_b.weight":   (gguf.MODEL_TENSOR.INDEXER_ATTN_Q_B, ".weight"),
-            "self_attn.indexer.wk.weight":     (gguf.MODEL_TENSOR.INDEXER_ATTN_K, ".weight"),
-            "self_attn.indexer.k_norm.weight": (gguf.MODEL_TENSOR.INDEXER_K_NORM, ".weight"),
-            "self_attn.indexer.k_norm.bias":   (gguf.MODEL_TENSOR.INDEXER_K_NORM, ".bias"),
-            "self_attn.indexer.weights_proj.weight": (gguf.MODEL_TENSOR.INDEXER_PROJ, ".weight"),
-            "hc_attn_layer.hc_pre.hc_fn":      (gguf.MODEL_TENSOR.HC_ATTN_FN, ".weight"),
-            "hc_attn_layer.hc_pre.hc_base":    (gguf.MODEL_TENSOR.HC_ATTN_BASE, ".weight"),
-            "hc_attn_layer.hc_pre.hc_scale":   (gguf.MODEL_TENSOR.HC_ATTN_SCALE, ".weight"),
-            "hc_mlp_layer.hc_pre.hc_fn":       (gguf.MODEL_TENSOR.HC_FFN_FN, ".weight"),
-            "hc_mlp_layer.hc_pre.hc_base":     (gguf.MODEL_TENSOR.HC_FFN_BASE, ".weight"),
-            "hc_mlp_layer.hc_pre.hc_scale":    (gguf.MODEL_TENSOR.HC_FFN_SCALE, ".weight"),
-            "mlp.gate.weight":                 (gguf.MODEL_TENSOR.FFN_GATE_INP, ".weight"),
-            "mlp.gate.e_score_correction.bias":(gguf.MODEL_TENSOR.FFN_EXP_PROBS_B, ".bias"),
-            "mlp.gate_proj.weight":            (gguf.MODEL_TENSOR.FFN_GATE, ".weight"),
-            "mlp.up_proj.weight":              (gguf.MODEL_TENSOR.FFN_UP, ".weight"),
-            "mlp.down_proj.weight":            (gguf.MODEL_TENSOR.FFN_DOWN, ".weight"),
-            "mlp.shared_experts.gate_proj.weight": (gguf.MODEL_TENSOR.FFN_GATE_SHEXP, ".weight"),
-            "mlp.shared_experts.up_proj.weight":   (gguf.MODEL_TENSOR.FFN_UP_SHEXP, ".weight"),
-            "mlp.shared_experts.down_proj.weight": (gguf.MODEL_TENSOR.FFN_DOWN_SHEXP, ".weight"),
-        }
-        if suffix in simple:
-            key, sfx = simple[suffix]
-            return [(tn(key, bid, sfx), data_torch)]
-
-        # kv_b_proj: split into k_b (transposed) and v_b
-        if suffix == "self_attn.kv_b_proj.weight":
-            k_b, v_b = split_kv_b_proj(data_torch, n_head, qk_nope, v_head_dim)
-            return [
-                (tn(gguf.MODEL_TENSOR.ATTN_K_B, bid), k_b),
-                (tn(gguf.MODEL_TENSOR.ATTN_V_B, bid), v_b),
-            ]
-
        # fused stacked experts: split gate_up into gate/up
-        if suffix == "mlp.experts.gate_up_proj":
+        if name.endswith("mlp.experts.gate_up_proj"):
            gate, up = split_gate_up(data_torch, moe_inter)
-            return [
-                (tn(gguf.MODEL_TENSOR.FFN_GATE_EXP, bid), gate),
-                (tn(gguf.MODEL_TENSOR.FFN_UP_EXP, bid), up),
-            ]
-        if suffix == "mlp.experts.down_proj":
-            return [(tn(gguf.MODEL_TENSOR.FFN_DOWN_EXP, bid), data_torch)]
+            yield from super().modify_tensors(gate, tn(gguf.MODEL_TENSOR.FFN_GATE_EXP, bid), bid)
+            yield from super().modify_tensors(up,   tn(gguf.MODEL_TENSOR.FFN_UP_EXP,   bid), bid)
+            return

-        raise ValueError(f"Unsupported HY_V4 tensor {name!r} (suffix {suffix!r})")
+        # add .weight suffixes
+        if name.endswith("mlp.experts.down_proj") or name.endswith(".self_attn.learnable_sink_param"):
+            name += ".weight"
+
+        if re.search(r"\.hc_head\.hc_head_(?:fn|base|scale)$", name):
+            name += ".weight"
+
+        if re.search(r"\.hc_(?:attn|mlp)_layer\.hc_pre\.hc_(?:fn|base|scale)$", name):
+            name += ".weight"
+
+        yield from super().modify_tensors(data_torch, name, bid)
@@ -128,7 +128,7 @@
          }:
          {
            # For standardised reproducible formatting with `nix fmt`
-            formatter = pkgs.nixfmt-rfc-style;
+            formatter = pkgs.nixfmt;

            # Unlike `.#packages`, legacyPackages may contain values of
            # arbitrary types (including nested attrsets) and may even throw
@@ -156,7 +156,7 @@
                windows = config.legacyPackages.llamaPackagesWindows.llama-cpp;
                python-scripts = config.legacyPackages.llamaPackages.python-scripts;
              }
-              // lib.optionalAttrs pkgs.stdenv.isLinux {
+              // lib.optionalAttrs pkgs.stdenv.hostPlatform.isLinux {
                cuda = config.legacyPackages.llamaPackagesCuda.llama-cpp;

                mpi-cpu = config.packages.default.override { useMpi = true; };
@@ -69,6 +69,8 @@
 #define GGML_CUDA_CC_GCN4       (GGML_CUDA_CC_OFFSET_AMD + 0x803)  // Tonga, Fiji, Polaris, minimum for fast fp16
 #define GGML_CUDA_CC_VEGA       (GGML_CUDA_CC_OFFSET_AMD + 0x900)  // Vega56/64, minimum for fp16 dual issue
 #define GGML_CUDA_CC_VEGA20     (GGML_CUDA_CC_OFFSET_AMD + 0x906)  // MI50/Radeon VII, minimum for dp4a
+#define GGML_CUDA_CC_GFX909     (GGML_CUDA_CC_OFFSET_AMD + 0x909)  // GCN APU
+#define GGML_CUDA_CC_GFX90C     (GGML_CUDA_CC_OFFSET_AMD + 0x90c)  // GCN APU
 #define GGML_CUDA_CC_CDNA1      (GGML_CUDA_CC_OFFSET_AMD + 0x908)  // MI100, minimum for MFMA, acc registers
 #define GGML_CUDA_CC_CDNA2      (GGML_CUDA_CC_OFFSET_AMD + 0x90a)  // MI210 (gfx90a), minimum acc register renaming
 #define GGML_CUDA_CC_CDNA3      (GGML_CUDA_CC_OFFSET_AMD + 0x942)  // MI300
@@ -89,12 +91,13 @@
 #define GGML_CUDA_CC_IS_RDNA3_5(cc) (cc >= GGML_CUDA_CC_RDNA3_5 && cc < GGML_CUDA_CC_RDNA4)
 #define GGML_CUDA_CC_IS_RDNA3(cc)   (GGML_CUDA_CC_IS_RDNA3_0(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc))
 #define GGML_CUDA_CC_IS_RDNA4(cc)   (cc >= GGML_CUDA_CC_RDNA4)
-#define GGML_CUDA_CC_IS_GCN(cc)     (cc > GGML_CUDA_CC_OFFSET_AMD && cc < GGML_CUDA_CC_CDNA1)
-#define GGML_CUDA_CC_IS_CDNA(cc)    (cc >= GGML_CUDA_CC_CDNA1 && cc < GGML_CUDA_CC_RDNA1)
-#define GGML_CUDA_CC_IS_CDNA1(cc)   (cc >= GGML_CUDA_CC_CDNA1 && cc < GGML_CUDA_CC_CDNA2)
-#define GGML_CUDA_CC_IS_CDNA2(cc)   (cc >= GGML_CUDA_CC_CDNA2 && cc < GGML_CUDA_CC_CDNA3)
-#define GGML_CUDA_CC_IS_CDNA3(cc)   (cc >= GGML_CUDA_CC_CDNA3 && cc < GGML_CUDA_CC_CDNA4)
-#define GGML_CUDA_CC_IS_CDNA4(cc)   (cc >= GGML_CUDA_CC_CDNA4 && cc < GGML_CUDA_CC_RDNA1)
+#define GGML_CUDA_CC_IS_GCN_APU(cc) ((cc) == GGML_CUDA_CC_GFX909 || (cc) == GGML_CUDA_CC_GFX90C)
+#define GGML_CUDA_CC_IS_GCN(cc)     ((cc > GGML_CUDA_CC_OFFSET_AMD && cc < GGML_CUDA_CC_CDNA1) || GGML_CUDA_CC_IS_GCN_APU(cc))
+#define GGML_CUDA_CC_IS_CDNA(cc)    (!GGML_CUDA_CC_IS_GCN_APU(cc) && cc >= GGML_CUDA_CC_CDNA1 && cc < GGML_CUDA_CC_RDNA1)
+#define GGML_CUDA_CC_IS_CDNA1(cc)   (GGML_CUDA_CC_IS_CDNA(cc) && cc >= GGML_CUDA_CC_CDNA1 && cc < GGML_CUDA_CC_CDNA2)
+#define GGML_CUDA_CC_IS_CDNA2(cc)   (GGML_CUDA_CC_IS_CDNA(cc) && cc >= GGML_CUDA_CC_CDNA2 && cc < GGML_CUDA_CC_CDNA3)
+#define GGML_CUDA_CC_IS_CDNA3(cc)   (GGML_CUDA_CC_IS_CDNA(cc) && cc >= GGML_CUDA_CC_CDNA3 && cc < GGML_CUDA_CC_CDNA4)
+#define GGML_CUDA_CC_IS_CDNA4(cc)   (GGML_CUDA_CC_IS_CDNA(cc) && cc >= GGML_CUDA_CC_CDNA4 && cc < GGML_CUDA_CC_RDNA1)

 // Moore Threads
 #define MUSART_HMASK 40300 // MUSA rc4.3, min. ver. for half2 -> uint mask comparisons
@@ -212,6 +212,7 @@ static int ggml_cuda_parse_id(char devName[]) {
    }
    archNum += archMajor * 0x100;
    archNum += archMinor;
+
    return archNum;
 }
 #endif // defined(GGML_USE_HIP)
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_ampere(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = false;
    CASE(GGML_TYPE_Q1_0, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
    CASE(GGML_TYPE_Q1_0, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
    CASE(GGML_TYPE_Q1_0, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
@@ -379,5 +380,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false);
    CASE(GGML_TYPE_NVFP4, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_blackwell(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = false;
    CASE(GGML_TYPE_MXFP4, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);
    CASE(GGML_TYPE_MXFP4, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);
    CASE(GGML_TYPE_MXFP4, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_FP4, MMQ_ITER_K_FP4, true, true);
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_cdna(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = false;
    CASE(GGML_TYPE_Q1_0, 512, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
    CASE(GGML_TYPE_Q1_0, 512, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
    CASE(GGML_TYPE_Q1_0, 512, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
@@ -181,5 +182,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 512, 1, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false);
    CASE(GGML_TYPE_NVFP4, 512, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, true, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 512, 1, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_pascal_dp4a(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = false;
    CASE(GGML_TYPE_Q1_0, 256, 2, 64,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 256, 2, 64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 256, 2, 64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
@@ -269,5 +270,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 256, 2, 64,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
    CASE(GGML_TYPE_NVFP4, 256, 2, 64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_pascal_older(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = false;
    CASE(GGML_TYPE_Q1_0, 256, 2, 64,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 256, 2, 64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 256, 2, 64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
@@ -269,5 +270,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 256, 2, 64,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
    CASE(GGML_TYPE_NVFP4, 256, 2, 64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna2(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = false;
    CASE(GGML_TYPE_Q1_0, 256, 2, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 256, 2, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 256, 2, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
@@ -269,5 +270,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 256, 2, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
    CASE(GGML_TYPE_NVFP4, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna3_5(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = false;
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
@@ -286,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
    CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna3(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = true;
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
@@ -270,5 +271,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 256, 2, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
    CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -1,4 +1,5 @@
 static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_rdna4(ggml_type type, int J, bool fallback) {
+    constexpr bool use_typical_moe_ncols = true;
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
    CASE(GGML_TYPE_Q1_0, 128, 2,  64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
@@ -286,5 +287,5 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
    CASE(GGML_TYPE_NVFP4, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
    CASE(GGML_TYPE_NVFP4, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);

-    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
+    return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, use_typical_moe_ncols, false, true);
 }
@@ -375,10 +375,10 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
        return true;
    }

-    // gfx900 (Vega 10) lacks native dp4a, loses to dequant + hipBLAS
+    // gfx900 (Vega 10), gfx909, and gfx90c lack native dp4a, losing to dequant + hipBLAS
    // for dense matrices; keep MMQ only for MoE, where the
    // hipBLAS path is much slower.
-    if (cc == GGML_CUDA_CC_VEGA) {
+    if (cc == GGML_CUDA_CC_VEGA || GGML_CUDA_CC_IS_GCN_APU(cc)) {
        return n_experts > 0;
    }

@@ -170,12 +170,13 @@ struct ggml_cuda_mmq_config {
    int                       J;           // SRAM tile width in src1->ne[1]/dst->ne[1] direction.
    ggml_cuda_mmq_sram_layout sram_layout; // SRAM tile length in src0->ne[0]/src1->ne[0] direction (physical 32 bit elements).
    int                       K_vram;      // VRAM tile length in src0->ne[0]/src1->ne[0] direction (logical elements).
+    bool                      use_typical_moe_ncols;
    bool                      stream_k;    // Whether or not to use stream-k decomposition.
    bool                      fallback;    // Whether a fallback for out-of-bounds check in src0->ne[1] direction is needed.

    constexpr __host__ __device__ ggml_cuda_mmq_config(
-            ggml_type type, int nthreads, int occupancy, int I, int J, ggml_cuda_mmq_sram_layout sram_layout, int K_vram, bool stream_k, bool fallback) :
-        type(type), nthreads(nthreads), occupancy(occupancy), I(I), J(J), sram_layout(sram_layout), K_vram(K_vram), stream_k(stream_k), fallback(fallback) {}
+            ggml_type type, int nthreads, int occupancy, int I, int J, ggml_cuda_mmq_sram_layout sram_layout, int K_vram, bool use_typical_moe_ncols, bool stream_k, bool fallback) :
+        type(type), nthreads(nthreads), occupancy(occupancy), I(I), J(J), sram_layout(sram_layout), K_vram(K_vram), use_typical_moe_ncols(use_typical_moe_ncols), stream_k(stream_k), fallback(fallback) {}

    constexpr __device__ int rows_per_warp() const {
 #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
@@ -210,7 +211,7 @@ struct ggml_cuda_mmq_config {
        static_assert((I_)        %  32 == 0,                             "bad I");                                                       \
        static_assert((J_)        %   8 == 0,                             "bad J");                                                       \
        static_assert((K_vram_)   % 256 == 0,                             "bad K_vram");                                                  \
-        return ggml_cuda_mmq_config((type_), (nthreads_), (occupancy_), (I_), (J_), (sram_layout_), (K_vram_), (stream_k_), (fallback_)); \
+        return ggml_cuda_mmq_config((type_), (nthreads_), (occupancy_), (I_), (J_), (sram_layout_), (K_vram_), use_typical_moe_ncols, (stream_k_), (fallback_)); \
    }                                                                                                                                     \

 #include "mmq-config-pascal-older.cuh"
@@ -1473,6 +1474,20 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args,
    const int    cc    = ggml_cuda_info().devices[id].cc;
    const size_t smpbo = ggml_cuda_info().devices[id].smpbo;

+    int64_t ncols_picker = args.ncols_max;
+    if (args.expert_bounds != nullptr && args.nchannels_x > 0) {
+        const int J_max = ggml_cuda_mmq_get_J_max(type, fallback, cc, 128);
+        const ggml_cuda_mmq_config config_max = ggml_cuda_mmq_get_config(type, J_max, fallback, cc);
+        if (config_max.use_typical_moe_ncols) {
+            // Use the typical expert width only for tile selection.
+            // The launch grid still uses args.ncols_max.
+            const int64_t ncols_typical = (args.ncols_dst + args.nchannels_x - 1) / args.nchannels_x;
+            if (ncols_typical >= 1 && ncols_typical < J_max && ncols_typical < ncols_picker) {
+                ncols_picker = ncols_typical;
+            }
+        }
+    }
+
    int J_best        = 0;
    int ntiles_J_best = INT_MAX;

@@ -1486,7 +1501,7 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args,
            continue;
        }

-        const int ntiles_x = (args.ncols_max + config.J - 1) / config.J;
+        const int ntiles_x = (ncols_picker + config.J - 1) / config.J;

        if (ntiles_x < ntiles_J_best) {
            J_best = J;
@@ -176,9 +176,9 @@

 #define __CUDA_ARCH__ 1300

-#if defined(__gfx900__) || defined(__gfx906__)
+#if defined(__gfx900__) || defined(__gfx906__) || defined(__gfx909__) || defined(__gfx90c__)
 #define GCN5
-#endif // defined(__gfx900__) || defined(__gfx906__)
+#endif // defined(__gfx900__) || defined(__gfx906__) || defined(__gfx909__) || defined(__gfx90c__)

 #if defined(__gfx803__)
 #define GCN4
@@ -4858,6 +4858,78 @@ static bool ggml_sycl_mul_mat_glu_mmvq_fused(ggml_backend_sycl_context & ctx, gg
                                               /*stride_col_dst=*/(int) glu->ne[0], stream);
 }

+// Batch the run of consecutive L2_NORM siblings starting at node_idx into one launch.
+// Returns the number of extra graph nodes consumed, or 0 if the run is shorter than two
+// (the caller then runs the norm through the per-tensor kernel).
+static int ggml_sycl_l2_norm_batch_fused(ggml_backend_sycl_context & ctx, ggml_cgraph * cgraph, int node_idx) {
+    const ggml_tensor * node = cgraph->nodes[node_idx];
+    if (ggml_sycl_info().device_count != 1 || node->type != GGML_TYPE_F32 ||
+        node->src[0]->type != GGML_TYPE_F32 || node->src[0]->ne[0] >= 1024) {
+        return 0;
+    }
+
+    ggml_tensor * batch[GGML_SYCL_L2_BATCH_MAX];
+    int           count = 0;
+    int           last  = node_idx;
+    float         eps0;
+    memcpy(&eps0, node->op_params, sizeof(float));
+
+    // Conservative aliasing test: the batched norms run concurrently in one kernel,
+    // so none may read what another writes, and none may write where another writes.
+    auto overlaps = [](const ggml_tensor * a, const ggml_tensor * b) {
+        const char * ab = (const char *) a->data;
+        const char * bb = (const char *) b->data;
+        return ab < bb + ggml_nbytes(b) && bb < ab + ggml_nbytes(a);
+    };
+
+    for (int j = node_idx; j < cgraph->n_nodes && count < GGML_SYCL_L2_BATCH_MAX; ++j) {
+        ggml_tensor * nj = cgraph->nodes[j];
+        if (ggml_is_empty(nj) || nj->op == GGML_OP_RESHAPE || nj->op == GGML_OP_TRANSPOSE ||
+            nj->op == GGML_OP_VIEW || nj->op == GGML_OP_PERMUTE || nj->op == GGML_OP_NONE ||
+            (nj->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) {
+            continue;  // not a launch; cannot break a run of adjacent norms
+        }
+        if (nj->op != GGML_OP_L2_NORM || nj->type != GGML_TYPE_F32 ||
+            nj->src[0]->type != GGML_TYPE_F32 || !ggml_are_same_shape(nj, node) ||
+            !ggml_are_same_shape(nj->src[0], node->src[0])) {
+            break;  // any other launch ends the run
+        }
+        bool same_nb = true;
+        for (int d = 0; d < GGML_MAX_DIMS; ++d) {
+            if (nj->nb[d] != node->nb[d] || nj->src[0]->nb[d] != node->src[0]->nb[d]) {
+                same_nb = false;
+                break;
+            }
+        }
+        if (!same_nb) {
+            break;  // one nb[] stride set is shared by the whole batch
+        }
+        float epsj;
+        memcpy(&epsj, nj->op_params, sizeof(float));
+        if (epsj != eps0) {
+            break;  // eps mismatch ends the run
+        }
+        bool indep = true;
+        for (int k = 0; k < count; ++k) {
+            if (overlaps(nj->src[0], batch[k]) || overlaps(nj, batch[k])) {
+                indep = false;
+                break;
+            }
+        }
+        if (!indep) {
+            break;  // an overlapping tensor would race inside one launch
+        }
+        batch[count++] = nj;
+        last           = j;
+    }
+    if (count < 2) {
+        return 0;  // a lone norm falls through to the per-tensor kernel
+    }
+    ggml_sycl_l2_norm_batch(ctx, batch, count);
+    return last - node_idx;
+}
+
+
 __dpct_inline__ static void k_copy_src1_to_contiguous(
    const char *__restrict__ src1_original, char *__restrict__ src1_contiguous,
    const mmid_row_mapping *__restrict__ row_mapping,
@@ -5908,6 +5980,17 @@ static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * syc
            continue;
        }

+        // Batch consecutive independent same-shape F32 L2_NORM siblings (the GDN q/k
+        // norms) into one launch; sources are strided views of the fused qkv buffer, so
+        // the scan skips the interleaved view nodes instead of breaking on them.
+        if (node->op == GGML_OP_L2_NORM) {
+            const int l2_batch_skip = ggml_sycl_l2_norm_batch_fused(*sycl_ctx, cgraph, i);
+            if (l2_batch_skip > 0) {
+                i += l2_batch_skip;
+                continue;
+            }
+        }
+
        if (node->op == GGML_OP_MUL_MAT && ggml_sycl_mul_mat_glu_mmvq_fused(*sycl_ctx, cgraph, i)) {
            i += 2;
            continue;
@@ -543,6 +543,62 @@ static void l2_norm_f32_sycl(const float *   x,
    }
 }

+// Batched L2 norm: N independent same-shape F32 tensors in one launch; the tensor
+// index is folded into grid dim0 and each row's reduction is identical to the
+// single-tensor kernel, so the result is bit-exact.
+struct l2_batch_ptrs {
+    const float * src[GGML_SYCL_L2_BATCH_MAX];
+    float *       dst[GGML_SYCL_L2_BATCH_MAX];
+};
+
+// One stride set shared by the whole batch: the caller only groups tensors whose nb[]
+// all match, so per-tensor state stays two pointers.
+struct l2_batch_strides {
+    int     ne1, ne2;
+    int64_t ss0, ss1, ss2, ss3;
+    int64_t ds0, ds1, ds2, ds3;
+};
+
+template <int warp_size>
+static void l2_norm_f32_batch(l2_batch_ptrs p, l2_batch_strides st, const int ncols, const float eps,
+                              const sycl::nd_item<3> & item_ct1) {
+    const int t   = item_ct1.get_group(0);  // tensor index
+    const int r   = item_ct1.get_group(2);  // flattened row over ne1*ne2*ne3
+    const int tid = item_ct1.get_local_id(2);
+
+    const int i1 = r % st.ne1;
+    const int i2 = (r / st.ne1) % st.ne2;
+    const int i3 = r / (st.ne1 * st.ne2);
+
+    const float * x   = p.src[t] + i3 * st.ss3 + i2 * st.ss2 + i1 * st.ss1;
+    float *       dst = p.dst[t] + i3 * st.ds3 + i2 * st.ds2 + i1 * st.ds1;
+
+    float tmp = 0.0f;
+    for (int col = tid; col < ncols; col += warp_size) {
+        const float xi = x[col * st.ss0];
+        tmp += xi * xi;
+    }
+    tmp = block_reduce<block_reduce_method::SUM, warp_size>(tmp, (float *) nullptr, warp_size);
+    const float scale = sycl::rsqrt(sycl::fmax(tmp, eps * eps));
+    for (int col = tid; col < ncols; col += warp_size) {
+        dst[col * st.ds0] = scale * x[col * st.ss0];
+    }
+}
+
+template <int warp_size>
+static void l2_norm_f32_batch_sycl(l2_batch_ptrs p, l2_batch_strides st, const int n_tensors,
+                                   const int ncols, const int nrows_total, const float eps,
+                                   queue_ptr stream) {
+    const dpct::dim3 blocks_num(nrows_total, 1, n_tensors);
+    const dpct::dim3 block_dims(warp_size, 1, 1);
+    stream->submit([&](sycl::handler & cgh) {
+        cgh.parallel_for(sycl::nd_range<3>(blocks_num * block_dims, block_dims),
+            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(warp_size)]] {
+                l2_norm_f32_batch<warp_size>(p, st, ncols, eps, item_ct1);
+            });
+    });
+}
+
 void ggml_sycl_op_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
    const ggml_tensor * src0 = dst->src[0];

@@ -961,3 +1017,30 @@ void ggml_sycl_op_l2_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst) {
    l2_norm_f32_sycl<WARP_SIZE>(src0_d, dst_d, ne00, ne01, ne02, ne03,
            ss0, ss1, ss2, ss3, ds0, ds1, ds2, ds3, eps, stream, ctx.device);
 }
+
+// nodes[0..count) are independent, same-shape, same-eps, same-nb L2_NORM ops validated
+// by the caller; requires ncols < 1024 (the warp reduction path).
+void ggml_sycl_l2_norm_batch(ggml_backend_sycl_context & ctx, ggml_tensor ** nodes, int count) {
+    const ggml_tensor * s0 = nodes[0]->src[0];
+    const int ncols       = (int) s0->ne[0];
+    const int nrows_total = (int) ggml_nrows(s0);
+    float eps;
+    memcpy(&eps, nodes[0]->op_params, sizeof(float));
+    GGML_ASSERT(eps >= 0.0f);
+
+    l2_batch_ptrs p{};
+    for (int t = 0; t < count; ++t) {
+        p.src[t] = (const float *) nodes[t]->src[0]->data;
+        p.dst[t] = (float *) nodes[t]->data;
+    }
+
+    const ggml_tensor * d0 = nodes[0];
+    const size_t        ts = ggml_type_size(GGML_TYPE_F32);
+    l2_batch_strides    st{};
+    st.ne1 = (int) s0->ne[1];
+    st.ne2 = (int) s0->ne[2];
+    st.ss0 = s0->nb[0] / ts; st.ss1 = s0->nb[1] / ts; st.ss2 = s0->nb[2] / ts; st.ss3 = s0->nb[3] / ts;
+    st.ds0 = d0->nb[0] / ts; st.ds1 = d0->nb[1] / ts; st.ds2 = d0->nb[2] / ts; st.ds3 = d0->nb[3] / ts;
+
+    l2_norm_f32_batch_sycl<WARP_SIZE>(p, st, count, ncols, nrows_total, eps, ctx.stream());
+}
@@ -29,4 +29,7 @@ void ggml_sycl_op_group_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst);

 void ggml_sycl_op_l2_norm(ggml_backend_sycl_context& ctx, ggml_tensor* dst);

+#define GGML_SYCL_L2_BATCH_MAX 8
+void ggml_sycl_l2_norm_batch(ggml_backend_sycl_context & ctx, ggml_tensor ** nodes, int count);
+
 #endif // GGML_SYCL_NORM_HPP
@@ -1110,6 +1110,9 @@ struct vk_device_struct {
    vk_pipeline pipeline_cumsum_multipass2_f32;
    vk_pipeline pipeline_argmax_f32;
    vk_pipeline pipeline_count_equal_i32;
+    vk_pipeline pipeline_dsv4_hc_comb_f32;
+    vk_pipeline pipeline_dsv4_hc_pre_f32;
+    vk_pipeline pipeline_dsv4_hc_post_f32;
    std::map<vk_solve_tri_pipeline_state, vk_pipeline> pipeline_solve_tri_f32;
    vk_pipeline pipeline_im2col_f32, pipeline_im2col_f32_f16;
    vk_pipeline pipeline_im2col_3d_f32, pipeline_im2col_3d_f32_f16;
@@ -1467,6 +1470,53 @@ struct vk_op_fwht_push_constants {
    float scale;
 };

+struct vk_op_dsv4_hc_comb_push_constants {
+    uint32_t n_tokens;
+
+    uint32_t nbm0; uint32_t nbm1;
+    uint32_t nbs0;
+    uint32_t nbb0;
+    uint32_t nbd0; uint32_t nbd1; uint32_t nbd2;
+
+    uint32_t m_offset;
+    uint32_t s_offset;
+    uint32_t b_offset;
+    uint32_t d_offset;
+
+    float eps;
+    uint32_t n_iter;
+};
+
+struct vk_op_dsv4_hc_pre_push_constants {
+    uint32_t n_embd;
+    uint32_t n_tokens;
+
+    uint32_t nbx0; uint32_t nbx1; uint32_t nbx2;
+    uint32_t nbw0; uint32_t nbw1;
+    uint32_t nbd0; uint32_t nbd1;
+
+    uint32_t x_offset;
+    uint32_t w_offset;
+    uint32_t d_offset;
+};
+
+struct vk_op_dsv4_hc_post_push_constants {
+    uint32_t n_embd;
+    uint32_t n_tokens;
+
+    uint32_t nbx0; uint32_t nbx1;
+    uint32_t nbr0; uint32_t nbr1; uint32_t nbr2;
+    uint32_t nbp0; uint32_t nbp1;
+    uint32_t nbc0; uint32_t nbc1; uint32_t nbc2;
+    uint32_t nbd0; uint32_t nbd1; uint32_t nbd2;
+
+    uint32_t x_offset;
+    uint32_t r_offset;
+    uint32_t p_offset;
+    uint32_t c_offset;
+    uint32_t d_offset;
+};
+
 struct vk_op_count_experts_push_constants {
    uint32_t ne00;
    uint32_t ne01;
@@ -2631,6 +2681,32 @@ template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk
    GGML_UNUSED(src3);
 }

+template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk_op_dsv4_hc_comb_push_constants &p, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst) {
+    p.m_offset = get_misalign_bytes(ctx, src0) / ggml_type_size(src0->type);
+    p.s_offset = get_misalign_bytes(ctx, src1) / ggml_type_size(src1->type);
+    p.b_offset = get_misalign_bytes(ctx, src2) / ggml_type_size(src2->type);
+    p.d_offset = get_misalign_bytes(ctx, dst)  / ggml_type_size(dst->type);
+
+    GGML_UNUSED(src3);
+}
+
+template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk_op_dsv4_hc_pre_push_constants &p, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst) {
+    p.x_offset = get_misalign_bytes(ctx, src0) / ggml_type_size(src0->type);
+    p.w_offset = get_misalign_bytes(ctx, src1) / ggml_type_size(src1->type);
+    p.d_offset = get_misalign_bytes(ctx, dst)  / ggml_type_size(dst->type);
+
+    GGML_UNUSED(src2);
+    GGML_UNUSED(src3);
+}
+
+template <> void init_pushconst_tensor_offsets(ggml_backend_vk_context * ctx, vk_op_dsv4_hc_post_push_constants &p, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * src2, const ggml_tensor * src3, ggml_tensor * dst) {
+    p.x_offset = get_misalign_bytes(ctx, src0) / ggml_type_size(src0->type);
+    p.r_offset = get_misalign_bytes(ctx, src1) / ggml_type_size(src1->type);
+    p.p_offset = get_misalign_bytes(ctx, src2) / ggml_type_size(src2->type);
+    p.c_offset = get_misalign_bytes(ctx, src3) / ggml_type_size(src3->type);
+    p.d_offset = get_misalign_bytes(ctx, dst)  / ggml_type_size(dst->type);
+}
+
 struct ggml_backend_vk_buffer_context {
    vk_device_ref device;
    vk_buffer dev_buffer;
@@ -5977,6 +6053,16 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
        ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_len, count_experts_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true);
    }

+    // comb holds a token's 4x4 matrix in one 16-lane slice of a subgroup, so it
+    // needs at least 16 lanes, pinned to a known size.
+    if (device->subgroup_basic && device->subgroup_shuffle && device->subgroup_require_full_support && device->subgroup_size >= 16) {
+        const uint32_t tokens_per_workgroup = 4 * (device->subgroup_size / 16);
+        ggml_vk_create_pipeline(device, device->pipeline_dsv4_hc_comb_f32, "dsv4_hc_comb_f32", dsv4_hc_comb_f32_len, dsv4_hc_comb_f32_data, "main", 4, sizeof(vk_op_dsv4_hc_comb_push_constants), {tokens_per_workgroup, 1, 1}, { device->subgroup_size }, 1, true, true, device->subgroup_size);
+    }
+
+    ggml_vk_create_pipeline(device, device->pipeline_dsv4_hc_pre_f32,  "dsv4_hc_pre_f32",  dsv4_hc_pre_f32_len,  dsv4_hc_pre_f32_data,  "main", 3, sizeof(vk_op_dsv4_hc_pre_push_constants),  {256, 1, 1}, { 256 }, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_dsv4_hc_post_f32, "dsv4_hc_post_f32", dsv4_hc_post_f32_len, dsv4_hc_post_f32_data, "main", 5, sizeof(vk_op_dsv4_hc_post_push_constants), {256, 1, 1}, { 256 }, 1);
+
    for (auto &s : device->pipeline_solve_tri_f32) {
        const vk_solve_tri_pipeline_state &state = s.first;

@@ -10204,6 +10290,98 @@ static void ggml_vk_fwht(ggml_backend_vk_context * ctx, vk_context& subctx, cons
    ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { src_buf, dst_buf }, pc, { workgroups_x, 1, 1 });
 }

+static uint32_t ggml_vk_nb_elem(const ggml_tensor * t, int i) {
+    return (uint32_t)(t->nb[i] / ggml_type_size(t->type));
+}
+
+static void ggml_vk_dsv4_hc_comb(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * mixes, const ggml_tensor * scale, const ggml_tensor * base, ggml_tensor * dst) {
+    VK_LOG_DEBUG("ggml_vk_dsv4_hc_comb(" << mixes << ", " << scale << ", " << base << ", " << dst << ")");
+
+    vk_pipeline pipeline = ctx->device->pipeline_dsv4_hc_comb_f32;
+    GGML_ASSERT(pipeline != nullptr);
+
+    const uint32_t n_tokens = (uint32_t)mixes->ne[1];
+
+    ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1);
+
+    const vk_subbuffer mixes_buf = ggml_vk_tensor_subbuffer(ctx, mixes, true);
+    const vk_subbuffer scale_buf = ggml_vk_tensor_subbuffer(ctx, scale, true);
+    const vk_subbuffer base_buf  = ggml_vk_tensor_subbuffer(ctx, base,  true);
+    const vk_subbuffer dst_buf   = ggml_vk_tensor_subbuffer(ctx, dst,   true);
+
+    vk_op_dsv4_hc_comb_push_constants pc = {
+        n_tokens,
+        ggml_vk_nb_elem(mixes, 0), ggml_vk_nb_elem(mixes, 1),
+        ggml_vk_nb_elem(scale, 0),
+        ggml_vk_nb_elem(base,  0),
+        ggml_vk_nb_elem(dst,   0), ggml_vk_nb_elem(dst, 1), ggml_vk_nb_elem(dst, 2),
+        0, 0, 0, 0,
+        ggml_get_op_params_f32(dst, 0),
+        (uint32_t)ggml_get_op_params_i32(dst, 1),
+    };
+    init_pushconst_tensor_offsets(ctx, pc, mixes, scale, base, nullptr, dst);
+
+    ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { mixes_buf, scale_buf, base_buf, dst_buf }, pc, { n_tokens, 1, 1 });
+}
+
+static void ggml_vk_dsv4_hc_pre(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * weights, ggml_tensor * dst) {
+    VK_LOG_DEBUG("ggml_vk_dsv4_hc_pre(" << x << ", " << weights << ", " << dst << ")");
+
+    vk_pipeline pipeline = ctx->device->pipeline_dsv4_hc_pre_f32;
+    GGML_ASSERT(pipeline != nullptr);
+
+    const uint32_t n_embd   = (uint32_t)x->ne[0];
+    const uint32_t n_tokens = (uint32_t)x->ne[2];
+
+    ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1);
+
+    const vk_subbuffer x_buf = ggml_vk_tensor_subbuffer(ctx, x,       true);
+    const vk_subbuffer w_buf = ggml_vk_tensor_subbuffer(ctx, weights, true);
+    const vk_subbuffer d_buf = ggml_vk_tensor_subbuffer(ctx, dst,     true);
+
+    vk_op_dsv4_hc_pre_push_constants pc = {
+        n_embd, n_tokens,
+        ggml_vk_nb_elem(x, 0), ggml_vk_nb_elem(x, 1), ggml_vk_nb_elem(x, 2),
+        ggml_vk_nb_elem(weights, 0), ggml_vk_nb_elem(weights, 1),
+        ggml_vk_nb_elem(dst, 0), ggml_vk_nb_elem(dst, 1),
+        0, 0, 0,
+    };
+    init_pushconst_tensor_offsets(ctx, pc, x, weights, nullptr, nullptr, dst);
+
+    ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { x_buf, w_buf, d_buf }, pc, { n_embd, n_tokens, 1 });
+}
+
+static void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * residual, const ggml_tensor * post, const ggml_tensor * comb, ggml_tensor * dst) {
+    VK_LOG_DEBUG("ggml_vk_dsv4_hc_post(" << x << ", " << residual << ", " << post << ", " << comb << ", " << dst << ")");
+
+    vk_pipeline pipeline = ctx->device->pipeline_dsv4_hc_post_f32;
+    GGML_ASSERT(pipeline != nullptr);
+
+    const uint32_t n_embd   = (uint32_t)x->ne[0];
+    const uint32_t n_tokens = (uint32_t)x->ne[1];
+
+    ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1);
+
+    const vk_subbuffer x_buf = ggml_vk_tensor_subbuffer(ctx, x,        true);
+    const vk_subbuffer r_buf = ggml_vk_tensor_subbuffer(ctx, residual, true);
+    const vk_subbuffer p_buf = ggml_vk_tensor_subbuffer(ctx, post,     true);
+    const vk_subbuffer c_buf = ggml_vk_tensor_subbuffer(ctx, comb,     true);
+    const vk_subbuffer d_buf = ggml_vk_tensor_subbuffer(ctx, dst,      true);
+
+    vk_op_dsv4_hc_post_push_constants pc = {
+        n_embd, n_tokens,
+        ggml_vk_nb_elem(x, 0), ggml_vk_nb_elem(x, 1),
+        ggml_vk_nb_elem(residual, 0), ggml_vk_nb_elem(residual, 1), ggml_vk_nb_elem(residual, 2),
+        ggml_vk_nb_elem(post, 0), ggml_vk_nb_elem(post, 1),
+        ggml_vk_nb_elem(comb, 0), ggml_vk_nb_elem(comb, 1), ggml_vk_nb_elem(comb, 2),
+        ggml_vk_nb_elem(dst,  0), ggml_vk_nb_elem(dst,  1), ggml_vk_nb_elem(dst,  2),
+        0, 0, 0, 0, 0,
+    };
+    init_pushconst_tensor_offsets(ctx, pc, x, residual, post, comb, dst);
+
+    ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { x_buf, r_buf, p_buf, c_buf, d_buf }, pc, { n_embd, n_tokens, 1 });
+}
+
 static void ggml_vk_mul_mat(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx) {
    ggml_tensor * dst = cgraph->nodes[node_idx];
    ggml_tensor * src0 = dst->src[0];
@@ -16222,6 +16400,18 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr
    case GGML_OP_CUMSUM:
        ggml_vk_cumsum(ctx, compute_ctx, src0, node);

+        break;
+    case GGML_OP_DSV4_HC_COMB:
+        ggml_vk_dsv4_hc_comb(ctx, compute_ctx, src0, src1, src2, node);
+
+        break;
+    case GGML_OP_DSV4_HC_PRE:
+        ggml_vk_dsv4_hc_pre(ctx, compute_ctx, src0, src1, node);
+
+        break;
+    case GGML_OP_DSV4_HC_POST:
+        ggml_vk_dsv4_hc_post(ctx, compute_ctx, src0, src1, src2, src3, node);
+
        break;
    case GGML_OP_MEAN:
        ggml_vk_mean(ctx, compute_ctx, src0, node);
@@ -19289,6 +19479,31 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                }
                return false;
            }
+        case GGML_OP_DSV4_HC_COMB:
+        case GGML_OP_DSV4_HC_PRE:
+        case GGML_OP_DSV4_HC_POST:
+            {
+                if (op->type != GGML_TYPE_F32) {
+                    return false;
+                }
+                for (uint32_t i = 0; i < GGML_MAX_SRC; ++i) {
+                    if (op->src[i] && op->src[i]->type != GGML_TYPE_F32) {
+                        return false;
+                    }
+                }
+                // hc is hardcoded to 4 in the shaders. ggml only constrains it
+                // to 4 for COMB, so PRE/POST have to be checked here.
+                if (op->op == GGML_OP_DSV4_HC_PRE && op->src[0]->ne[1] != 4) {
+                    return false;
+                }
+                if (op->op == GGML_OP_DSV4_HC_POST && op->src[1]->ne[1] != 4) {
+                    return false;
+                }
+                if (op->op == GGML_OP_DSV4_HC_COMB) {
+                    return device->pipeline_dsv4_hc_comb_f32 != nullptr;
+                }
+                return true;
+            }
        case GGML_OP_SOLVE_TRI:
            {
                if (op->type != GGML_TYPE_F32 || op->src[0]->type != GGML_TYPE_F32) {
@@ -20277,6 +20492,13 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph *
            tensor_clone = ggml_sum_rows(ggml_ctx, src_clone[0]);
        } else if (tensor->op == GGML_OP_CUMSUM) {
            tensor_clone = ggml_cumsum(ggml_ctx, src_clone[0]);
+        } else if (tensor->op == GGML_OP_DSV4_HC_COMB) {
+            tensor_clone = ggml_dsv4_hc_comb(ggml_ctx, src_clone[0], src_clone[1], src_clone[2],
+                ggml_get_op_params_f32(tensor, 0), ggml_get_op_params_i32(tensor, 1));
+        } else if (tensor->op == GGML_OP_DSV4_HC_PRE) {
+            tensor_clone = ggml_dsv4_hc_pre(ggml_ctx, src_clone[0], src_clone[1]);
+        } else if (tensor->op == GGML_OP_DSV4_HC_POST) {
+            tensor_clone = ggml_dsv4_hc_post(ggml_ctx, src_clone[0], src_clone[1], src_clone[2], src_clone[3]);
        } else if (tensor->op == GGML_OP_MEAN) {
            tensor_clone = ggml_mean(ggml_ctx, src_clone[0]);
        } else if (tensor->op == GGML_OP_ARGMAX) {
@@ -0,0 +1,90 @@
+#version 450
+
+#extension GL_EXT_control_flow_attributes : require
+#extension GL_KHR_shader_subgroup_basic : require
+#extension GL_KHR_shader_subgroup_shuffle : require
+
+// 16 lanes per token, indexed idst + hc*isrc: idst in bits 0..1, isrc in bits 2..3,
+// so subgroupShuffleXor by 1|2 reduces a row and by 4|8 a column.
+
+layout(constant_id = 0) const uint SUBGROUP_SIZE = 32;
+
+layout(local_size_x_id = 0, local_size_y = 4, local_size_z = 1) in;
+
+layout(push_constant) uniform parameter
+{
+    uint n_tokens;
+
+    uint nbm0; uint nbm1;   // mixes
+    uint nbs0;              // scale
+    uint nbb0;              // base
+    uint nbd0; uint nbd1; uint nbd2;   // dst
+
+    uint m_offset;
+    uint s_offset;
+    uint b_offset;
+    uint d_offset;
+
+    float eps;
+    uint n_iter;
+};
+
+layout(binding = 0, std430) readonly buffer M { float data_m[]; };
+layout(binding = 1, std430) readonly buffer S { float data_s[]; };
+layout(binding = 2, std430) readonly buffer B { float data_b[]; };
+layout(binding = 3, std430) writeonly buffer D { float data_d[]; };
+
+const uint hc          = 4;
+const uint comb_offset = 2 * hc;
+
+const uint TOKENS_PER_SUBGROUP = SUBGROUP_SIZE / 16;
+
+void main() {
+    const uint lane = gl_SubgroupInvocationID;
+    const uint blk  = lane >> 4;    // which 16-lane block, i.e. which token
+    const uint idx  = lane & 15;    // idst + hc*isrc
+
+    const uint sg = gl_WorkGroupID.x * gl_WorkGroupSize.y + gl_SubgroupID;
+    const uint it = sg * TOKENS_PER_SUBGROUP + blk;
+
+    // no early return, the shuffles need every lane; out-of-range blocks compute a discarded value
+    const bool in_range = it < n_tokens;
+
+    const float scale_comb = data_s[s_offset + 2 * nbs0];
+
+    float v = 0.0f;
+    if (in_range) {
+        v = data_m[m_offset + (comb_offset + idx) * nbm0 + it * nbm1] * scale_comb
+          + data_b[b_offset + (comb_offset + idx) * nbb0];
+    }
+
+    // Softmax across destinations: the four lanes sharing an isrc.
+    float vmax = max(v, subgroupShuffleXor(v, 1));
+    vmax = max(vmax, subgroupShuffleXor(vmax, 2));
+    v = exp(v - vmax);
+
+    float sum = v + subgroupShuffleXor(v, 1);
+    sum += subgroupShuffleXor(sum, 2);
+    v = v / sum + eps;
+
+    // Normalize columns: equal destination indices are four lanes apart.
+    sum = v + subgroupShuffleXor(v, 4);
+    sum += subgroupShuffleXor(sum, 8);
+    v /= sum + eps;
+
+    for (uint i = 1; i < n_iter; ++i) {
+        sum = v + subgroupShuffleXor(v, 1);
+        sum += subgroupShuffleXor(sum, 2);
+        v /= sum + eps;
+
+        sum = v + subgroupShuffleXor(v, 4);
+        sum += subgroupShuffleXor(sum, 8);
+        v /= sum + eps;
+    }
+
+    if (in_range) {
+        const uint idst = idx & 3;
+        const uint isrc = idx >> 2;
+        data_d[d_offset + idst * nbd0 + isrc * nbd1 + it * nbd2] = v;
+    }
+}
@@ -0,0 +1,83 @@
+#version 450
+
+#extension GL_EXT_control_flow_attributes : require
+
+// Fan one stream back out to hc streams and add the combination-weighted
+// residuals:
+//
+//   dst[i0, idst, it] = x[i0, it]*post[idst, it]
+//                     + sum_isrc residual[i0, isrc, it]*comb[idst, isrc, it]
+
+layout(constant_id = 0) const uint BLOCK_SIZE = 256;
+
+layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
+
+layout(push_constant) uniform parameter
+{
+    uint n_embd;
+    uint n_tokens;
+
+    uint nbx0; uint nbx1;              // x
+    uint nbr0; uint nbr1; uint nbr2;   // residual
+    uint nbp0; uint nbp1;              // post
+    uint nbc0; uint nbc1; uint nbc2;   // comb
+    uint nbd0; uint nbd1; uint nbd2;   // dst
+
+    uint x_offset;
+    uint r_offset;
+    uint p_offset;
+    uint c_offset;
+    uint d_offset;
+};
+
+layout(binding = 0, std430) readonly buffer X { float data_x[]; };
+layout(binding = 1, std430) readonly buffer R { float data_r[]; };
+layout(binding = 2, std430) readonly buffer P { float data_p[]; };
+layout(binding = 3, std430) readonly buffer C { float data_c[]; };
+layout(binding = 4, std430) writeonly buffer D { float data_d[]; };
+
+const uint hc = 4;
+
+shared float post_s[hc];
+shared float comb_s[hc * hc];
+
+void main() {
+    const uint tid = gl_LocalInvocationID.x;
+    const uint it  = gl_WorkGroupID.y;
+
+    if (tid < hc) {
+        post_s[tid] = data_p[p_offset + tid * nbp0 + it * nbp1];
+    }
+    if (tid < hc * hc) {
+        const uint idst = tid & 3;
+        const uint isrc = tid >> 2;
+        comb_s[tid] = data_c[c_offset + idst * nbc0 + isrc * nbc1 + it * nbc2];
+    }
+    barrier();
+
+    // After the barrier, so every invocation reaches it.
+    const uint i0 = gl_WorkGroupID.x * BLOCK_SIZE + tid;
+    if (i0 >= n_embd) {
+        return;
+    }
+
+    const float xv = data_x[x_offset + i0 * nbx0 + it * nbx1];
+
+    const uint rb = r_offset + i0 * nbr0 + it * nbr2;
+
+    float r[hc];
+    [[unroll]]
+    for (uint isrc = 0; isrc < hc; ++isrc) {
+        r[isrc] = data_r[rb + isrc * nbr1];
+    }
+
+    [[unroll]]
+    for (uint idst = 0; idst < hc; ++idst) {
+        float result = xv * post_s[idst];
+        [[unroll]]
+        for (uint isrc = 0; isrc < hc; ++isrc) {
+            result = fma(r[isrc], comb_s[idst + hc * isrc], result);
+        }
+        data_d[d_offset + i0 * nbd0 + idst * nbd1 + it * nbd2] = result;
+    }
+}
@@ -0,0 +1,59 @@
+#version 450
+
+#extension GL_EXT_control_flow_attributes : require
+
+// Collapse the hc residual streams of a token into one, weighted per stream:
+//
+//   dst[i0, it] = sum_ih x[i0, ih, it] * weights[ih, it]
+
+layout(constant_id = 0) const uint BLOCK_SIZE = 256;
+
+layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
+
+layout(push_constant) uniform parameter
+{
+    uint n_embd;
+    uint n_tokens;
+
+    uint nbx0; uint nbx1; uint nbx2;   // x
+    uint nbw0; uint nbw1;              // weights
+    uint nbd0; uint nbd1;              // dst
+
+    uint x_offset;
+    uint w_offset;
+    uint d_offset;
+};
+
+layout(binding = 0, std430) readonly buffer X { float data_x[]; };
+layout(binding = 1, std430) readonly buffer W { float data_w[]; };
+layout(binding = 2, std430) writeonly buffer D { float data_d[]; };
+
+const uint hc = 4;
+
+shared float w[hc];
+
+void main() {
+    const uint tid = gl_LocalInvocationID.x;
+    const uint it  = gl_WorkGroupID.y;
+
+    if (tid < hc) {
+        w[tid] = data_w[w_offset + tid * nbw0 + it * nbw1];
+    }
+    barrier();
+
+    // After the barrier, so every invocation reaches it.
+    const uint i0 = gl_WorkGroupID.x * BLOCK_SIZE + tid;
+    if (i0 >= n_embd) {
+        return;
+    }
+
+    const uint xb = x_offset + i0 * nbx0 + it * nbx2;
+
+    float result = 0.0f;
+    [[unroll]]
+    for (uint ih = 0; ih < hc; ++ih) {
+        result = fma(data_x[xb + ih * nbx1], w[ih], result);
+    }
+
+    data_d[d_offset + i0 * nbd0 + it * nbd1] = result;
+}
@@ -1042,6 +1042,9 @@ void process_shaders() {
    string_to_spv("fwht_f32", "fwht.comp", {});
    string_to_spv("fwht_shmem_f32", "fwht.comp", {{"FWHT_SHMEM", "1"}});
    string_to_spv("count_equal_i32", "count_equal.comp", merge_maps(base_dict, {{"A_TYPE", "int"}, {"B_TYPE", "int"}, {"D_TYPE", "int"}}));
+    string_to_spv("dsv4_hc_comb_f32", "dsv4_hc_comb.comp", {});
+    string_to_spv("dsv4_hc_pre_f32",  "dsv4_hc_pre.comp",  {});
+    string_to_spv("dsv4_hc_post_f32", "dsv4_hc_post.comp", {});
    string_to_spv("cumsum_f32", "cumsum.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
    string_to_spv("cumsum_multipass1_f32", "cumsum_multipass1.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
    string_to_spv("cumsum_multipass2_f32", "cumsum_multipass2.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
@@ -4323,21 +4323,23 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
                            op->type == GGML_TYPE_Q4_0) &&
                           src0->type == GGML_TYPE_F32 && (src1->type == GGML_TYPE_I64 || src1->type == GGML_TYPE_I32));
            break;
-        case GGML_OP_GET_ROWS: {
-            const size_t storage_alignment =
-                ctx->webgpu_global_ctx->capabilities.limits.minStorageBufferOffsetAlignment;
-            const size_t src_address_unit =
-                src0->type == GGML_TYPE_F32 && op->ne[0] % 4 == 0 ? 4 * sizeof(float) : ggml_type_size(src0->type);
-            if (ggml_webgpu_tensor_misalignment(src0, storage_alignment) % src_address_unit != 0) {
+        case GGML_OP_GET_ROWS:
+            {
+                const size_t storage_alignment =
+                    ctx->webgpu_global_ctx->capabilities.limits.minStorageBufferOffsetAlignment;
+                const size_t src_address_unit =
+                    src0->type == GGML_TYPE_F32 && op->ne[0] % 4 == 0 ? 4 * sizeof(float) : ggml_type_size(src0->type);
+                if (ggml_webgpu_tensor_misalignment(src0, storage_alignment) % src_address_unit != 0) {
+                    break;
+                }
+                if (src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 ||
+                    ggml_webgpu_supported_qtype(src0->type)) {
+                    supports_op = (op->type == GGML_TYPE_F32);
+                } else if (src0->type == GGML_TYPE_I32) {
+                    supports_op = op->type == GGML_TYPE_I32;
+                }
                break;
            }
-            if (src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16 || ggml_webgpu_supported_qtype(src0->type)) {
-                supports_op = (op->type == GGML_TYPE_F32);
-            } else if (src0->type == GGML_TYPE_I32) {
-                supports_op = op->type == GGML_TYPE_I32;
-            }
-            break;
-        }
        case GGML_OP_MUL_MAT:
            {
                switch (src1->type) {
@@ -385,6 +385,7 @@ class TensorNameMap:
        MODEL_TENSOR.ATTN_SINKS: (
            "model.layers.{bid}.self_attn.sinks", # openai-moe
            "model.layers.{bid}.self_attn.attention_sink_bias", # mimov2
+            "model.layers.{bid}.self_attn.learnable_sink_param", # hy-v4
        ),

        MODEL_TENSOR.ATTN_GATE: (
@@ -392,6 +393,7 @@ class TensorNameMap:
            "model.layers.{bid}.linear_attn.in_proj_z",  # qwen3.5
            "model.layers.{bid}.self_attn.g_proj",    # step3.5 head-wise attention gate
            "model.layers.{bid}.self_attn.output_gate",  # minimax-01
+            "model.layers.{bid}.self_attn.linear_gate",  # hy-v4
        ),

        # Feed-forward norm
@@ -1329,6 +1331,42 @@ class TensorNameMap:
            "model.layers.{bid}.self_attn.index_q_norm", # MSA
        ),

+        MODEL_TENSOR.HC_ATTN_FN: (
+            "model.layers.{bid}.hc_attn_layer.hc_pre.hc_fn", # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_ATTN_BASE: (
+            "model.layers.{bid}.hc_attn_layer.hc_pre.hc_base", # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_ATTN_SCALE: (
+            "model.layers.{bid}.hc_attn_layer.hc_pre.hc_scale", # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_FFN_FN: (
+            "model.layers.{bid}.hc_mlp_layer.hc_pre.hc_fn", # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_FFN_BASE: (
+            "model.layers.{bid}.hc_mlp_layer.hc_pre.hc_base", # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_FFN_SCALE: (
+            "model.layers.{bid}.hc_mlp_layer.hc_pre.hc_scale", # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_HEAD_FN: (
+            "model.hc_head.hc_head_fn",  # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_HEAD_BASE: (
+            "model.hc_head.hc_head_base", # hy-v4
+        ),
+
+        MODEL_TENSOR.HC_HEAD_SCALE: (
+            "model.hc_head.hc_head_scale", # hy-v4
+        ),
+
        ############################################################################
        # TODO: these do not belong to block_mappings_cfg - move them to mappings_cfg
        MODEL_TENSOR.ENC_OUTPUT_NORM: (
@@ -7206,6 +7206,49 @@ struct test_group_norm_mul_add : public test_case {
    }
 };

+// GGML_OP_L2_NORM x N: independent same-shape norms in one graph (strided qkv views or
+// contiguous), consuming adds nested so the norms stay adjacent in the graph.
+struct test_l2_norm_batch : public test_case {
+    const ggml_type              type;
+    const std::array<int64_t, 4> ne;
+    const int                    n_norms;
+    const float                  eps;
+    const bool                   strided;
+
+    std::string vars() override { return VARS_TO_STR5(type, ne, n_norms, eps, strided); }
+    std::string op_desc(ggml_tensor * t) override { GGML_UNUSED(t); return "L2_NORM_BATCH"; }
+    bool run_whole_graph() override { return true; }
+
+    test_l2_norm_batch(ggml_type type = GGML_TYPE_F32, std::array<int64_t, 4> ne = { 128, 16, 16, 1 },
+                       int n_norms = 4, float eps = 1e-12f, bool strided = true)
+        : type(type), ne(ne), n_norms(n_norms), eps(eps), strided(strided) {}
+
+    ggml_tensor * build_graph(ggml_context * ctx) override {
+        GGML_ASSERT(n_norms >= 2 && n_norms <= 8);
+        ggml_tensor * parent = nullptr;
+        if (strided) {
+            parent = ggml_new_tensor_4d(ctx, type, ne[0], ne[1] * n_norms, ne[2], ne[3]);  // qkv buffer
+        }
+        ggml_tensor * norms[8];
+        for (int t = 0; t < n_norms; ++t) {
+            ggml_tensor * src;
+            if (strided) {
+                src = ggml_view_4d(ctx, parent, ne[0], ne[1], ne[2], ne[3], parent->nb[1], parent->nb[2],
+                                   parent->nb[3], t * ne[1] * parent->nb[1]);
+            } else {
+                src = ggml_new_tensor(ctx, type, 4, ne.data());
+            }
+            norms[t] = ggml_l2_norm(ctx, src, eps);
+        }
+        ggml_tensor * out = norms[n_norms - 1];
+        for (int t = n_norms - 2; t >= 0; --t) {
+            out = ggml_add(ctx, norms[t], out);
+        }
+        ggml_set_name(out, "out");
+        return out;
+    }
+};
+
 // GGML_OP_L2_NORM
 struct test_l2_norm : public test_case {
    const ggml_type type;
@@ -8807,6 +8850,11 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
    test_cases.emplace_back(new test_dsv4_hc_comb(17, 4));
    test_cases.emplace_back(new test_dsv4_hc_comb(257, 8));
    test_cases.emplace_back(new test_dsv4_hc_comb(17, 20));
+    // production n_iter (DeepSeek-V4 uses 20) across batch sizes that cross
+    // subgroup and workgroup boundaries; 1 = single-token decode
+    for (int64_t n_tokens : {1, 256, 336, 512, 513, 1024, 2048}) {
+        test_cases.emplace_back(new test_dsv4_hc_comb(n_tokens, 20));
+    }

    test_cases.emplace_back(new test_dsv4_hc_pre(1, 1));
    test_cases.emplace_back(new test_dsv4_hc_pre(31, 17));
@@ -9490,6 +9538,10 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
            test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false));
            test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, true));
            test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 5, 4, 3 }, eps, false, true));
+            // sibling batching: strided (production shape) and contiguous, 2 and 4 wide
+            test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 5, 4, 3 }, 2, eps, true));
+            test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 5, 4, 3 }, 4, eps, true));
+            test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 5, 4, 3 }, 4, eps, false));
        }
        // row lengths that are not a multiple of 32, for the scalar (33) and float4 (132, 260) paths
        for (uint32_t n : { 33, 132, 260 }) {
@@ -11176,6 +11228,16 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
        }
    }

+    // launch-overhead isolation: single L2_NORM launch vs batched siblings at the GDN
+    // production shape (strided qkv views) -- perf-mode only, the eval list has its own
+    // 2/4-wide coverage
+    for (int n : { 128, 256 }) {
+        test_cases.emplace_back(new test_l2_norm(GGML_TYPE_F32, { n, 16, 16, 1 }, 1e-12f, false, false));
+        test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 16, 16, 1 }, 2, 1e-12f, true));
+        test_cases.emplace_back(new test_l2_norm_batch(GGML_TYPE_F32, { n, 16, 16, 1 }, 4, 1e-12f, true));
+    }
+
+
    return test_cases;
 }

@@ -15,6 +15,7 @@
 #include <cstdio>
 #include <cstdlib>
 #include <filesystem>
+#include <fstream>
 #include <map>
 #include <thread>
 #include <string>
@@ -139,7 +140,7 @@ static const std::vector<std::string> hole = {
    "dflash-model-Q8_0.gguf",
 };

-// unsloth-style naming with UD quants and a suffix MTP file
+// unsloth-style naming with UD quants and an uppercase infix MTP head
 static const std::vector<std::string> unsloth = {
    "model-UD-Q8_K_XL.gguf",
    "mmproj-BF16.gguf",
@@ -180,6 +181,13 @@ static const std::vector<std::string> spark = {
    "dspark-model-MXFP4.gguf",
 };

+// main weights plus a trailing-form mtp head, as in unsloth
+// gemma-4-E2B-it-GGUF (Model-Q4_0-mtp.gguf)
+static const std::vector<std::string> trailing = {
+    "model-BF16.gguf",
+    "model-BF16-mtp.gguf",
+};
+
 // dspark outranks dflash in the type auto-selection
 static const std::vector<std::string> dspark_dflash = {
    "model-Q8_0.gguf",
@@ -223,7 +231,7 @@ static const plan_case plan_cases[] = {
    // no tag and no default match falls back to the first model in the listing
    {"unsloth fallback", unsloth, "test/repo", "", true, true,
     "model-UD-Q8_K_XL.gguf", {"model-UD-Q8_K_XL.gguf"},
-     "mmproj-BF16.gguf", "", "", "", ""},
+     "mmproj-BF16.gguf", "model-MTP-BF16.gguf", "", "", ""},

    // explicit hf_file picks that exact file
    {"flat hf_file", flat, "test/repo", "model-BF16.gguf", false, false,
@@ -258,10 +266,11 @@ static const plan_case plan_cases[] = {
     "model-Q4_K_M.gguf", {"model-Q4_K_M.gguf"},
     "", "mtp-model-Q4_0.gguf", "dflash-model-Q8_0.gguf", "", ""},

-    // the mtp- keyword is case sensitive, a suffix -MTP file is not discovered
+    // the uppercase infix MTP head is a sidecar too; it resolves at the
+    // nearest quant to the tag when the exact one does not exist
    {"unsloth suffix mtp", unsloth, "test/repo:Q8_K_XL", "", true, false,
     "model-UD-Q8_K_XL.gguf", {"model-UD-Q8_K_XL.gguf"},
-     "mmproj-BF16.gguf", "", "", "", ""},
+     "mmproj-BF16.gguf", "model-MTP-BF16.gguf", "", "", ""},

    // vendor prefixes and the dot quant convention both match the tag,
    // first match wins between two files at the same quant
@@ -283,6 +292,45 @@ static const plan_case plan_cases[] = {
    {"spark tag sidecar", spark, "test/repo:BF16", "", true, false,
     "", {},
     "", "", "", "", "dspark-model-BF16.gguf"},
+
+    // a `<quant>-<sidecar>` tag resolves that sidecar alone, without a primary
+    {"hole quant-sidecar tag", hole, "test/repo:Q4_0-mtp", "", false, false,
+     "", {},
+     "", "mtp-model-Q4_0.gguf", "", "", ""},
+
+    {"spark quant-sidecar tag", spark, "test/repo:BF16-dspark", "", false, false,
+     "", {},
+     "", "", "", "", "dspark-model-BF16.gguf"},
+
+    // a bare sidecar tag resolves the sidecar at any quant
+    {"hole bare sidecar tag", hole, "test/repo:mtp", "", false, false,
+     "", {},
+     "", "mtp-model-Q4_0.gguf", "", "", ""},
+
+    // a trailing-form sidecar resolves as the sidecar alone, never as a primary
+    {"trailing quant-sidecar tag", trailing, "test/repo:BF16-mtp", "", false, false,
+     "", {},
+     "", "model-BF16-mtp.gguf", "", "", ""},
+
+    // the plain tag resolves the plain model; the trailing-token file is skipped
+    {"trailing plain tag", trailing, "test/repo:BF16", "", false, false,
+     "model-BF16.gguf", {"model-BF16.gguf"},
+     "", "", "", "", ""},
+
+    // a short-form sidecar (`mmproj-F16.gguf`) resolves by its bare quant tag
+    {"unsloth quant-sidecar tag", unsloth, "test/repo:BF16-mmproj", "", false, false,
+     "", {},
+     "mmproj-BF16.gguf", "", "", "", ""},
+
+    // an uppercase infix sidecar resolves through the lowercase tag
+    {"unsloth uppercase infix tag", unsloth, "test/repo:BF16-mtp", "", false, false,
+     "", {},
+     "", "model-MTP-BF16.gguf", "", "", ""},
+
+    // a sidecar token in the middle of the name resolves as the sidecar alone
+    {"subdir quant-sidecar tag", subdir, "test/repo:Q8_0-mtp", "", false, false,
+     "", {},
+     "", "model-mtp-Q8_0.gguf", "", "", ""},
 };

 static void check_plan(const plan_case & c) {
@@ -468,6 +516,90 @@ static void test_task_assembly() {
    g_repos.clear();
 }

+//
+// cache listing and removal against the isolated LLAMA_CACHE, using the
+// same filename grammar the plan tests exercise above
+//
+
+static void cache_put(const std::string & repo, const std::string & path) {
+    namespace fs = std::filesystem;
+    auto local = fs::path(cached(repo, path));
+    fs::create_directories(local.parent_path());
+    { std::ofstream(local) << "gguf"; }
+    auto repo_dir = local.parent_path().parent_path().parent_path();
+    auto refs = repo_dir / "refs";
+    fs::create_directories(refs);
+    auto ref = refs / "main";
+    if (!fs::exists(ref)) {
+        std::ofstream(ref) << COMMIT << "\n";
+    }
+}
+
+static bool cache_lists(const std::string & repo_tag) {
+    for (const auto & e : common_list_cached_models()) {
+        if (e.to_string() == repo_tag) {
+            return true;
+        }
+    }
+    return false;
+}
+
+static void test_cache_listing_and_remove() {
+    namespace fs = std::filesystem;
+    const std::string repo = "test/eh";
+
+    printf("test-model-resolution: cache listing and removal\n");
+
+    g_context = "cache fixture";
+    cache_put(repo, "gemma-4-E2B-it-BF16.gguf");        // main weights
+    cache_put(repo, "gemma-4-E2B-it-BF16-mtp.gguf");    // trailing-form draft head
+    cache_put(repo, "gemma-4-31B-it-MTP-BF16.gguf");    // uppercase infix draft head
+    cache_put(repo, "mmproj-gemma-4-E2B-it-BF16.gguf");  // mmproj sidecar
+    cache_put(repo, "mmproj-F16.gguf");                 // short-form mmproj
+    cache_put(repo, "model-mtp-Q8_0.gguf");             // mid-name mtp sidecar
+
+    // a sidecar is listed under `<quant>-<sidecar>` in every form and case it
+    // can be named; only the token-less main weights stay a loadable model
+    REQUIRE(cache_lists("test/eh:BF16"));
+    REQUIRE(cache_lists("test/eh:BF16-mtp"));
+    REQUIRE(cache_lists("test/eh:BF16-mmproj"));
+    REQUIRE(cache_lists("test/eh:F16-mmproj"));
+    REQUIRE(cache_lists("test/eh:Q8_0-mtp"));
+    REQUIRE(!cache_lists("test/eh:MTP"));
+
+    // a plain quant tag removes the model files and leaves every sidecar,
+    // whatever form or case its name carries
+    g_context = "remove plain quant";
+    REQUIRE(common_download_remove("test/eh:BF16"));
+    REQUIRE(!fs::exists(cached(repo, "gemma-4-E2B-it-BF16.gguf")));
+    REQUIRE(fs::exists(cached(repo, "gemma-4-E2B-it-BF16-mtp.gguf")));
+    REQUIRE(fs::exists(cached(repo, "gemma-4-31B-it-MTP-BF16.gguf")));
+    REQUIRE(fs::exists(cached(repo, "mmproj-gemma-4-E2B-it-BF16.gguf")));
+    REQUIRE(fs::exists(cached(repo, "mmproj-F16.gguf")));
+    REQUIRE(fs::exists(cached(repo, "model-mtp-Q8_0.gguf")));
+
+    // a `<quant>-<sidecar>` tag removes every sidecar listed under it,
+    // across forms and cases
+    g_context = "remove quant-sidecar";
+    REQUIRE(common_download_remove("test/eh:BF16-mtp"));
+    REQUIRE(!fs::exists(cached(repo, "gemma-4-E2B-it-BF16-mtp.gguf")));
+    REQUIRE(!fs::exists(cached(repo, "gemma-4-31B-it-MTP-BF16.gguf")));
+    REQUIRE(fs::exists(cached(repo, "mmproj-gemma-4-E2B-it-BF16.gguf")));
+
+    // the short form removes by its bare quant tag
+    REQUIRE(common_download_remove("test/eh:F16-mmproj"));
+    REQUIRE(!fs::exists(cached(repo, "mmproj-F16.gguf")));
+
+    REQUIRE(common_download_remove("test/eh:Q8_0-mtp"));
+    REQUIRE(!fs::exists(cached(repo, "model-mtp-Q8_0.gguf")));
+
+    // a bare sidecar tag is ambiguous across quants and is rejected
+    g_context = "remove bare sidecar";
+    cache_put(repo, "mtp-Model-Q4_0.gguf");
+    REQUIRE(!common_download_remove("test/eh:mtp"));
+    REQUIRE(fs::exists(cached(repo, "mtp-Model-Q4_0.gguf")));
+}
+
 int main(void) {
    // unbuffered, so a crash cannot swallow the reports already printed
    setvbuf(stdout, nullptr, _IONBF, 0);
@@ -496,6 +628,7 @@ int main(void) {

    test_plan_resolution();
    test_task_assembly();
+    test_cache_listing_and_remove();

    server.stop();
    server_thread.join();
@@ -80,18 +80,19 @@ struct server_lru_sched {
    }

    // returns "" if no model can be given up
-    std::string pick_victim(std::unique_lock<std::mutex> & lk, const std::string & exclude) {
+    std::string pick_victim(std::unique_lock<std::mutex> & lk) {
        check_lock(lk);
        std::string victim;
        int64_t victim_last_used = 0;
        for (const auto & m : models.mapping) {
-            if (m.first == exclude) {
-                continue;
-            }
            // a busy model is mid-request, one still coming up has no request to finish
            if (m.second.req_count != 0 || !m.second.meta.is_ready_or_sleep()) {
                continue;
            }
+            // already on its way out, or a queued request wants it
+            if (models.stopping_models.count(m.first) || find(m.first)) {
+                continue;
+            }
            if (victim.empty() || m.second.meta.last_used < victim_last_used) {
                victim           = m.first;
                victim_last_used = m.second.meta.last_used;
@@ -109,7 +110,7 @@ struct server_lru_sched {
            SRV_INF("request for name=%s joined the queue, %d waiting\n", model_id.c_str(), e->n_waiters);
            return;
        }
-        queue.push_back({ model_id, 1, false, false });
+        queue.push_back({ model_id, 1, false });
        SRV_INF("models_max reached, request for name=%s queued at position %zu\n",
                model_id.c_str(), queue.size());
    }
@@ -144,85 +145,67 @@ struct server_lru_sched {
        return true;
    }

-    // ok means the model is up: drop the entry, the other waiters just watch its status now
+    // on failure the entry is back in line; on success it stays until its waiters leave,
+    // so the model coming up is never picked as a victim before they use it
    void claim_done(std::unique_lock<std::mutex> & lk, const std::string & model_id, bool ok) {
        check_lock(lk);
+        if (ok) {
+            return;
+        }
        for (auto it = queue.begin(); it != queue.end(); ++it) {
            if (it->model_id == model_id) {
-                if (ok) {
-                    queue.erase(it);
-                } else {
-                    it->loading = false;
-                }
+                it->loading = false;
                return;
            }
        }
    }

-    // a model is on its way out for this entry, so other requests do not also give up one
-    void mark_slot_pending(std::unique_lock<std::mutex> & lk, const std::string & model_id) {
+    // evict idle models while queued requests outnumber the slots that are free or being freed
+    // caller must hold models.mutex; never blocks, so it is safe from any thread
+    void tick(std::unique_lock<std::mutex> & lk) {
        check_lock(lk);
-        if (entry_t * e = find(model_id)) {
-            e->slot_pending = true;
+        if (models.base_params.models_max <= 0 || queue.empty()) {
+            return;
        }
-    }
-
-    // model_id went idle: give up its slot if a queued request needs one
-    // thread-safe, caller must NOT hold models.mutex
-    void on_model_idle(const std::string & model_id) {
-        if (models.base_params.models_max <= 0) {
-            return; // no limit, nothing is ever queued
-        }
-        {
-            std::unique_lock<std::mutex> lk(models.mutex);
-            if (queue.empty()) {
-                return;
-            }
-            size_t promised     = 0;
-            bool   has_unserved = false;
-            for (const auto & e : queue) {
-                if (e.needs_slot()) {
-                    has_unserved = true;
-                } else {
-                    promised++;
-                }
-            }
-            if (!has_unserved) {
-                return;
-            }
-            if ((int) count_running() - (int) promised < models.base_params.models_max) {
-                return; // a slot is already on its way
-            }
-            // never give up a model that a queued request wants
-            for (const auto & e : queue) {
-                if (e.model_id == model_id) {
-                    return;
-                }
-            }
-            auto it = models.mapping.find(model_id);
-            if (it == models.mapping.end() || it->second.req_count != 0 || !it->second.meta.is_ready_or_sleep()) {
-                return;
-            }
-            for (auto & e : queue) {
-                if (!e.slot_pending) {
-                    e.slot_pending = true;
-                    break;
+        int n_running  = 0;
+        int n_stopping = 0;
+        for (const auto & m : models.mapping) {
+            if (m.second.meta.is_running()) {
+                n_running++;
+                if (models.stopping_models.count(m.first)) {
+                    n_stopping++;
                }
            }
        }
-        SRV_INF("model name=%s went idle, giving up its slot to a queued request\n", model_id.c_str());
-        models.unload(model_id);
+        int n_needed  = 0;
+        int n_claimed = 0; // claimed the slot, but load() has not spawned yet
+        for (const auto & e : queue) {
+            if (!e.loading) {
+                n_needed++;
+                continue;
+            }
+            auto it = models.mapping.find(e.model_id);
+            if (it != models.mapping.end() && !it->second.meta.is_running()) {
+                n_claimed++;
+            }
+        }
+        int n_free = models.base_params.models_max - n_running + n_stopping - n_claimed;
+        while (n_free < n_needed) {
+            std::string victim = pick_victim(lk);
+            if (victim.empty()) {
+                return; // all remaining models are busy, wait for a request to end
+            }
+            SRV_INF("evicting idle LRU name=%s for a queued request\n", victim.c_str());
+            models.request_stop(victim);
+            n_free++;
+        }
    }

  private:
    struct entry_t {
        std::string model_id;
-        int  n_waiters;    // requests waiting for this model
-        bool slot_pending; // a model is already being evicted for this entry
-        bool loading;      // one of the waiters is doing the load right now
-
-        // a slot is already coming, or already taken by the load in flight
-        bool needs_slot() const { return !slot_pending && !loading; }
+        int  n_waiters; // requests waiting for this model
+        bool loading;   // one of the waiters is doing the load right now
    };

    entry_t * find(const std::string & model_id) {
@@ -946,7 +929,7 @@ void server_models::unload_lru() {
        if (sched->has_capacity(lk)) {
            return;
        }
-        lru_model_name = sched->pick_victim(lk, "");
+        lru_model_name = sched->pick_victim(lk);
    }
    if (!lru_model_name.empty()) {
        SRV_INF("models_max limit reached, removing LRU name=%s\n", lru_model_name.c_str());
@@ -1169,6 +1152,11 @@ void server_models::load(const std::string & name, const load_options & opts) {
    cv.notify_all();
 }

+void server_models::request_stop(const std::string & name) {
+    stopping_models.insert(name);
+    cv_stop.notify_all();
+}
+
 void server_models::unload(const std::string & name) {
    std::unique_lock<std::mutex> lk(mutex);
    auto it = mapping.find(name);
@@ -1182,13 +1170,12 @@ void server_models::unload(const std::string & name) {
            });
        } else if (it->second.meta.is_running()) {
            SRV_INF("stopping model instance name=%s\n", name.c_str());
-            stopping_models.insert(name);
            if (it->second.meta.status == SERVER_MODEL_STATUS_LOADING) {
                // special case: if model is in loading state, unloading means force-killing it
                SRV_WRN("model name=%s is still loading, force-killing\n", name.c_str());
                it->second.subproc->terminate();
            }
-            cv_stop.notify_all();
+            request_stop(name);
            // status change will be handled by the managing thread
        } else {
            SRV_WRN("model instance name=%s is not running\n", name.c_str());
@@ -1206,8 +1193,7 @@ void server_models::unload_all() {
                inst.subproc->stopped.store(true, std::memory_order_relaxed);
            } else if (inst.meta.is_running()) {
                SRV_INF("stopping model instance name=%s\n", name.c_str());
-                stopping_models.insert(name);
-                cv_stop.notify_all();
+                request_stop(name);
                // status change will be handled by the managing thread
            }
            // moving the thread to join list to avoid deadlock
@@ -1234,6 +1220,8 @@ void server_models::update_status(const std::string & name, const update_status_
        if (!args.progress.is_null()) {
            meta.progress = args.progress;
        }
+        // a model that comes up idle or goes down changes the slot count for queued requests
+        sched->tick(lk);
    }
    // broadcast status change to SSE
    {
@@ -1332,14 +1320,22 @@ bool server_models::remove(const std::string & name) {
        return true;
    }

-    // join before erasing - thread no longer acquires this mutex
-    if (it->second.th.joinable()) {
-        it->second.th.join();
+    // on the cancelled-download path the status flips to DOWNLOADED while the
+    // monitoring thread still has a mutex-guarded step left, so joining under
+    // the lock would deadlock - join outside, as load_models() does
+    std::thread th = std::move(it->second.th);
+    mapping.erase(name);
+    lk.unlock();
+
+    // join first so the monitoring thread's final mutex-guarded cleanup cannot
+    // race the disk removal, then remove from disk without holding the lock
+    // (best-effort: cancelled downloads may have no cached files)
+    if (th.joinable()) {
+        th.join();
    }

-    // remove from disk (best-effort: cancelled downloads may have no cached files)
    bool ok = common_download_remove(name);
-    mapping.erase(name);
+
    if (!ok) {
        SRV_WRN("removing model name=%s from disk returned false (no cached files?)\n", name.c_str());
    }
@@ -1380,13 +1376,11 @@ bool server_models::ensure_model_ready(const std::string & name, const std::func

    bool queued   = false;
    bool did_load = false;
-    std::string victim;
    {
        std::unique_lock<std::mutex> lk(mutex);
        auto it = mapping.find(name);
        if (it != mapping.end() && it->second.meta.status == SERVER_MODEL_STATUS_UNLOADED) {
-            bool has_capacity = sched->has_capacity(lk);
-            if (has_capacity && sched->queue_empty(lk)) {
+            if (sched->has_capacity(lk) && sched->queue_empty(lk)) {
                lk.unlock();
                SRV_INF("model name=%s is not loaded, loading...\n", name.c_str());
                load(name);
@@ -1394,21 +1388,11 @@ bool server_models::ensure_model_ready(const std::string & name, const std::func
            } else {
                // also queue when a slot looks free but others wait already, else they starve
                sched->join(lk, name);
+                sched->tick(lk);
                queued = true;
-                if (!has_capacity) {
-                    // an idle model may sit here right now, do not wait for a request to end
-                    victim = sched->pick_victim(lk, name);
-                    if (!victim.empty()) {
-                        sched->mark_slot_pending(lk, name);
-                    }
-                }
            }
        }
    }
-    if (!victim.empty()) {
-        SRV_INF("evicting idle LRU name=%s to make room for name=%s\n", victim.c_str(), name.c_str());
-        unload(victim);
-    }

    // while queued, this is also where the load happens: the head of the queue does it
    SRV_INF("waiting until model name=%s is fully loaded...\n", name.c_str());
@@ -1470,9 +1454,7 @@ bool server_models::ensure_model_ready(const std::string & name, const std::func
                }
                lk.lock();
                sched->claim_done(lk, name, ok);
-                if (ok) {
-                    queued = false; // entry is gone, the other waiters watch the status now
-                }
+                sched->tick(lk);
                continue;
            }

@@ -1480,6 +1462,7 @@ bool server_models::ensure_model_ready(const std::string & name, const std::func
        }
    } catch (...) {
        leave_queue();
+        sched->tick(lk); // a slot freed for this waiter goes to the next one
        throw;
    }
    leave_queue();
@@ -1529,18 +1512,14 @@ server_http_res_ptr server_models::proxy_request(const server_http_req & req, co
            );

    proxy->cleanup = [this, name]() {
-        bool went_idle = false;
-        {
-            std::unique_lock<std::mutex> lk(mutex);
-            auto it = mapping.find(name);
-            if (it != mapping.end() && it->second.req_count > 0) {
-                it->second.req_count--;
-                went_idle = it->second.req_count == 0;
+        std::unique_lock<std::mutex> lk(mutex);
+        auto it = mapping.find(name);
+        if (it != mapping.end() && it->second.req_count > 0) {
+            it->second.req_count--;
+            if (it->second.req_count == 0) {
+                sched->tick(lk);
            }
        }
-        if (went_idle) {
-            sched->on_model_idle(name);
-        }
    };

    return proxy;
@@ -216,6 +216,10 @@ private:
    // not thread-safe, caller must hold mutex
    void add_model(server_model_meta && meta);

+    // ask the monitoring thread to stop a running instance
+    // not thread-safe, caller must hold mutex
+    void request_stop(const std::string & name);
+
    // notify SSE clients
    void notify_sse(const std::string & event, const std::string & model_id, const json & data = nullptr);

@@ -297,6 +297,26 @@ def test_router_queue_is_fifo():
    assert first.done_at < second.done_at, "queue was not served in arrival order"


+def test_router_queue_two_waiters_share_one_eviction():
+    """two requests that both find the same idle model must both be served in the end"""
+    global server
+    server.models_max = 1
+    server.start()
+
+    _load_model_and_wait(MODEL_A, timeout=120)
+
+    # both arrive while MODEL_A is idle, so both want its slot; only one eviction can happen
+    first = _Bg(lambda: _tokenize(MODEL_B)).start()
+    second = _Bg(lambda: _tokenize(MODEL_C)).start()
+
+    first.join(90)
+    second.join(90)
+
+    first.assert_ok("first queued request")
+    second.assert_ok("second queued request")
+    assert _get_model_status(MODEL_A) == "unloaded"
+
+
 def test_router_no_models_autoload():
    global server
    server.no_models_autoload = True
@@ -71,7 +71,7 @@
 				"svelte-check": "4.6.0",
 				"svelte-sonner": "1.1.1",
 				"tailwind-merge": "3.6.0",
-				"tailwind-variants": "3.2.2",
+				"tailwind-variants": "3.3.1",
 				"tailwindcss": "4.3.0",
 				"tw-animate-css": "1.4.0",
 				"typescript": "5.9.3",
@@ -16048,13 +16048,13 @@
 			}
 		},
 		"node_modules/tailwind-variants": {
-			"version": "3.2.2",
-			"resolved": "https://registry.npmjs.org/tailwind-variants/-/tailwind-variants-3.2.2.tgz",
-			"integrity": "sha512-Mi4kHeMTLvKlM98XPnK+7HoBPmf4gygdFmqQPaDivc3DpYS6aIY6KiG/PgThrGvii5YZJqRsPz0aPyhoFzmZgg==",
+			"version": "3.3.1",
+			"resolved": "https://registry.npmjs.org/tailwind-variants/-/tailwind-variants-3.3.1.tgz",
+			"integrity": "sha512-4pAvwUtM4HKBiRZftncAbpn6V9Hhwoa5Fl7O2u5zbp7Z5Cvu+/o/6+176WY3WCEES209543quG8zFIcXCsc5Jw==",
 			"dev": true,
 			"license": "MIT",
 			"engines": {
-				"node": ">=16.x",
+				"node": ">=16.9.x",
 				"pnpm": ">=7.x"
 			},
 			"peerDependencies": {
@@ -16064,6 +16064,9 @@
 			"peerDependenciesMeta": {
 				"tailwind-merge": {
 					"optional": true
+				},
+				"tailwindcss": {
+					"optional": true
 				}
 			}
 		},
@@ -13,6 +13,7 @@
 		"check:watch": "svelte-kit sync && svelte-check --tsconfig ./tsconfig.json --watch",
 		"reset": "rm -rf .svelte-kit node_modules",
 		"format": "eslint --fix . && prettier --write .",
+		"format:files": "sh -c 'eslint --fix \"$@\" && prettier --write \"$@\"' sh",
 		"lint": "prettier --check . && eslint .",
 		"test": "npm run test:ui -- --run && npm run test:client -- --run && npm run test:unit -- --run && npm run test:e2e",
 		"test:e2e": "playwright test",
@@ -90,7 +91,7 @@
 		"svelte-check": "4.6.0",
 		"svelte-sonner": "1.1.1",
 		"tailwind-merge": "3.6.0",
-		"tailwind-variants": "3.2.2",
+		"tailwind-variants": "3.3.1",
 		"tailwindcss": "4.3.0",
 		"tw-animate-css": "1.4.0",
 		"typescript": "5.9.3",
@@ -18,20 +18,16 @@ import type {
 	ApiErrorResponse,
 	ApiLlamaCppServerProps,
 	ApiModelDataEntry,
-	ApiModelListResponse,
 	ApiModelLoadStage,
+	ApiModelsDownloadRequest,
+	ApiModelsDownloadResponse,
+	ApiModelsListResponse,
+	ApiModelsLoadResponse,
 	ApiModelsSseData,
 	ApiModelsSseEvent,
 	ApiModelsSseProgress,
+	ApiModelsUnloadResponse,
 	ApiProcessingState,
-	ApiRouterModelMeta,
-	ApiRouterModelsListResponse,
-	ApiRouterModelsLoadRequest,
-	ApiRouterModelsLoadResponse,
-	ApiRouterModelsStatusRequest,
-	ApiRouterModelsStatusResponse,
-	ApiRouterModelsUnloadRequest,
-	ApiRouterModelsUnloadResponse,
 	ChatAttachmentDisplayItem,
 	// Chat types
 	ChatMessagePromptProgress,
@@ -86,19 +82,17 @@ declare global {
 		ApiLlamaCppServerProps,
 		ApiModelDataEntry,
 		ApiModelLoadStage,
+		ModelDownloadProgress,
 		ApiModelsSseProgress,
 		ApiModelsSseData,
+		ApiModelsSseDownloadProgressData,
 		ApiModelsSseEvent,
-		ApiModelListResponse,
+		ApiModelsListResponse,
+		ApiModelsLoadResponse,
+		ApiModelsDownloadRequest,
+		ApiModelsDownloadResponse,
+		ApiModelsUnloadResponse,
 		ApiProcessingState,
-		ApiRouterModelMeta,
-		ApiRouterModelsLoadRequest,
-		ApiRouterModelsLoadResponse,
-		ApiRouterModelsStatusRequest,
-		ApiRouterModelsStatusResponse,
-		ApiRouterModelsListResponse,
-		ApiRouterModelsUnloadRequest,
-		ApiRouterModelsUnloadResponse,
 		// Chat types
 		ChatAttachmentDisplayItem,
 		ChatMessagePromptProgress,
@@ -220,8 +220,7 @@ export { default as ChatFormActionModels } from './ChatForm/ChatFormActions/Chat
 */
 export { default as ChatFormActionAddToolsSubmenu } from './ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte';

-/**
- * Dropdown submenu for selecting reasoning effort level.
+/** Dropdown submenu for selecting reasoning effort level.
 *
 * Shows a "Reasoning" sub-menu item with a lightbulb icon indicating
 * thinking status, and a nested list of effort levels.
@@ -12,7 +12,7 @@
 	} from '@lucide/svelte';
 	import { ActionIcon, ModelId } from '$lib/components/app';
 	import { ICON_CLASS_DEFAULT } from '$lib/constants';
-	import { ServerModelStatus } from '$lib/enums';
+	import { ModelCapability, ServerModelStatus } from '$lib/enums';
 	import { modelsStore } from '$lib/stores';
 	import type { ModelOption } from '$lib/types/models';
 	import { modelLoadFraction, modelLoadProgressText } from '$lib/utils';
@@ -60,7 +60,8 @@
 	let loadTitle = $derived(modelLoadProgressText(loadProgress));
 	let modalities = $derived(option.modalities);
 	let capabilities = $derived.by(() => ({
-		reasoning: modelsStore.props.checkModelSupportsThinking(option.model)
+		reasoning: modelsStore.props.checkModelSupportsThinking(option.model),
+		tools: option.capabilities.includes(ModelCapability.TOOL_USE)
 	}));
 </script>

@@ -1,4 +1,7 @@
 export const API_MODELS = {
+	/** Download a model from HuggingFace (ROUTER mode, POST) or cancel/remove it (DELETE) */
+	DELETE: '/models',
+	DOWNLOAD: '/models',
 	LIST: '/v1/models',
 	LOAD: '/models/load',
 	SSE: '/models/sse',
@@ -2,6 +2,12 @@ export const CLI_FLAGS = {
 	AGENT: '--agent',
 	API_KEY: '--api-key',
 	MCP_PROXY: '--ui-mcp-proxy',
+	/** Multimodal projector path; unlocks vision/audio for the model. */
+	MMPROJ: '--mmproj',
+	/** Draft model weights path (long form); the router records it per model. */
+	MODEL_DRAFT: '--model-draft',
+	/** Draft model weights path (short form). */
+	MODEL_DRAFT_SHORT: '-md',
 	SLOTS: '--slots',
 	TOOLS: '--tools'
 } as const;
@@ -0,0 +1,232 @@
+/**
+ * HuggingFace Hub constants.
+ *
+ * URLs, parsing regexes and formatting units for the HuggingFaceService.
+ * Reference: https://huggingface.co/docs/huggingface_hub/package_reference/hf_api
+ */
+
+// API endpoints
+
+export const HF_BASE_URL = 'https://huggingface.co';
+export const HF_API_MODELS_URL = `${HF_BASE_URL}/api/models`;
+export const HF_AVATARS_URL = `${HF_BASE_URL}/api/avatars`;
+
+// Query params
+
+export const HF_FULL_DETAIL_PARAM = 'full=true';
+export const HF_RECURSIVE_TREE_PARAM = 'recursive=true';
+/** Search filter that restricts results to repos containing GGUF files. */
+export const HF_GGUF_FILTER = 'gguf';
+/** Repeatable `expand` query param selecting fields on the list endpoint. */
+export const HF_EXPAND_PARAM = 'expand';
+/**
+ * Fields the model list endpoint omits by default but the discover list rows
+ * render: `gguf` (chat template, context length, param count) drives the
+ * reasoning / tool-use icons and the context badge, `siblings` the vision and
+ * draft-sidecar badges. Without them those parts of a row stay empty.
+ */
+export const HF_MODEL_LIST_EXPAND: readonly string[] = [
+	'author',
+	'downloads',
+	'gguf',
+	'lastModified',
+	'likes',
+	'pipeline_tag',
+	'siblings',
+	// `base_model:` tags, so search rows can show the base org's avatar as the
+	// main avatar with the quant org as the corner badge, like catalog rows.
+	'tags'
+];
+
+// Repo file conventions
+
+export const HF_MAIN_BRANCH = 'main';
+export const HF_README_FILENAME = 'README.md';
+export const HF_RAW_PATH = 'raw';
+export const HF_TREE_PATH = 'tree';
+
+// Pagination
+
+export const HF_LINK_NEXT_REGEX = /<([^>]+)>;\s*rel="next"/;
+/** `Link` response header carrying the next page URL for cursor pagination. */
+export const HF_LINK_HEADER = 'Link';
+
+// Fetch retry
+
+export const HF_RETRY_ATTEMPTS = 3;
+export const HF_RETRY_DELAY_MS = 1000;
+export const HF_HTTP_NOT_FOUND = 404;
+export const HF_HTTP_SERVER_ERROR_MIN = 500;
+
+// Search limits
+
+export const HF_DEFAULT_LIMIT = 50;
+/** Safety cap on `/tree` pagination: more pages means a misbehaving endpoint. */
+export const HF_TREE_MAX_PAGES = 10;
+export const HF_MAX_LIMIT = 100;
+
+// GGUF shard files
+
+/** Matches a split-shard GGUF file name, e.g. `Model-00001-of-00015.gguf`. */
+export const HF_SHARD_REGEX = /-(\d{5})-of-(\d{5})\.gguf$/i;
+/** Index (1-based) of the first shard in a split-shard set. */
+export const HF_FIRST_SHARD = 1;
+/** Zero-padded width of the shard index in a split-shard file name. */
+export const HF_SHARD_PAD_WIDTH = 5;
+
+// Quantization tokens
+
+/** `UD-` (Unsloth Dynamic) custom quantization prefix, e.g. `UD-Q4_K_XL`. */
+export const HF_UD_QUANT_PREFIX = 'UD';
+export const HF_UD_QUANT_PREFIX_REGEX = /^UD-/i;
+/**
+ * Segment marking an Unsloth `shared-` draft head that borrows the target
+ * model's embedding/output weights, e.g. `...-shared-Q4_K_M.gguf`.
+ */
+export const HF_SHARED_DRAFT_TOKEN = 'shared';
+/**
+ * Extracts the leading precision digits from a quant token, e.g.
+ * `Q4_K_XL` -> 4, `IQ2_XXS` -> 2, `TQ1_0` -> 1, `BF16` -> 16.
+ */
+export const HF_QUANT_PRECISION_REGEX = /^(?:I?Q|TQ|BF|F|MXFP)?(\d+)/i;
+
+// Model card tags
+
+/** Matches the `base_model:` tag (plain or `quantized:`), capturing the repo id. */
+export const HF_BASE_MODEL_TAG_REGEX = /^base_model:(?:quantized:)?(.+)$/;
+export const HF_LICENSE_TAG_PREFIX = 'license:';
+export const HF_GATED_TAG = 'gated';
+export const HF_GGUF_TAG = 'gguf';
+export const HF_SAFETENSORS_TAG = 'safetensors';
+
+// Pipeline tasks (logic use only - matching `pipeline_tag` values against tags)
+
+/**
+ * `pipeline_tag` values grouped by the input/output modality they imply, used
+ * to derive a discover row's modality icons. A tag in more than one group (e.g.
+ * `image-to-video`) lights up each modality it belongs to.
+ */
+export const HF_MODALITY_PIPELINE_TAGS: Readonly<
+	Record<'audio' | 'video' | 'vision', readonly string[]>
+> = {
+	audio: [
+		'audio-classification',
+		'audio-to-audio',
+		'automatic-speech-recognition',
+		'text-to-speech',
+		'voice-activity-detection'
+	],
+	video: ['text-to-video', 'image-to-video', 'video-to-video'],
+	vision: ['image-text-to-text', 'image-to-text', 'text-to-image', 'image-to-video']
+};
+
+/** Filename token marking an mmproj sidecar sibling (unlocks vision / audio). */
+export const HF_MMPROJ_FILENAME_TOKEN = 'mmproj';
+
+export const HF_TASK_TAGS: readonly string[] = [
+	'audio-classification',
+	'audio-to-audio',
+	'automatic-speech-recognition',
+	'conversational',
+	'depth-estimation',
+	'feature-extraction',
+	'fill-mask',
+	'image-classification',
+	'image-feature-extraction',
+	'image-segmentation',
+	'image-text-to-text',
+	'image-to-text',
+	'image-to-video',
+	'object-detection',
+	'question-answering',
+	'reinforcement-learning',
+	'robotics',
+	'sentence-similarity',
+	'summarization',
+	'text2text-generation',
+	'text-classification',
+	'text-generation',
+	'text-to-image',
+	'text-to-speech',
+	'text-to-video',
+	'token-classification',
+	'translation',
+	'video-to-video',
+	'voice-activity-detection',
+	'zero-shot-classification'
+];
+
+// Formatting
+
+export const BYTE = 1;
+export const KILOBYTE = 1_000;
+export const MEGABYTE = 1_000_000;
+export const GIGABYTE = 1_000_000_000;
+export const TERABYTE = 1_000_000_000_000;
+
+/**
+ * Matches a human size string (`177GB`, `1.2 TB`, `500MB`), capturing the
+ * numeric value and its unit suffix. Used by `parseSizeBytes`.
+ */
+export const HF_SIZE_STRING_REGEX = /^\s*([\d.]+)\s*([a-z]+)\s*$/i;
+
+/**
+ * Byte multiplier for a size suffix (`k` kilobyte, `m` megabyte, ...) as used by
+ * the llama.app catalog `size` strings, whose suffix is lowercase.
+ */
+export const HF_SIZE_SUFFIX_BYTES: Readonly<Record<string, number>> = {
+	b: BYTE,
+	g: GIGABYTE,
+	k: KILOBYTE,
+	m: MEGABYTE,
+	t: TERABYTE
+};
+
+export const BYTE_LABEL = 'B';
+export const KILOBYTE_LABEL = 'KB';
+export const MEGABYTE_LABEL = 'MB';
+export const GIGABYTE_LABEL = 'GB';
+
+/** Count suffixes for compact number formatting, e.g. `1.5K`, `2.0M`. */
+export const KILO_LABEL = 'K';
+export const MEGA_LABEL = 'M';
+export const GIGA_LABEL = 'B';
+
+// Relative time
+
+export const MS_PER_DAY = 1000 * 60 * 60 * 24;
+export const DAYS_PER_WEEK = 7;
+/** Rough month length in days, used to bucket relative timestamps. */
+export const DAYS_PER_MONTH = 30;
+export const DAYS_PER_YEAR = 365;
+
+export const TODAY_LABEL = 'Today';
+export const YESTERDAY_LABEL = 'Yesterday';
+export const DAYS_AGO_LABEL = 'days ago';
+export const WEEKS_AGO_LABEL = 'weeks ago';
+export const MONTHS_AGO_LABEL = 'months ago';
+export const YEARS_AGO_LABEL = 'years ago';
+
+// Cache paths
+
+/**
+ * Matches a local HF cache file path
+ * (`.../models--<org>--<name>/snapshots/<sha>/<file>`), capturing the repo
+ * directory name and the repo-relative file path.
+ */
+export const HF_CACHE_PATH_REGEX = /models--(.+?)\/snapshots\/[^/]+\/(.+)$/;
+/** Separator between org and name segments in an HF cache directory name. */
+export const HF_CACHE_DIR_SEPARATOR = '--';
+
+// README
+
+/** Matches a leading YAML frontmatter block (--- ... ---) in a markdown document. */
+export const HF_FRONTMATTER_REGEX = /^---\r?\n[\s\S]*?\r?\n---\r?\n?/;
+
+// Param counts
+
+/**
+ * Best-effort parameter count token in a model id/name, e.g. `27B` from
+ * `Qwen3.8-27B-GGUF` or `300M` from `embeddinggemma-300M-GGUF`.
+ */
+export const HF_PARAM_COUNT_REGEX = /(?:^|[^a-z0-9])(\d+(?:[._]\d+)?)\s*([bm])(?![a-z0-9])/i;
@@ -10,7 +10,8 @@ import {
 	Image as ImageIcon,
 	Lightbulb as ReasoningIcon,
 	Mic as AudioIcon,
-	Video as VideoIcon
+	Video as VideoIcon,
+	Wrench as ToolUseIcon
 } from '@lucide/svelte';
 import { FileTypeCategory, ModelCapability, ModelModality } from '$lib/enums';
 import type { ModelCapabilities, ModelModalities } from '$lib/types/models';
@@ -49,16 +50,19 @@ export const MODALITY_FLAG_KEYS: Record<
 };

 export const CAPABILITY_ICONS: Record<ModelCapability, Component> = {
-	[ModelCapability.REASONING]: ReasoningIcon
+	[ModelCapability.REASONING]: ReasoningIcon,
+	[ModelCapability.TOOL_USE]: ToolUseIcon
 } as const;

 export const CAPABILITY_LABELS: Record<ModelCapability, string> = {
-	[ModelCapability.REASONING]: 'Reasoning'
+	[ModelCapability.REASONING]: 'Reasoning',
+	[ModelCapability.TOOL_USE]: 'Tool use'
 } as const;

 /** Maps a ModelCapability to the boolean flag it drives on the ModelCapabilities type */
 export const CAPABILITY_FLAG_KEYS: Record<ModelCapability, keyof ModelCapabilities> = {
-	[ModelCapability.REASONING]: 'reasoning'
+	[ModelCapability.REASONING]: 'reasoning',
+	[ModelCapability.TOOL_USE]: 'tools'
 };

 // Shared SVG icon strings for copy and preview buttons
@@ -45,6 +45,9 @@ export * from './message-export.constants';
 export * from './path-display.constants';
 export * from './model-id.constants';
 export * from './model-loading.constants';
+export * from './models-discover.constants';
+export * from './model-compatibility.constants';
+export * from './huggingface.constants';
 export * from './precision.constants';
 export * from './pwa.constants';
 export * from './routes.constants';
@@ -0,0 +1,32 @@
+/**
+ * Model memory-fit constants.
+ *
+ * Mirrors the app's compatibility check (Model+Compatibility.swift):
+ *   budget      = RAM x RAM_BUDGET_RATIO - RAM_OVERHEAD_MB
+ *   weightBytes = fileBytes x QUANT_WEIGHT
+ * a file fits when weightBytes <= budget. Kept here so the estimation util and
+ * any caller share one source.
+ */
+
+/** Bytes in one mebibyte (MiB), used to convert a file size to MB. */
+export const MIB_BYTES = 1_048_576;
+
+/** MB in one GB. */
+export const MB_PER_GB = 1024;
+
+/** Overhead multiplier applied to the file size when estimating weight memory. */
+export const QUANT_WEIGHT = 1.05;
+
+/** Share of RAM the app allows the model to occupy. */
+export const RAM_BUDGET_RATIO = 0.75;
+
+/** Fixed RAM overhead (MB) reserved for the system and KV cache. */
+export const RAM_OVERHEAD_MB = 2048;
+
+/**
+ * Memory tiers (GB) covering the RAM sizes common machines ship with, in
+ * small enough steps that the requirement reads honestly. Device-agnostic on
+ * purpose: the server exposes no host RAM, so the UI presents the tier and
+ * lets the user judge.
+ */
+export const MEM_TIERS = [4, 6, 8, 12, 16, 24, 32, 48, 64, 96, 128, 192, 256, 384, 512, 768, 1024];
@@ -2,35 +2,48 @@
 * Parsing of `org/ModelName[-tag][:quant]` style model IDs.
 */

+import { ModelAuxSidecar, ModelDraftSidecar } from '$lib/enums';
+
+/** Any sidecar file type: a draft variant or an auxiliary sidecar like mmproj. */
+export type ModelSidecar = ModelDraftSidecar | ModelAuxSidecar;
+
+/** All sidecar filename tokens: the bare lowercase enum values, e.g. `mtp`, `mmproj`. */
+export const SIDECAR_TOKENS: string[] = [
+	...Object.values(ModelDraftSidecar),
+	...Object.values(ModelAuxSidecar)
+];
+
+/** Separator between token alternatives in the sidecar regexes. */
+const REGEX_ALTERNATION_SEPARATOR = '|';
+const SIDECAR_TOKEN_ALTERNATION = SIDECAR_TOKENS.join(REGEX_ALTERNATION_SEPARATOR);
+
 export const MODEL_ID = {
 	/**
 	 * Matches an activated-parameter-count segment, e.g. `A10B`, `a2.4b`.
 	 * The leading `A`/`a` distinguishes it from a regular params segment.
 	 */
-	ACTIVATED_PARAMS_RE: /^[Aa]\d+(\.\d+)?[BbMmKkTt]$/,
+	ACTIVATED_PARAMS_REGEX: /^[Aa]\d+(\.\d+)?[BbMmKkTt]$/,

 	/** Matches prefix for custom quantization types, e.g. `UD-Q8_K_XL`. */
-	CUSTOM_QUANTIZATION_PREFIX_RE: /^UD$/i,
+	CUSTOM_QUANTIZATION_PREFIX_REGEX: /^UD$/i,
 	/** Container format segments to exclude from tags (every model uses these). */
 	IGNORED_SEGMENTS: new Set(['GGUF', 'GGML']),
 	/** Sentinel value returned by `indexOf` when a substring is not found. */
 	NOT_FOUND: -1,
-
 	/** Separates `<org>` from `<model>` in a model ID, e.g. `org/ModelName`. */
 	ORG_SEPARATOR: '/',
-
 	/**
 	 * Matches a parameter-count segment, e.g. `7B`, `1.5b`, `120M`.
 	 * The optional leading `E` covers effective-parameter sizes, e.g. Gemma's
 	 * `E2B`/`E4B` (MatFormer models sized by resident params).
 	 */
-	PARAMS_RE: /^[Ee]?\d+(\.\d+)?[BbMmKkTt]$/,
+	PARAMS_REGEX: /^[Ee]?\d+(\.\d+)?[BbMmKkTt]$/,

 	/**
 	 * Matches a quantization/precision segment, e.g. `Q4_K_M`, `IQ4_XS`, `F16`, `BF16`, `MXFP4`.
 	 * Case-insensitive to handle both uppercase and lowercase inputs.
 	 */
-	QUANTIZATION_SEGMENT_RE: /^(I?Q\d+(_[A-Z0-9]+)*|F\d+|BF\d+|MXFP\d+(_[A-Z0-9]+)*)$/i,
+	QUANTIZATION_SEGMENT_REGEX: /^(I?Q\d+(_[A-Z0-9]+)*|F\d+|BF\d+|MXFP\d+(_[A-Z0-9]+)*)$/i,

 	/** Separates the model path from the quantization tag, e.g. `model:Q4_K_M`. */
 	QUANTIZATION_SEPARATOR: ':',
@@ -38,6 +51,36 @@ export const MODEL_ID = {
 	/** Separates named segments within the model path, e.g. `ModelName-7B-GGUF`. */
 	SEGMENT_SEPARATOR: '-',

+	/**
+	 * Sidecar token between name segments, e.g. `Model-mtp-Q4_0.gguf`,
+	 * `model-eagle3-BF16.gguf`. Captures the name head and tail around the
+	 * token; same case-insensitive rule as the prefix form.
+	 */
+	SIDECAR_INFIX_REGEX: new RegExp(`^(.*)-(${SIDECAR_TOKEN_ALTERNATION})-(.+)$`, 'i'),
+
+	/**
+	 * Sidecar prefix that wraps a model id with a sidecar type, e.g.
+	 * `mtp-<name>.gguf`, `dflash-<name>.gguf`, `dspark-<name>.gguf`,
+	 * `eagle3-<name>.gguf`, `mmproj-<name>.gguf`. Captures the bare type
+	 * token for typed lookup.
+	 *
+	 * The token matches case-insensitively (real repos ship uppercase
+	 * heads, e.g. `Model-MTP-BF16.gguf`) and is normalized through
+	 * `sidecarFromFileToken`; the server's filename grammar
+	 * (common/download.cpp) matches the same segments.
+	 */
+	SIDECAR_PREFIX_REGEX: new RegExp(`^(${SIDECAR_TOKEN_ALTERNATION})-(.*)$`, 'i'),
+
+	/**
+	 * Trailing `-<type>` suffix marking a GGUF with an embedded draft in the
+	 * same weight file (MTP) or a sidecar download entry, e.g.
+	 * `Hy3-IQ1_M-mtp.gguf`, `Q4_K_M-dspark`. An optional `-draft` tail covers
+	 * standalone sidecar files, e.g. `Model-mtp-draft.gguf`. The captured
+	 * prefix is the candidate model id; the caller decides whether it looks
+	 * quantized. Case-insensitive, like the prefix form.
+	 */
+	SIDECAR_SUFFIX_REGEX: new RegExp(`^(.*)-(${SIDECAR_TOKEN_ALTERNATION})(-draft)?$`, 'i'),
+
 	/** Matches a trailing weight file extension, e.g. `model.gguf` -> `model`. */
-	WEIGHT_EXTENSION_RE: /\.(gguf|ggml)$/i
+	WEIGHT_EXTENSION_REGEX: /\.(gguf|ggml)$/i
 };
@@ -0,0 +1,14 @@
+/**
+ * Models discover constants.
+ *
+ * Endpoints and settings for the Models Discover dialog.
+ */
+
+/** llama.app model catalog used as the default model list. Online-only source; the discover feature requires an internet connection anyway. */
+export const MODELS_DISCOVER_CATALOG_URL = 'https://llama.app/v1/catalog.json';
+
+/**
+ * Catalog repos fetched in parallel per batch; small on purpose so the HF API
+ * is not hit with the whole catalog at once.
+ */
+export const MODELS_DISCOVER_CATALOG_BATCH = 4;
@@ -15,6 +15,9 @@ export const STORAGE_APP_NAME_DEPRECATED = 'LlamaCppWebui';
 export const DB_APP_NAME_DEPRECATED = 'LlamacppWebui';

 export const ALWAYS_ALLOWED_TOOLS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.alwaysAllowedTools`;
+
+/** Paused model download ids (`<repo>:<tag>`), restored on the next page load. */
+export const PAUSED_MODEL_DOWNLOADS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.pausedModelDownloads`;
 export const CONFIG_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.config`;
 export const DISABLED_TOOLS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.disabledTools`;

@@ -0,0 +1,35 @@
+/**
+ * HuggingFace Hub enums.
+ *
+ * Values mirror the strings used by the HF REST API
+ * (https://huggingface.co/docs/huggingface_hub/package_reference/hf_api)
+ * so they can be sent and compared directly.
+ */
+
+/** Sort field for /api/models search queries. */
+export enum HfModelSort {
+	CREATED_AT = 'createdAt',
+	DOWNLOADS = 'downloads',
+	LAST_MODIFIED = 'lastModified',
+	LIKES = 'likes',
+	TRENDING_SCORE = 'trendingScore'
+}
+
+/**
+ * Where the sidecar token (`mtp` / `dflash` / `mmproj` / ...) sits in the
+ * filename.
+ * - `prefix`  sidecar file that lives next to the main weights, e.g. `mtp-Q4_0.gguf`
+ * - `suffix`  embedded draft baked into the main weights, e.g. `Hy3-IQ1_M-mtp.gguf`
+ * - `infix`   standalone sidecar named between head and quant, e.g. `model-mtp-Q8_0.gguf`
+ */
+export enum SidecarForm {
+	INFIX = 'infix',
+	PREFIX = 'prefix',
+	SUFFIX = 'suffix'
+}
+
+/** Entry type in a model repository file tree (`/tree` responses). */
+export enum HfEntryType {
+	DIRECTORY = 'directory',
+	FILE = 'file'
+}
@@ -57,6 +57,8 @@ export {
 	SpecialFileType
 } from './files.enums';

+export { HfEntryType, HfModelSort, SidecarForm } from './huggingface.enums';
+
 export {
 	MCPConnectionPhase,
 	MCPLogLevel,
@@ -67,7 +69,15 @@ export {
 	JsonSchemaType
 } from './mcp.enums';

-export { ModelCapability, ModelModality } from './model.enums';
+export {
+	ModelAuxSidecar,
+	ModelCapability,
+	ModelDraftSidecar,
+	ModelModality,
+	ModelSelectableFileKind
+} from './model.enums';
+
+export { ModelDownloadStopRequest } from './model.enums';

 export { ServerRole, ServerModelStatus, ServerModelsSseEventType } from './server.enums';

@@ -6,5 +6,52 @@ export enum ModelModality {
 }

 export enum ModelCapability {
-	REASONING = 'REASONING'
+	REASONING = 'reasoning',
+	TOOL_USE = 'tools'
+}
+
+/**
+ * Speculative-decoding draft sidecars (server spec-type draft-*).
+ * Filenames use the lowercase token, e.g. `mtp-<name>.gguf` or `-mtp` suffix.
+ */
+export enum ModelDraftSidecar {
+	/** DFlash block-diffusion draft (spec-type draft-dflash). */
+	DFLASH = 'dflash',
+	/** DSpark block-diffusion draft (spec-type draft-dspark). */
+	DSPARK = 'dspark',
+	/** EAGLE-3 speculative draft (spec-type draft-eagle3). */
+	EAGLE3 = 'eagle3',
+	/** Multi-token-prediction draft head (spec-type draft-mtp). */
+	MTP = 'mtp'
+}
+
+/**
+ * Non-draft sidecar file types. A sidecar is any auxiliary GGUF file
+ * accompanying the main model weights.
+ */
+export enum ModelAuxSidecar {
+	/** Importance-matrix data used to build imatrix quants; not loaded at serve time. */
+	IMATRIX = 'imatrix',
+	/** Multimodal projector: unlocks vision and/or audio input modalities. */
+	MMPROJ = 'mmproj'
+}
+
+/**
+ * Role of a selectable GGUF in the download options: the main weights, a
+ * speculative-decoding draft sidecar, or an auxiliary sidecar (mmproj).
+ */
+export enum ModelSelectableFileKind {
+	AUX = 'aux',
+	DRAFT = 'draft',
+	MAIN = 'main'
+}
+
+/**
+ * Why an in-flight download is being stopped, so the terminal `download_failed`
+ * feed event can be attributed: a user pause (resumable) or a user cancel
+ * (discard). Distinguishes these from a genuine download failure.
+ */
+export enum ModelDownloadStopRequest {
+	CANCEL = 'cancel',
+	PAUSE = 'pause'
 }
@@ -13,6 +13,10 @@ export enum ServerRole {
 * Used as the `value` field in the status object from /models endpoint
 */
 export enum ServerModelStatus {
+	DOWNLOAD_FAILED = 'download_failed',
+	DOWNLOAD_FINISHED = 'download_finished',
+	DOWNLOADED = 'downloaded',
+	DOWNLOADING = 'downloading',
 	FAILED = 'failed',
 	LOADED = 'loaded',
 	LOADING = 'loading',
@@ -26,6 +30,8 @@ export enum ServerModelStatus {
 * tools/server/server-models.cpp from the C++ server.
 */
 export enum ServerModelsSseEventType {
+	DOWNLOAD_FAILED = 'download_failed',
+	DOWNLOAD_FINISHED = 'download_finished',
 	DOWNLOAD_PROGRESS = 'download_progress',
 	MODEL_REMOVE = 'model_remove',
 	MODEL_STATUS = 'model_status',
@@ -0,0 +1,794 @@
+import { PATH_SEPARATOR } from '$lib/constants';
+import {
+	BYTE,
+	BYTE_LABEL,
+	DAYS_AGO_LABEL,
+	DAYS_PER_MONTH,
+	DAYS_PER_WEEK,
+	DAYS_PER_YEAR,
+	GIGA_LABEL,
+	GIGABYTE,
+	GIGABYTE_LABEL,
+	HF_API_MODELS_URL,
+	HF_AVATARS_URL,
+	HF_BASE_MODEL_TAG_REGEX,
+	HF_BASE_URL,
+	HF_CACHE_DIR_SEPARATOR,
+	HF_CACHE_PATH_REGEX,
+	HF_DEFAULT_LIMIT,
+	HF_FIRST_SHARD,
+	HF_FRONTMATTER_REGEX,
+	HF_FULL_DETAIL_PARAM,
+	HF_GATED_TAG,
+	HF_GGUF_FILTER,
+	HF_GGUF_TAG,
+	HF_HTTP_NOT_FOUND,
+	HF_HTTP_SERVER_ERROR_MIN,
+	HF_LICENSE_TAG_PREFIX,
+	HF_LINK_HEADER,
+	HF_LINK_NEXT_REGEX,
+	HF_MAIN_BRANCH,
+	HF_MAX_LIMIT,
+	HF_MODEL_LIST_EXPAND,
+	HF_PARAM_COUNT_REGEX,
+	HF_QUANT_PRECISION_REGEX,
+	HF_RAW_PATH,
+	HF_README_FILENAME,
+	HF_RECURSIVE_TREE_PARAM,
+	HF_RETRY_ATTEMPTS,
+	HF_RETRY_DELAY_MS,
+	HF_SAFETENSORS_TAG,
+	HF_SHARD_PAD_WIDTH,
+	HF_SHARD_REGEX,
+	HF_SHARED_DRAFT_TOKEN,
+	HF_SIZE_STRING_REGEX,
+	HF_SIZE_SUFFIX_BYTES,
+	HF_TASK_TAGS,
+	HF_TREE_MAX_PAGES,
+	HF_TREE_PATH,
+	HF_UD_QUANT_PREFIX,
+	HF_UD_QUANT_PREFIX_REGEX,
+	KILO_LABEL,
+	KILOBYTE,
+	KILOBYTE_LABEL,
+	MEGA_LABEL,
+	MEGABYTE,
+	MEGABYTE_LABEL,
+	MODELS_DISCOVER_CATALOG_URL,
+	MONTHS_AGO_LABEL,
+	MS_PER_DAY,
+	TODAY_LABEL,
+	WEEKS_AGO_LABEL,
+	YEARS_AGO_LABEL,
+	YESTERDAY_LABEL
+} from '$lib/constants';
+import { MODEL_ID, type ModelSidecar } from '$lib/constants';
+import { HfEntryType, HfModelSort, SidecarForm } from '$lib/enums';
+import type {
+	HfCatalogEntry,
+	HfModelDetailInfo,
+	HfModelInfo,
+	HfModelSearchParams,
+	HfModelSibling
+} from '$lib/types/huggingface';
+import { sidecarFromFileToken } from '$lib/utils';
+
+/**
+ * HuggingFaceService - Service for browsing and searching GGUF models on Hugging Face Hub
+ */
+export class HuggingFaceService {
+	private static readonly BASE_URL = HF_API_MODELS_URL;
+
+	// Cached base model lookups keyed by repo id, so repeated selector opens
+	// never re-hit the HF API for the same repo.
+	private static baseModelCache = new Map<string, { org: string; name: string } | null>();
+
+	private static baseModelPending = new Map<
+		string,
+		Promise<{ org: string; name: string } | null>
+	>();
+
+	/**
+	 * Map of quant token to its average bit-depth in bits-per-weight (bpw).
+	 */
+	private static readonly QUANT_BIT_DEPTH: Record<string, number> = {
+		BF16: 16,
+		F16: 16,
+		IQ1_M: 1,
+		IQ1_S: 1,
+		IQ1_XS: 1,
+		IQ1_XXS: 1,
+		IQ2_M: 2,
+		IQ2_S: 2,
+		IQ2_XS: 2,
+		IQ2_XXS: 2,
+		IQ3_M: 3,
+		IQ3_S: 3,
+		IQ3_XS: 3,
+		IQ3_XXS: 3,
+		Q2_K: 2,
+		Q2_K_M: 2,
+		Q2_K_S: 2,
+		Q3_K: 3,
+		Q3_K_L: 3,
+		Q3_K_M: 3,
+		Q3_K_S: 3,
+		Q4_0: 4,
+		Q4_1: 4,
+		Q4_K: 4,
+		Q4_K_M: 4,
+		Q4_K_S: 4,
+		Q5_0: 5,
+		Q5_1: 5,
+		Q5_K: 5,
+		Q5_K_M: 5,
+		Q5_K_S: 5,
+		Q6_K: 6,
+		Q8_0: 8
+	};
+
+	/**
+	 * Collapse split GGUF shard sets (`-00001-of-00015.gguf`, ...) to their first
+	 * shard, summing every shard's size so the kept entry reflects the whole
+	 * quant. Non-sharded files pass through unchanged. Downloads are tag-based
+	 * (`repo:quant`), so the first shard is enough to represent the set.
+	 */
+	static collapseGgufShards(siblings: HfModelSibling[]): HfModelSibling[] {
+		const sizeByPath = new Map(siblings.map((f) => [f.path, f.size ?? 0]));
+		const result: HfModelSibling[] = [];
+
+		for (const file of siblings) {
+			const match = HF_SHARD_REGEX.exec(file.path);
+
+			if (!match) {
+				result.push(file);
+
+				continue;
+			}
+
+			// Keep only the first shard; its size becomes the whole shard set's.
+			if (Number(match[1]) !== HF_FIRST_SHARD) continue;
+
+			const total = Number(match[2]);
+			const stem = file.path.slice(0, file.path.length - match[0].length);
+
+			let size = 0;
+
+			for (let i = HF_FIRST_SHARD; i <= total; i++) {
+				const shard = HuggingFaceService.shardPath(stem, i, total);
+
+				size += sizeByPath.get(shard) ?? 0;
+			}
+
+			result.push({ ...file, size });
+		}
+
+		return result;
+	}
+
+	// GGUF Model Browsing
+
+	/**
+	 * Extract the GGUF quantization token (e.g. `Q4_K_M`) and any sidecar type
+	 * (`mtp`, `dflash`, `mmproj`, ...) from a `.gguf` filename. The sidecar token
+	 * shows up either as a sidecar prefix (`mtp-<name>.gguf`, `dflash-<name>.gguf`,
+	 * `mmproj-<name>.gguf`), as a `-mtp` suffix, or as the whole filename
+	 * (`imatrix.gguf`); a `-draft` tail marks a standalone sidecar file
+	 * (`Model-MTP-draft.gguf`).
+	 *
+	 * `sidecarForm` records which side of the filename the sidecar token sat
+	 * on so callers can render badges differently (e.g. prefix on the left of
+	 * the quant label, suffix appended to it).
+	 * `quant` is `null` for files that don't carry a bit-depth token
+	 * (e.g. `*-BF16.gguf`); `sidecar` is `null` if no sidecar flag is present.
+	 * Returns `null` only when the filename doesn't end in `.gguf`.
+	 */
+	static extractQuantMeta(filename: string): {
+		quant: string | null;
+		/** Draft-head-only variant borrowing embed/output weights from the target model. */
+		shared: boolean;
+		sidecar: ModelSidecar | null;
+		sidecarForm: SidecarForm | null;
+	} | null {
+		if (!MODEL_ID.WEIGHT_EXTENSION_REGEX.test(filename)) return null;
+
+		// HF repos may nest sidecars in a folder (e.g. `MTP/mtp-Model-Q4_0.gguf`);
+		// parse the file name only, the folder adds no quant information.
+		let source = (filename.split(PATH_SEPARATOR).pop() ?? filename).replace(
+			MODEL_ID.WEIGHT_EXTENSION_REGEX,
+			''
+		);
+		let sidecar: ModelSidecar | null = null;
+		let sidecarForm: SidecarForm | null = null;
+
+		// A file named just the sidecar token (`imatrix.gguf`) is the sidecar
+		// itself: no name or quant segments to parse.
+		const bareSidecar = sidecarFromFileToken(source.toLowerCase());
+
+		if (bareSidecar) {
+			return { quant: null, shared: false, sidecar: bareSidecar, sidecarForm: SidecarForm.PREFIX };
+		}
+
+		const prefixMatch = source.match(MODEL_ID.SIDECAR_PREFIX_REGEX);
+
+		if (prefixMatch) {
+			sidecar = sidecarFromFileToken(prefixMatch[1].toLowerCase());
+			sidecarForm = SidecarForm.PREFIX;
+			source = prefixMatch[2];
+		} else {
+			const suffixMatch = source.match(MODEL_ID.SIDECAR_SUFFIX_REGEX);
+
+			if (suffixMatch) {
+				// Take the suffix sidecar even when the head carries no quant:
+				// embedded drafts end in one (`Hy3-IQ1_M-mtp`), standalone sidecar
+				// files do not (`Model-mtp-draft`, `Model-imatrix`).
+				sidecar = sidecarFromFileToken(suffixMatch[2].toLowerCase());
+				sidecarForm = SidecarForm.SUFFIX;
+				source = suffixMatch[1];
+			} else {
+				const infixMatch = source.match(MODEL_ID.SIDECAR_INFIX_REGEX);
+
+				if (infixMatch) {
+					sidecar = sidecarFromFileToken(infixMatch[2].toLowerCase());
+					sidecarForm = SidecarForm.INFIX;
+					source = `${infixMatch[1]}-${infixMatch[3]}`;
+				}
+			}
+		}
+
+		// Scan dash-separated segments left-to-right for the first quant match.
+		// - For sidecars like `mtp-Q4_0-180MB.gguf` the quant is `Q4_0`.
+		// - For embedded MTP like `Hy3-IQ1_M-mtp.gguf` we have `Hy3-IQ1_M` and `IQ1_M` matches.
+		// - For main files like `Llama-3-8B-Q4_K_M.gguf` we land on the trailing quant.
+		const segments = source.split(MODEL_ID.SEGMENT_SEPARATOR);
+		const quantIdx = segments.findIndex((seg) => MODEL_ID.QUANTIZATION_SEGMENT_REGEX.test(seg));
+		// Unsloth ships draft heads in two layouts: `shared-` files borrow the
+		// embedding/output weights from the target model, others are self-contained.
+		const shared = segments.some((seg) => seg.toLowerCase() === HF_SHARED_DRAFT_TOKEN);
+
+		let quant = quantIdx >= 0 ? segments[quantIdx].toUpperCase() : null;
+
+		// Recombine a `UD-` (Unsloth Dynamic) prefix, e.g. `...-UD-Q4_K_XL.gguf`.
+		// The prefix must be the whole previous segment, matching the server's
+		// `UD-<quant>` custom-quant convention (e.g. not `-mtp-Q4_K_M`).
+		const udPrefixIdx = quantIdx - 1;
+
+		if (quant && quantIdx > 0 && segments[udPrefixIdx].toUpperCase() === HF_UD_QUANT_PREFIX) {
+			quant = `${HF_UD_QUANT_PREFIX}-${quant}`;
+		}
+
+		return { quant, shared, sidecar, sidecarForm };
+	}
+
+	/**
+	 * Filter raw siblings by file extension and sort by size descending.
+	 */
+	static filterByExtension(siblings: HfModelSibling[], ext: string): HfModelSibling[] {
+		return siblings
+			.filter((f) => f.path.toLowerCase().endsWith(ext.toLowerCase()) && (f.size ?? 0) > 0)
+			.sort((a, b) => (b.size ?? 0) - (a.size ?? 0));
+	}
+
+	/**
+	 * Format model downloads count with K/M/B suffix
+	 */
+	static formatDownloads(downloads: number): string {
+		if (downloads >= GIGABYTE) {
+			return `${(downloads / GIGABYTE).toFixed(1)}${GIGA_LABEL}`;
+		}
+
+		if (downloads >= MEGABYTE) {
+			return `${(downloads / MEGABYTE).toFixed(1)}${MEGA_LABEL}`;
+		}
+
+		if (downloads >= KILOBYTE) {
+			return `${(downloads / KILOBYTE).toFixed(1)}${KILO_LABEL}`;
+		}
+
+		return downloads.toString();
+	}
+
+	/**
+	 * Format file size in bytes to human-readable string
+	 */
+	static formatFileSize(bytes: number): string {
+		if (bytes >= GIGABYTE) {
+			return `${(bytes / GIGABYTE).toFixed(1)} ${GIGABYTE_LABEL}`;
+		}
+
+		if (bytes >= MEGABYTE) {
+			return `${(bytes / MEGABYTE).toFixed(1)} ${MEGABYTE_LABEL}`;
+		}
+
+		if (bytes >= KILOBYTE) {
+			return `${(bytes / KILOBYTE).toFixed(1)} ${KILOBYTE_LABEL}`;
+		}
+
+		return `${bytes} ${BYTE_LABEL}`;
+	}
+
+	/**
+	 * Format likes count with K suffix if applicable
+	 */
+	static formatLikes(likes: number): string {
+		if (likes >= KILOBYTE) {
+			return `${(likes / KILOBYTE).toFixed(1)}${KILO_LABEL}`;
+		}
+
+		return likes.toString();
+	}
+
+	/**
+	 * Format timestamp to relative time
+	 */
+	static formatRelativeTime(timestamp: string): string {
+		const date = new Date(timestamp);
+		const now = new Date();
+		const diffMs = now.getTime() - date.getTime();
+		// timestamps can lie in the future (clock skew); clamp so they read as today
+		const diffDays = Math.max(0, Math.floor(diffMs / MS_PER_DAY));
+
+		if (diffDays === 0) return TODAY_LABEL;
+
+		if (diffDays === 1) return YESTERDAY_LABEL;
+
+		if (diffDays < DAYS_PER_WEEK) return `${diffDays} ${DAYS_AGO_LABEL}`;
+
+		if (diffDays < DAYS_PER_MONTH) {
+			return `${Math.floor(diffDays / DAYS_PER_WEEK)} ${WEEKS_AGO_LABEL}`;
+		}
+
+		if (diffDays < DAYS_PER_YEAR) {
+			return `${Math.floor(diffDays / DAYS_PER_MONTH)} ${MONTHS_AGO_LABEL}`;
+		}
+
+		return `${Math.floor(diffDays / DAYS_PER_YEAR)} ${YEARS_AGO_LABEL}`;
+	}
+
+	/**
+	 * Format a min-max size range with a single shared unit and no spaces
+	 * around the dash, e.g. `19.0-28.6 GB`.
+	 */
+	static formatSizeRange(min: number, max: number): string {
+		const unit =
+			max >= GIGABYTE
+				? GIGABYTE_LABEL
+				: max >= MEGABYTE
+					? MEGABYTE_LABEL
+					: max >= KILOBYTE
+						? KILOBYTE_LABEL
+						: BYTE_LABEL;
+		const div =
+			unit === GIGABYTE_LABEL
+				? GIGABYTE
+				: unit === MEGABYTE_LABEL
+					? MEGABYTE
+					: unit === KILOBYTE_LABEL
+						? KILOBYTE
+						: BYTE;
+		const fmt = (n: number) => (div === BYTE ? `${n}` : `${(n / div).toFixed(1)}`);
+
+		return `${fmt(min)}-${fmt(max)} ${unit}`;
+	}
+
+	// Model Details & Files
+
+	/**
+	 * Avatar URL for an author (org or user). 404s when the author does not
+	 * exist, so callers should provide a fallback.
+	 */
+	static getAvatarUrl(author: string): string {
+		return `${HF_AVATARS_URL}${PATH_SEPARATOR}${author}`;
+	}
+
+	/**
+	 * Resolve the original (non-GGUF) base model `{ org, name }` for a GGUF repo
+	 * from its HF card (`cardData.base_model`). Returns null when the card has no
+	 * base model. Results are cached per repo.
+	 */
+	static getBaseModel(repoId: string): Promise<{ org: string; name: string } | null> {
+		const cached = this.baseModelCache.get(repoId);
+
+		if (cached !== undefined) return Promise.resolve(cached);
+
+		const pending = this.baseModelPending.get(repoId);
+
+		if (pending) return pending;
+
+		const promise = (async () => {
+			const details = await this.getDetails(repoId);
+			const base = this.getBaseModels(details)[0];
+
+			if (!base) return null;
+
+			const [org, ...rest] = base.split(PATH_SEPARATOR);
+
+			return { name: rest.join(PATH_SEPARATOR), org };
+		})();
+
+		this.baseModelPending.set(repoId, promise);
+
+		promise
+			.then((result) => this.baseModelCache.set(repoId, result))
+			.finally(() => this.baseModelPending.delete(repoId));
+
+		return promise;
+	}
+
+	/**
+	 * Extract the original (non-GGUF) base model ids for a repo, from
+	 * `cardData.base_model` (string or list) and the `base_model:` tags.
+	 */
+	static getBaseModels(model: HfModelDetailInfo | null): string[] {
+		if (!model) return [];
+
+		const cardBase = model.cardData?.base_model;
+		const fromCard: string[] = Array.isArray(cardBase) ? cardBase : cardBase ? [cardBase] : [];
+		const fromTags = (model.tags ?? [])
+			.map((t) => HF_BASE_MODEL_TAG_REGEX.exec(t)?.[1])
+			.filter((v): v is string => Boolean(v));
+
+		return Array.from(new Set([...fromCard, ...fromTags]));
+	}
+
+	/**
+	 * Look up the average bit-depth for a known GGUF quantization.
+	 * Returns `null` for unrecognized tokens.
+	 */
+	static getBitDepth(quant: string): number | null {
+		// Strip a leading `UD-` (Unsloth Dynamic) prefix before lookup.
+		const base = quant.replace(HF_UD_QUANT_PREFIX_REGEX, '');
+		const direct = HuggingFaceService.QUANT_BIT_DEPTH[base];
+
+		if (direct !== undefined) return direct;
+
+		// Fall back to the leading precision digits for variants missing from the
+		// map, e.g. `Q4_K_XL` -> 4, `IQ2_XXS` -> 2, `TQ1_0` -> 1, `BF16` -> 16.
+		const match = HF_QUANT_PRECISION_REGEX.exec(base);
+
+		return match ? parseInt(match[1], 10) : null;
+	}
+
+	/**
+	 * Get GGUF models by pipeline task
+	 */
+	static async getByTask(
+		pipelineTag: string,
+		params: Omit<HfModelSearchParams, 'pipeline_tag'> = {}
+	): Promise<HfModelInfo[]> {
+		return this.search({
+			...params,
+			pipeline_tag: pipelineTag
+		});
+	}
+
+	/**
+	 * Fetch the llama.app model catalog. Returns an empty array on failure so
+	 * callers can fall back gracefully.
+	 */
+	static async getCatalog(): Promise<HfCatalogEntry[]> {
+		const response = await fetch(MODELS_DISCOVER_CATALOG_URL);
+
+		if (!response.ok) throw new Error(`Failed to fetch catalog: ${response.status}`);
+
+		return (await response.json()) as HfCatalogEntry[];
+	}
+
+	static async getDetails(modelId: string): Promise<HfModelDetailInfo | null> {
+		// Do not encode the modelId, it contains slashes for author/name.
+		// `full=true` includes cardData (description, base_model) and safetensors.
+		const url = `${HF_API_MODELS_URL}${PATH_SEPARATOR}${modelId}?${HF_FULL_DETAIL_PARAM}`;
+
+		try {
+			const response = await fetch(url);
+
+			if (response.status === HF_HTTP_NOT_FOUND) return null;
+
+			if (!response.ok) throw new Error(`Failed to fetch model details: ${response.status}`);
+
+			const data = (await response.json()) as HfModelDetailInfo;
+
+			return data;
+		} catch (error) {
+			console.error(`Error fetching details for ${modelId}:`, error);
+
+			return null;
+		}
+	}
+
+	/**
+	 * Get model URL on Hugging Face Hub
+	 */
+	static getModelUrl(modelId: string): string {
+		return `${HF_BASE_URL}${PATH_SEPARATOR}${modelId}`;
+	}
+
+	// Utility Methods
+
+	/**
+	 * Get most liked GGUF models
+	 */
+	static async getMostLiked(limit: number = HF_DEFAULT_LIMIT): Promise<HfModelInfo[]> {
+		return this.search({ limit, sort: HfModelSort.LIKES });
+	}
+
+	/**
+	 * Get newly released GGUF models
+	 */
+	static async getNew(limit: number = HF_DEFAULT_LIMIT): Promise<HfModelInfo[]> {
+		return this.search({ limit, sort: HfModelSort.CREATED_AT });
+	}
+
+	/**
+	 * Get most popular GGUF models by downloads
+	 */
+	static async getPopular(limit: number = HF_DEFAULT_LIMIT): Promise<HfModelInfo[]> {
+		return this.search({ limit, sort: HfModelSort.DOWNLOADS });
+	}
+
+	/**
+	 * Fetch the raw README.md for a repo, with the YAML frontmatter stripped.
+	 */
+	static async getReadme(modelId: string): Promise<string | null> {
+		// Do not encode the modelId, it contains slashes for author/name
+		const url = `${HF_BASE_URL}${PATH_SEPARATOR}${modelId}${PATH_SEPARATOR}${HF_RAW_PATH}${PATH_SEPARATOR}${HF_MAIN_BRANCH}${PATH_SEPARATOR}${HF_README_FILENAME}`;
+
+		try {
+			const response = await fetch(url);
+
+			if (response.status === HF_HTTP_NOT_FOUND) return null;
+
+			if (!response.ok) throw new Error(`Failed to fetch README: ${response.status}`);
+
+			return HuggingFaceService.stripFrontmatter(await response.text());
+		} catch (error) {
+			console.error(`Error fetching README for ${modelId}:`, error);
+
+			return null;
+		}
+	}
+
+	/**
+	 * Get repository file tree to list available GGUF variants. Recursive so
+	 * repos that keep quants in per-quant subdirectories (e.g. `UD-Q4_K_XL/`)
+	 * are included; follows cursor pagination for repos over one page.
+	 */
+	static async getTree(modelId: string): Promise<HfModelSibling[]> {
+		const files: HfModelSibling[] = [];
+		const firstUrl =
+			`${HF_API_MODELS_URL}${PATH_SEPARATOR}${modelId}${PATH_SEPARATOR}${HF_TREE_PATH}` +
+			`${PATH_SEPARATOR}${HF_MAIN_BRANCH}?${HF_RECURSIVE_TREE_PARAM}`;
+
+		let url: string | null = firstUrl;
+
+		try {
+			for (let page = 0; url && page < HF_TREE_MAX_PAGES; page++) {
+				const response: Response = await fetch(url);
+
+				if (!response.ok) return files;
+
+				const data = (await response.json()) as HfModelSibling[];
+
+				files.push(...data.filter((f) => f.type !== HfEntryType.DIRECTORY));
+
+				url = HuggingFaceService.parseNextPageUrl(response.headers.get(HF_LINK_HEADER));
+			}
+		} catch {
+			// Return whatever was fetched before the failure.
+		}
+
+		return files;
+	}
+
+	/**
+	 * Get trending GGUF models
+	 */
+	static async getTrending(limit: number = HF_DEFAULT_LIMIT): Promise<HfModelInfo[]> {
+		return this.search({ limit, sort: HfModelSort.TRENDING_SCORE });
+	}
+
+	/**
+	 * Parse a local HF cache file path
+	 * (`.../models--<org>--<name>/snapshots/<sha>/<file>`) into its repo id and
+	 * repo-relative file path. Returns null when the path is not an HF cache path.
+	 */
+	static parseCachePath(path: string): { repo: string; file: string } | null {
+		// the paths come from the server's CLI args, which use native separators
+		const match = HF_CACHE_PATH_REGEX.exec(path.replace(/\\/g, PATH_SEPARATOR));
+
+		if (!match) return null;
+
+		const parts = match[1].split(HF_CACHE_DIR_SEPARATOR);
+
+		if (parts.length < 2) return null;
+
+		return {
+			file: match[2],
+			repo: `${parts[0]}${PATH_SEPARATOR}${parts.slice(1).join(HF_CACHE_DIR_SEPARATOR)}`
+		};
+	}
+
+	/**
+	 * Best-effort parameter count parsed from a model id/name, e.g. `27B` from
+	 * `Qwen3.8-27B-GGUF` or `300M` from `embeddinggemma-300M-GGUF`. Returns null
+	 * when no size token is present.
+	 */
+	static parseParamCount(name: string): string | null {
+		const match = HF_PARAM_COUNT_REGEX.exec(name);
+
+		if (!match) return null;
+
+		return `${match[1]}${match[2].toUpperCase()}`;
+	}
+
+	/**
+	 * Parse a human size string (`177GB`, `1.2 TB`, `500MB`) to bytes. Returns
+	 * null when it carries no number or no known suffix, so callers can fall
+	 * back to another source instead of showing a wrong size.
+	 */
+	static parseSizeBytes(size: string): number | null {
+		const match = HF_SIZE_STRING_REGEX.exec(size);
+
+		if (!match) return null;
+
+		const value = parseFloat(match[1]);
+		const multiplier = HF_SIZE_SUFFIX_BYTES[match[2].toLowerCase()];
+
+		if (!Number.isFinite(value) || multiplier === undefined) return null;
+
+		return value * multiplier;
+	}
+
+	/**
+	 * Parse model tags to extract useful information
+	 */
+	static parseTags(tags: string[]): {
+		license: string | null;
+		isGated: boolean;
+		isGguf: boolean;
+		isSafetensors: boolean;
+		tasks: string[];
+	} {
+		const license =
+			tags
+				.find((tag) => tag.startsWith(HF_LICENSE_TAG_PREFIX))
+				?.replace(HF_LICENSE_TAG_PREFIX, '') || null;
+		const isGated = tags.includes(HF_GATED_TAG);
+		const isGguf = tags.includes(HF_GGUF_TAG);
+		const isSafetensors = tags.includes(HF_SAFETENSORS_TAG);
+		const tasks = tags.filter((tag) => HF_TASK_TAGS.includes(tag));
+
+		return { isGated, isGguf, isSafetensors, license, tasks };
+	}
+
+	/**
+	 * Search GGUF models with various filters and options.
+	 *
+	 * Always expands the fields the discover rows render (chat template, context
+	 * length, siblings, ...) so a search result carries the same badges as a
+	 * catalog entry; caller-provided `expand` entries are merged in.
+	 */
+	static async search(params: HfModelSearchParams = {}): Promise<HfModelInfo[]> {
+		const { expand, limit = HF_DEFAULT_LIMIT, ...restParams } = params;
+		const url = this.buildUrl({
+			...restParams,
+			expand: [...new Set([...HF_MODEL_LIST_EXPAND, ...(expand ?? [])])],
+			filter: HF_GGUF_FILTER,
+			limit: Math.min(limit, HF_MAX_LIMIT)
+		});
+
+		return this.fetchWithRetry(url);
+	}
+
+	/**
+	 * Search models by query string
+	 */
+	static async searchByQuery(
+		query: string,
+		params: Omit<HfModelSearchParams, 'search'> = {}
+	): Promise<HfModelInfo[]> {
+		return this.search({
+			...params,
+			search: query
+		});
+	}
+
+	/**
+	 * Build API URL from search parameters
+	 */
+	private static buildUrl(params: HfModelSearchParams): string {
+		const url = new URL(this.BASE_URL);
+
+		Object.entries(params).forEach(([key, value]) => {
+			if (value !== undefined && value !== null && value !== '') {
+				if (Array.isArray(value)) {
+					value.forEach((v) => url.searchParams.append(key, v));
+				} else {
+					url.searchParams.set(key, String(value));
+				}
+			}
+		});
+
+		return url.toString();
+	}
+
+	/**
+	 * Delay helper for retry logic
+	 */
+	private static delay(ms: number): Promise<void> {
+		return new Promise((resolve) => setTimeout(resolve, ms));
+	}
+
+	/**
+	 * Fetch data with retry logic for resilience
+	 */
+	private static async fetchWithRetry(url: string, attempt: number = 1): Promise<HfModelInfo[]> {
+		try {
+			const response = await fetch(url);
+
+			if (!response.ok) {
+				if (response.status === HF_HTTP_NOT_FOUND) {
+					return [];
+				}
+
+				if (response.status >= HF_HTTP_SERVER_ERROR_MIN && attempt < HF_RETRY_ATTEMPTS) {
+					await this.delay(HF_RETRY_DELAY_MS * attempt);
+
+					return this.fetchWithRetry(url, attempt + 1);
+				}
+
+				throw new Error(`API request failed: ${response.status} ${response.statusText}`);
+			}
+
+			const data = await response.json();
+
+			if (Array.isArray(data)) {
+				return data as HfModelInfo[];
+			}
+
+			if (data && Array.isArray(data.data)) {
+				return data.data as HfModelInfo[];
+			}
+
+			throw new Error('Unexpected API response format');
+		} catch (error) {
+			// only transient failures are retried; anything else fails the search
+			const transient =
+				error instanceof TypeError ||
+				(error instanceof Error && error.message.startsWith('API request failed: 5'));
+
+			if (transient && attempt < HF_RETRY_ATTEMPTS) {
+				await this.delay(HF_RETRY_DELAY_MS * attempt);
+
+				return this.fetchWithRetry(url, attempt + 1);
+			}
+
+			throw error;
+		}
+	}
+
+	// Internal Methods
+
+	/** Extract the `rel="next"` URL from an RFC 5988 `Link` header, if present. */
+	private static parseNextPageUrl(linkHeader: string | null): string | null {
+		if (!linkHeader) return null;
+
+		const match = HF_LINK_NEXT_REGEX.exec(linkHeader);
+
+		return match ? match[1] : null;
+	}
+
+	/** Full path of one shard in a split-shard GGUF set. */
+	private static shardPath(stem: string, index: number, total: number): string {
+		const pad = (n: number) => String(n).padStart(HF_SHARD_PAD_WIDTH, '0');
+
+		return `${stem}-${pad(index)}-of-${pad(total)}.gguf`;
+	}
+
+	/** Strip a leading YAML frontmatter block (--- ... ---) from a markdown document. */
+	private static stripFrontmatter(text: string): string {
+		const match = text.match(HF_FRONTMATTER_REGEX);
+
+		return match ? text.slice(match[0].length) : text;
+	}
+}
@@ -136,7 +136,7 @@ export { ConversationTransferService } from './conversation-transfer.service';
 *
 * **Server Mode Behavior:**
 * - **MODEL mode**: Only `list()` is relevant — single model always loaded
- * - **ROUTER mode**: Full lifecycle — `list()`, `listRouter()`, `load()`, `unload()`
+ * - **ROUTER mode**: Full lifecycle — `list()`, `load()`, `unload()`
 *
 * **Endpoints:**
 * - `GET /v1/models` — OpenAI-compatible model list (both modes)
@@ -147,6 +147,16 @@ export { ConversationTransferService } from './conversation-transfer.service';
 */
 export { ModelsService } from './models.service';

+/**
+ * **HuggingFaceService** - Hugging Face Hub browsing and searching
+ *
+ * Stateless HTTP client for the HF REST API (`/api/models`, `/tree`, raw
+ * README) and the llama.app model catalog. Provides GGUF file analysis
+ * (quant metadata, shard collapsing, size formatting) used by the models
+ * discover UI.
+ */
+export { HuggingFaceService } from './huggingface.service';
+
 /**
 * **PropsService** - Server properties and capabilities retrieval
 *
@@ -7,14 +7,16 @@
 */

 import { base } from '$app/paths';
-import { API_MODELS, MODEL_ID } from '$lib/constants';
+import { API_MODELS, MODEL_ID, type ModelSidecar, SIDECAR_TOKENS } from '$lib/constants';
 import { ServerModelStatus } from '$lib/enums';
 import type { ParsedModelId } from '$lib/types/models';
 import {
+	apiDelete,
 	apiFetch,
 	apiPost,
 	extractSseDataPayload,
 	normalizeModelName,
+	sidecarFromFileToken,
 	splitSseRecords
 } from '$lib/utils';
 import { getAuthHeaders } from '$lib/utils/api-headers';
@@ -22,6 +24,70 @@ import { getAuthHeaders } from '$lib/utils/api-headers';
 export class ModelsService {
 	private static readonly SSE_RECONNECT_MS = 1000;

+	/**
+	 * Build the `<repo>:<tag>` string expected by POST /models from a parsed
+	 * filename quant + optional sidecar type. Used by the model download
+	 * dialog so callers don't have to know about the tag conventions.
+	 *
+	 * @param repoId - HuggingFace repo id (e.g. `ggml-org/gemma-3-4b-it-GGUF`)
+	 * @param quant - Quantization token, e.g. `Q4_K_M`
+	 * @param sidecar - Sidecar type, as its lowercase filename token (e.g. `mtp`)
+	 * @returns Repo id possibly suffixed with `:tag`
+	 */
+	static buildDownloadTag(
+		repoId: string,
+		quant: string | null,
+		sidecar: ModelSidecar | null
+	): string {
+		if (!quant && !sidecar) return repoId;
+
+		if (!quant) return `${repoId}:${sidecar}`;
+
+		const tag = sidecar ? `${quant}-${sidecar}` : quant;
+
+		return `${repoId}:${tag}`;
+	}
+
+	/**
+	 * Cancel an in-flight download or remove a previously downloaded/failed
+	 * entry from the server's model cache (ROUTER mode only).
+	 *
+	 * Sends DELETE `/models?model=<hfRepoWithTag>`:
+	 * - while a download is running, the child subprocess is asked to exit
+	 *   and any partial `.tmp` files are removed;
+	 * - once the entry has finished downloading or has failed, the cached
+	 *   files are removed from disk.
+	 *
+	 * @param hfRepoWithTag - HuggingFace repo id in the same `<repo>:<tag>`
+	 *                        format returned by `buildDownloadTag`.
+	 * @returns Server acknowledgement containing the success flag
+	 */
+	static async cancelDownload(hfRepoWithTag: string): Promise<ApiModelsDownloadResponse> {
+		return apiDelete<ApiModelsDownloadResponse>(API_MODELS.DELETE, {
+			model: hfRepoWithTag
+		});
+	}
+
+	/**
+	 * Trigger a model download from HuggingFace (ROUTER mode only).
+	 *
+	 * Sends a POST request to `/models`. The response returns immediately; the
+	 * actual download runs in the background and tracks progress through
+	 * `/models/sse`. The server picks the file that matches the supplied tag
+	 * (when present) and additionally pulls mmproj / draft sidecar weights as
+	 * appropriate for the model.
+	 *
+	 * @param hfRepoWithTag - HuggingFace repo id, optionally suffixed with
+	 *                        `:<tag>` (e.g. `ggml-org/gemma-3-4b-it-GGUF:Q4_K_M`
+	 *                        or `:IQ1_M-mtp` for an embedded-draft GGUF).
+	 * @returns Server acknowledgement containing the success flag
+	 */
+	static async downloadModel(hfRepoWithTag: string): Promise<ApiModelsDownloadResponse> {
+		const payload: ApiModelsDownloadRequest = { model: hfRepoWithTag };
+
+		return apiPost<ApiModelsDownloadResponse>(API_MODELS.DOWNLOAD, payload);
+	}
+
 	/**
 	 * Check if a model is loaded based on its metadata.
 	 *
@@ -32,14 +98,6 @@ export class ModelsService {
 		return model.status.value === ServerModelStatus.LOADED;
 	}

-	/**
-	 *
-	 *
-	 * Load/Unload
-	 *
-	 *
-	 */
-
 	/**
 	 * Check if a model is currently loading.
 	 *
@@ -50,25 +108,39 @@ export class ModelsService {
 		return model.status.value === ServerModelStatus.LOADING;
 	}

+	/**
+	 *
+	 *
+	 * Load/Unload
+	 *
+	 *
+	 */
+
+	/**
+	 * True when a router entry id is a sidecar-only entry, e.g. `org/model:Q4_0-mtp`
+	 * or `org/model:mmproj`. Such entries mark a downloaded sidecar file, not a
+	 * loadable model, so the selector skips them.
+	 */
+	static isSidecarEntry(modelId: string): boolean {
+		const idx = modelId.indexOf(MODEL_ID.QUANTIZATION_SEPARATOR);
+
+		if (idx === MODEL_ID.NOT_FOUND) return false;
+
+		const tag = modelId.slice(idx + 1).toLowerCase();
+		const dash = tag.lastIndexOf(MODEL_ID.SEGMENT_SEPARATOR);
+		const token = dash === -1 ? tag : tag.slice(dash + 1);
+
+		return SIDECAR_TOKENS.includes(token);
+	}
+
 	/**
 	 * Fetch list of models from OpenAI-compatible endpoint.
 	 * Works in both MODEL and ROUTER modes.
 	 *
 	 * @returns List of available models with basic metadata
 	 */
-	static async list(): Promise<ApiModelListResponse> {
-		return apiFetch<ApiModelListResponse>(API_MODELS.LIST);
-	}
-
-	/**
-	 * Fetch list of all models with detailed metadata (ROUTER mode).
-	 * Returns models with load status, paths, and other metadata
-	 * beyond what the OpenAI-compatible endpoint provides.
-	 *
-	 * @returns List of models with detailed status and configuration info
-	 */
-	static async listRouter(): Promise<ApiRouterModelsListResponse> {
-		return apiFetch<ApiRouterModelsListResponse>(API_MODELS.LIST);
+	static async list(): Promise<ApiModelsListResponse> {
+		return apiFetch<ApiModelsListResponse>(API_MODELS.LIST);
 	}

 	/**
@@ -80,14 +152,14 @@ export class ModelsService {
 	 * @param extraArgs - Optional additional arguments to pass to the model instance
 	 * @returns Load response from the server
 	 */
-	static async load(modelId: string, extraArgs?: string[]): Promise<ApiRouterModelsLoadResponse> {
+	static async load(modelId: string, extraArgs?: string[]): Promise<ApiModelsLoadResponse> {
 		const payload: { model: string; extra_args?: string[] } = { model: modelId };

 		if (extraArgs && extraArgs.length > 0) {
 			payload.extra_args = extraArgs;
 		}

-		return apiPost<ApiRouterModelsLoadResponse>(API_MODELS.LOAD, payload);
+		return apiPost<ApiModelsLoadResponse>(API_MODELS.LOAD, payload);
 	}

 	/**
@@ -108,11 +180,46 @@ export class ModelsService {
 			params: null,
 			quantization: null,
 			raw: modelId,
+			sidecar: null,
 			tags: []
 		};
+
 		// strip directory path and weight extension so a bare `-m /path/file.gguf`
 		// parses like a clean repo id; the HF `org/model` form is preserved
-		const source = normalizeModelName(modelId).replace(MODEL_ID.WEIGHT_EXTENSION_RE, '');
+		let source = normalizeModelName(modelId).replace(MODEL_ID.WEIGHT_EXTENSION_REGEX, '');
+
+		// 0. Detect sidecar prefix (mtp-, dflash-, mmproj-) before any other
+		//    splitting so the inner id parses cleanly.
+		const prefixMatch = source.match(MODEL_ID.SIDECAR_PREFIX_REGEX);
+
+		if (prefixMatch) {
+			result.sidecar = sidecarFromFileToken(prefixMatch[1].toLowerCase());
+			source = prefixMatch[2];
+
+			// a sidecar filename's remainder may be just the quant token,
+			// e.g. `mtp-Q4_0.gguf` or `mmproj-F16.gguf`
+			if (MODEL_ID.QUANTIZATION_SEGMENT_REGEX.test(source)) {
+				result.quantization = source.toUpperCase();
+				source = '';
+			}
+		} else {
+			// 0b. Detect `-<type>` suffix (`-mtp`, `-dflash`, `-dspark`, `-eagle3`).
+			//     Only strip it when the segment preceding it looks like a real quant
+			//     token, so a model literally named `MyModel-mtp` is not mistaken for a
+			//     draft one.
+			const suffixMatch = source.match(MODEL_ID.SIDECAR_SUFFIX_REGEX);
+
+			if (suffixMatch) {
+				const candidate = suffixMatch[1];
+				const headSeg = candidate.split(MODEL_ID.SEGMENT_SEPARATOR).pop();
+
+				if (headSeg && MODEL_ID.QUANTIZATION_SEGMENT_REGEX.test(headSeg)) {
+					result.sidecar = sidecarFromFileToken(suffixMatch[2].toLowerCase());
+					source = candidate;
+				}
+			}
+		}
+
 		// 1. Extract colon-separated quantization (e.g. `model:Q4_K_M`)
 		const colonIdx = source.indexOf(MODEL_ID.QUANTIZATION_SEPARATOR);

@@ -143,7 +250,7 @@ export class ModelsService {
 		if (dotIdx !== MODEL_ID.NOT_FOUND && !result.quantization) {
 			const afterDot = modelStr.slice(dotIdx + 1);

-			if (MODEL_ID.QUANTIZATION_SEGMENT_RE.test(afterDot)) {
+			if (MODEL_ID.QUANTIZATION_SEGMENT_REGEX.test(afterDot)) {
 				result.quantization = afterDot;
 				modelStr = modelStr.slice(0, dotIdx);
 			}
@@ -158,8 +265,8 @@ export class ModelsService {
 			const last = segments[segments.length - 1];
 			const secondLast = segments.length > 2 ? segments[segments.length - 2] : null;

-			if (MODEL_ID.QUANTIZATION_SEGMENT_RE.test(last)) {
-				if (secondLast && MODEL_ID.CUSTOM_QUANTIZATION_PREFIX_RE.test(secondLast)) {
+			if (MODEL_ID.QUANTIZATION_SEGMENT_REGEX.test(last)) {
+				if (secondLast && MODEL_ID.CUSTOM_QUANTIZATION_PREFIX_REGEX.test(secondLast)) {
 					result.quantization = `${secondLast}-${last}`;
 					segments.splice(segments.length - 2, 2);
 				} else {
@@ -176,10 +283,10 @@ export class ModelsService {
 		for (let i = 0; i < segments.length; i++) {
 			const seg = segments[i];

-			if (paramsIdx === MODEL_ID.NOT_FOUND && MODEL_ID.PARAMS_RE.test(seg)) {
+			if (paramsIdx === MODEL_ID.NOT_FOUND && MODEL_ID.PARAMS_REGEX.test(seg)) {
 				paramsIdx = i;
 				result.params = seg.toUpperCase();
-			} else if (paramsIdx !== MODEL_ID.NOT_FOUND && MODEL_ID.ACTIVATED_PARAMS_RE.test(seg)) {
+			} else if (paramsIdx !== MODEL_ID.NOT_FOUND && MODEL_ID.ACTIVATED_PARAMS_REGEX.test(seg)) {
 				activatedParamsIdx = i;
 				result.activatedParams = seg.toUpperCase();
 			}
@@ -220,8 +327,8 @@ export class ModelsService {
 	 * @param modelId - Model identifier to unload
 	 * @returns Unload response from the server
 	 */
-	static async unload(modelId: string): Promise<ApiRouterModelsUnloadResponse> {
-		return apiPost<ApiRouterModelsUnloadResponse>(API_MODELS.UNLOAD, { model: modelId });
+	static async unload(modelId: string): Promise<ApiModelsUnloadResponse> {
+		return apiPost<ApiModelsUnloadResponse>(API_MODELS.UNLOAD, { model: modelId });
 	}

 	/**
@@ -40,6 +40,9 @@ export { mcpStore } from './mcp/index.svelte';
 // MODELS
 export { modelsStore } from './models/index.svelte';

+// MODELS DISCOVER (HuggingFace browse state for the discover dialog)
+export { modelsDiscoverStore } from './models-discover/index.svelte';
+
 // SERVER
 export { serverStore } from './server.svelte';

@@ -0,0 +1,275 @@
+/**
+ * modelsDiscoverStore - Models Discover browse state
+ *
+ * Owns the HuggingFace GGUF model list shown in the discover sidebar
+ * (DialogModelsDiscover). By default the list is the curated catalog set;
+ * search replaces it with matches across all of HuggingFace. Both paths fetch
+ * the same fields, so a row renders the same badges and sizes either way.
+ */
+
+import { MODELS_DISCOVER_CATALOG_BATCH } from '$lib/constants';
+import { HuggingFaceService } from '$lib/services';
+import type {
+	HfCatalogBuild,
+	HfCatalogEntry,
+	HfModelInfo,
+	HfModelSibling
+} from '$lib/types/huggingface';
+import { isAuxSidecar } from '$lib/utils';
+import { SvelteMap } from 'svelte/reactivity';
+
+/** Min/max GGUF file size (bytes) across the quants of one repo. */
+export interface ModelsDiscoverSizeRange {
+	max: number;
+	min: number;
+}
+
+class ModelsDiscoverStore {
+	error = $state<string | null>(null);
+	models = $state<HfModelInfo[]>([]);
+	/** First model in the list - discover auto-opens this one. */
+	firstModel = $derived(this.models[0] ?? null);
+
+	loading = $state(false);
+	/** True while a search is in flight; only the newest request owns this flag. */
+	searching = $state(false);
+
+	private catalog: HfCatalogEntry[] = [];
+	/** Repo id -> size range, for catalog rows and lazily measured search rows. */
+	private catalogSizeRanges = new SvelteMap<string, ModelsDiscoverSizeRange>();
+	private defaultModels: HfModelInfo[] = [];
+	private fetched = false;
+	private searchRequestId = 0;
+	/** In-flight `sizeRange()` lookups, keyed by repo id. */
+	private sizeRangePending = new Map<string, Promise<ModelsDiscoverSizeRange | undefined>>();
+
+	/**
+	 * Cached size range for a repo, without measuring: the synchronous part of
+	 * `sizeRange()`, for rendering a row before its measurement resolves.
+	 */
+	cachedSizeRangeFor(modelId: string): ModelsDiscoverSizeRange | undefined {
+		return this.catalogSizeRanges.get(modelId);
+	}
+
+	/**
+	 * Catalog family description for a repo id, or undefined when the repo is
+	 * not part of the catalog (e.g. a search result outside the curated list).
+	 */
+	descriptionFor(modelId: string): string | undefined {
+		return this.catalog.find((entry) =>
+			entry.sizes.some((size) => size.builds.some((build) => build.repo === modelId))
+		)?.description;
+	}
+
+	/**
+	 * Fetch the default list from the llama.app catalog, one repo per catalog
+	 * size in display order (newest family first). Every repo is fetched by ID
+	 * with its file tree, so the rows carry chat-template capabilities, context
+	 * length and a real size range - the same data a search result gets.
+	 * No-op when already loaded or in flight.
+	 */
+	async fetch(): Promise<void> {
+		if (this.loading || this.fetched) return;
+
+		this.loading = true;
+		this.error = null;
+
+		try {
+			const catalog = await HuggingFaceService.getCatalog();
+
+			this.catalog = catalog;
+
+			const builds = this.catalogBuilds();
+
+			this.defaultModels = [];
+
+			// fetch in small batches: not every repo hits the HF API at once,
+			// and the rows land as they arrive instead of all at the end
+			for (let i = 0; i < builds.length; i += MODELS_DISCOVER_CATALOG_BATCH) {
+				const batch = await Promise.all(
+					builds.slice(i, i + MODELS_DISCOVER_CATALOG_BATCH).map(async (build) => {
+						const [info, tree] = await Promise.all([
+							HuggingFaceService.getDetails(build.repo),
+							HuggingFaceService.getTree(build.repo)
+						]);
+
+						return { build, info, tree };
+					})
+				);
+
+				for (const { build, info, tree } of batch) {
+					if (!info) continue;
+
+					this.catalogSizeRanges.set(build.repo, this.sizeRangeFor(build, tree));
+
+					// the catalog repo id, not the HF response `id`, drives selection
+					this.defaultModels.push({ ...info, id: build.repo, modelId: build.repo } as HfModelInfo);
+				}
+
+				// show what has landed, unless a search owns the list right now
+				if (!this.searching) {
+					this.models = [...this.defaultModels];
+				}
+			}
+
+			this.fetched = true;
+		} catch (err) {
+			this.error = err instanceof Error ? err.message : 'Failed to fetch models';
+		} finally {
+			this.loading = false;
+		}
+	}
+
+	/** Replace the list with search results; an empty query restores the default list. */
+	async search(query: string): Promise<void> {
+		const trimmed = query.trim();
+		const requestId = ++this.searchRequestId;
+
+		if (!trimmed) {
+			this.searching = false;
+			this.models = this.defaultModels;
+			this.error = null;
+
+			return;
+		}
+
+		this.searching = true;
+
+		try {
+			const results = await HuggingFaceService.searchByQuery(trimmed, { limit: 50 });
+
+			if (requestId !== this.searchRequestId) return;
+
+			this.models = results;
+			this.error = null;
+		} catch (err) {
+			if (requestId !== this.searchRequestId) return;
+
+			this.error = err instanceof Error ? err.message : 'Search failed';
+		} finally {
+			if (requestId === this.searchRequestId) this.searching = false;
+		}
+	}
+
+	/**
+	 * Size range for a repo not measured yet - a search result, or a catalog
+	 * repo whose tree came back empty. Fetches the file tree once per repo and
+	 * caches it, so remounting a row (scrolling, searching back) is free.
+	 */
+	sizeRange(modelId: string): Promise<ModelsDiscoverSizeRange | undefined> {
+		const cached = this.catalogSizeRanges.get(modelId);
+
+		if (cached) return Promise.resolve(cached);
+
+		const pending = this.sizeRangePending.get(modelId);
+
+		if (pending) return pending;
+
+		const request = (async () => {
+			const tree = await HuggingFaceService.getTree(modelId);
+			const range = this.sizeRangeOfMainQuants(tree);
+
+			if (range) this.catalogSizeRanges.set(modelId, range);
+
+			return range;
+		})()
+			.catch(() => undefined)
+			.finally(() => this.sizeRangePending.delete(modelId));
+
+		this.sizeRangePending.set(modelId, request);
+
+		return request;
+	}
+
+	/**
+	 * Bytes of every quant the catalog lists under this repo, parsed from the
+	 * `size` strings when a build carries no `sizeBytes`. Never empty: an
+	 * unparsable entry contributes the build's own size.
+	 */
+	private buildSizeBytes(build: HfCatalogBuild): number[] {
+		const sizes = this.catalog
+			.flatMap((entry) => entry.sizes)
+			.flatMap((size) => size.builds.filter((b) => b.repo === build.repo))
+			.map((b) => b.sizeBytes ?? HuggingFaceService.parseSizeBytes(b.size))
+			.filter((bytes): bytes is number => Boolean(bytes) && bytes > 0);
+
+		return sizes.length > 0 ? sizes : [build.sizeBytes ?? 0];
+	}
+
+	/**
+	 * One build per catalog size, newest family first (by release date).
+	 * Prefers the official ggml-org repo, falling back to the first build so
+	 * families published only by other orgs (mistralai, unsloth) still show up.
+	 * Returns an empty array when the catalog is empty.
+	 */
+	private catalogBuilds(): HfCatalogBuild[] {
+		return [...this.catalog]
+			.sort((a, b) => b.released.localeCompare(a.released))
+			.flatMap((entry) =>
+				entry.sizes.flatMap((size) => {
+					const build = size.builds.find((b) => b.repo.startsWith('ggml-org/')) ?? size.builds[0];
+
+					return build ? [build] : [];
+				})
+			);
+	}
+
+	/** Byte sizes of every non-sidecar quant file in a tree, shards collapsed. */
+	private quantSizesOf(tree: HfModelSibling[]): number[] {
+		return HuggingFaceService.collapseGgufShards(
+			HuggingFaceService.filterByExtension(tree, '.gguf')
+		)
+			.filter((f) => {
+				const { quant, sidecar } = HuggingFaceService.extractQuantMeta(f.path) ?? {};
+
+				return Boolean(quant) && (sidecar === null || sidecar === undefined);
+			})
+			.map((f) => f.size ?? 0)
+			.filter((size) => size > 0);
+	}
+
+	/**
+	 * Size range of one catalog build: every quant in the repo's file tree, plus
+	 * the draft sidecars those files carry (mtp, dflash, ...) so the downloaded
+	 * model fits within the range. Falls back to the catalog `size` / `sizeBytes`
+	 * strings when the tree yielded nothing (partial fetch, sharded-only repo).
+	 */
+	private sizeRangeFor(build: HfCatalogBuild, tree: HfModelSibling[]): ModelsDiscoverSizeRange {
+		const quantSizes = this.quantSizesOf(tree);
+
+		if (quantSizes.length === 0) {
+			const listed = this.buildSizeBytes(build);
+
+			return { max: Math.max(...listed), min: Math.min(...listed) };
+		}
+
+		const draftSizes = tree
+			.filter((f) => {
+				const sidecar = HuggingFaceService.extractQuantMeta(f.path)?.sidecar;
+
+				return sidecar !== null && sidecar !== undefined && !isAuxSidecar(sidecar);
+			})
+			.map((f) => f.size ?? 0)
+			.filter((size) => size > 0);
+		const extra = draftSizes.length > 0 ? Math.max(...draftSizes) : 0;
+
+		return {
+			max: Math.max(...quantSizes) + extra,
+			min: Math.min(...quantSizes) + (draftSizes.length > 0 ? Math.min(...draftSizes) : 0)
+		};
+	}
+
+	/**
+	 * Size range across the main-model quants of a file tree, draft sidecars
+	 * excluded (they only widen the range when a row advertises them).
+	 */
+	private sizeRangeOfMainQuants(tree: HfModelSibling[]): ModelsDiscoverSizeRange | undefined {
+		const sizes = this.quantSizesOf(tree);
+
+		if (sizes.length === 0) return undefined;
+
+		return { max: Math.max(...sizes), min: Math.min(...sizes) };
+	}
+}
+
+export const modelsDiscoverStore = new ModelsDiscoverStore();
@@ -193,17 +193,20 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 	}

 	/**
-	 * Fetch router models with full metadata (ROUTER mode only).
-	 * No-op in router mode — fetch() already calls listRouter() internally.
+	 * Fetch models with full metadata (ROUTER mode only).
+	 * No-op in MODEL mode - fetch() already calls list() internally.
 	 * Kept for API compatibility (e.g. handleOpenChange dropdown open handler).
 	 */
 	async fetchRouterModels(): Promise<void> {
 		if (!serverStore.isRouterMode) return;

 		try {
-			const response = await ModelsService.listRouter();
+			const response = await ModelsService.list();

 			this.routerModels = response.data;
+			// keep the selector options in sync: a downloaded / deleted model shows
+			// up here too, not only in the router model rows
+			this.models = this.buildModelOptions(response);
 			await this.props.fetchModalitiesForLoadedModels();

 			const visible = this.getVisibleModels();
@@ -358,30 +361,45 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 	 * Both MODEL and ROUTER modes share the same mapping logic;
 	 * they differ only in which endpoint is called.
 	 */
-	private buildModelOptions(
-		response: ApiModelListResponse | ApiRouterModelsListResponse
-	): ModelOption[] {
-		return response.data.map((item: ApiModelDataEntry, index: number) => {
-			const details = response.models?.[index];
-			const rawCapabilities = Array.isArray(details?.capabilities) ? details?.capabilities : [];
-			const displayNameSource =
-				details?.name && details.name.trim().length > 0 ? details.name : item.id;
-			const modelId = details?.model || item.id;
+	private buildModelOptions(response: ApiModelsListResponse): ModelOption[] {
+		const entries: {
+			details?: ApiModelsListResponse['models'][number];
+			item: ApiModelDataEntry;
+		}[] = response.data.map((item: ApiModelDataEntry, index: number) => ({
+			details: response.models?.[index],
+			item
+		}));

-			return {
-				aliases: item.aliases ?? [],
-				capabilities: rawCapabilities.filter((value: unknown): value is string => Boolean(value)),
-				description: details?.description,
-				details: details?.details,
-				id: item.id,
-				meta: item.meta ?? null,
-				modalities: this.props.buildArchitectureModalities(item.architecture),
-				model: modelId,
-				name: this.toDisplayName(displayNameSource),
-				parsedId: ModelsService.parseModelId(modelId),
-				tags: item.tags ?? []
-			};
-		});
+		return (
+			entries
+				// sidecar entries mark downloaded sidecar files, not loadable models
+				.filter(({ item }) => !ModelsService.isSidecarEntry(item.id))
+				// in-flight downloads are not usable models yet; the selector tracks
+				// them in its "Download in progress" section instead
+				.filter(({ item }) => item.status?.value !== ServerModelStatus.DOWNLOADING)
+				.map(({ details, item }) => {
+					const rawCapabilities = Array.isArray(details?.capabilities) ? details?.capabilities : [];
+					const displayNameSource =
+						details?.name && details.name.trim().length > 0 ? details.name : item.id;
+					const modelId = details?.model || item.id;
+
+					return {
+						aliases: item.aliases ?? [],
+						capabilities: rawCapabilities.filter((value: unknown): value is string =>
+							Boolean(value)
+						),
+						description: details?.description,
+						details: details?.details,
+						id: item.id,
+						meta: item.meta ?? null,
+						modalities: this.props.buildArchitectureModalities(item.architecture),
+						model: modelId,
+						name: this.toDisplayName(displayNameSource),
+						parsedId: ModelsService.parseModelId(modelId),
+						tags: item.tags ?? []
+					};
+				})
+		);
 	}

 	/** Fetch models in MODEL mode (single model, standard OpenAI-compatible). */
@@ -390,7 +408,6 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {

 		return this.buildModelOptions(response);
 	}
-
 	/**
 	 * Filter to models visible in the UI (ui !== false).
 	 */
@@ -422,7 +439,7 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 			const router = serverStore.isRouterMode;

 			if (router) {
-				const response = await ModelsService.listRouter();
+				const response = await ModelsService.list();

 				this.routerModels = response.data;
 				this.models = this.buildModelOptions(response);
@@ -7,12 +7,22 @@
 * modelsStore; the host owns the router model rows the feed updates.
 */

-import { ServerModelsSseEventType, ServerModelStatus } from '$lib/enums';
+import {
+	CLI_FLAGS,
+	HF_UD_QUANT_PREFIX_REGEX,
+	MODEL_ID,
+	PATH_SEPARATOR,
+	PAUSED_MODEL_DOWNLOADS_LOCALSTORAGE_KEY
+} from '$lib/constants';
+import { ModelDownloadStopRequest, ServerModelsSseEventType, ServerModelStatus } from '$lib/enums';
+import { HuggingFaceService } from '$lib/services/huggingface.service';
 import { ModelsService } from '$lib/services/models.service';
 import type { ModelPropsManager } from '$lib/stores/models/props.svelte';
 // direct imports between stores, not via the barrel, to avoid circular deps
 import { serverStore } from '$lib/stores/server.svelte';
-import { SvelteMap } from 'svelte/reactivity';
+// explicit type imports: the app.d.ts globals resolve to `any`, so import the real types
+import type { ApiModelsSseDownloadProgressData, ModelDownloadProgress } from '$lib/types';
+import { SvelteMap, SvelteSet } from 'svelte/reactivity';
 import { toast } from 'svelte-sonner';

 /**
@@ -30,9 +40,59 @@ export interface ModelStatusHost {
 	toDisplayName(id: string): string;
 }

+/**
+ * Comparison key of a `<repo>:<tag>` download identifier: uppercased, with the
+ * `UD-` quant prefix stripped. The router derives cached model names from the
+ * actual file, which drops the prefix, so `repo:UD-Q4_K_XL` and `repo:Q4_K_XL`
+ * must compare equal.
+ */
+function downloadIdKey(repoWithTag: string): string {
+	const idx = repoWithTag.indexOf(MODEL_ID.QUANTIZATION_SEPARATOR);
+	const repo = idx === -1 ? repoWithTag : repoWithTag.slice(0, idx);
+	const tag = idx === -1 ? '' : repoWithTag.slice(idx + 1);
+
+	return `${repo.toUpperCase()}:${tag.toUpperCase().replace(HF_UD_QUANT_PREFIX_REGEX, '')}`;
+}
+
 export class ModelStatusManager {
+	/**
+	 * Sidecar files pulled by registered models, as `<repo>/<file>` keys.
+	 * Sidecars are not separate /v1/models entries - the router pulls them as
+	 * sidecars of a main model and records them in its `--model-draft` /
+	 * `--mmproj` args.
+	 */
+	private downloadedSidecars = $derived.by(() => {
+		const result = new SvelteSet<string>();
+
+		for (const m of this.host.routerModels) {
+			const args = m.status?.args;
+
+			if (!args) continue;
+
+			for (let i = 0; i < args.length - 1; i++) {
+				if (
+					args[i] !== CLI_FLAGS.MODEL_DRAFT &&
+					args[i] !== CLI_FLAGS.MODEL_DRAFT_SHORT &&
+					args[i] !== CLI_FLAGS.MMPROJ
+				) {
+					continue;
+				}
+
+				const parsed = HuggingFaceService.parseCachePath(args[i + 1]);
+
+				if (parsed) result.add(`${parsed.repo}${PATH_SEPARATOR}${parsed.file}`);
+			}
+		}
+
+		return result;
+	});
+	private downloadProgress = new SvelteMap<string, ModelDownloadProgress>();
+	/** `<repo>:<tag>` strings whose most recent download attempt failed (download_failed). */
+	private failedDownloads = new SvelteSet<string>();
 	private loadingStates = new SvelteMap<string, boolean>();
 	private loadProgress = new SvelteMap<string, ModelLoadProgress>();
+	/** Paused downloads with their last reported progress, or null when none arrived before the pause. */
+	private pausedDownloads = new SvelteMap<string, ModelDownloadProgress | null>();
 	// /models/sse feed state, the single source of truth for status and load progress
 	private statusAbort: AbortController | null = null;
 	private statusReaderActive = false;
@@ -40,8 +100,128 @@ export class ModelStatusManager {
 		string,
 		{ target: ServerModelStatus; resolve: () => void; reject: (e: Error) => void }
 	>();
+	/** Tags the user asked to stop (pause or cancel); the download_failed the stop triggers is intentional, not a failure. */
+	private stopRequests = new SvelteMap<string, ModelDownloadStopRequest>();

-	constructor(private host: ModelStatusHost) {}
+	/**
+	 * Cancel an in-flight download or remove a previously downloaded/failed model
+	 * from the server cache (ROUTER mode only). The cached row is dropped via the
+	 * feed's model_remove event.
+	 */
+	async cancelDownload(repoWithTag: string): Promise<boolean> {
+		if (!serverStore.isRouterMode) {
+			toast.error('Model downloads are only available in router mode');
+
+			return false;
+		}
+
+		this.subscribe();
+
+		// in-flight: the kill triggers download_failed over the feed; mark it as a
+		// user cancel so it settles silently instead of toasting a failure
+		if (this.downloadProgress.has(repoWithTag)) {
+			this.stopRequests.set(repoWithTag, ModelDownloadStopRequest.CANCEL);
+		}
+
+		// a downloaded model registers under the name the router derived from the
+		// cached file (e.g. the UD- quant prefix is dropped), so resolve the tag to
+		// the registered id before asking the server to remove it
+		const registeredId =
+			this.host.routerModels.find((m) => downloadIdKey(m.id) === downloadIdKey(repoWithTag))?.id ??
+			repoWithTag;
+
+		try {
+			const res = await ModelsService.cancelDownload(registeredId);
+			const ok = res.success === true;
+
+			if (ok) {
+				this.downloadProgress.delete(repoWithTag);
+				this.failedDownloads.delete(repoWithTag);
+				this.deletePausedDownload(repoWithTag);
+			}
+
+			return ok;
+		} catch (error) {
+			toast.error(`Failed to cancel: ${error instanceof Error ? error.message : 'unknown error'}`);
+
+			return false;
+		}
+	}
+
+	/**
+	 * Cancel an in-flight load (ROUTER mode only). The server force-kills a
+	 * LOADING model on unload; the feed reports the settled status, so no
+	 * waiter is registered here.
+	 */
+	async cancelLoad(modelId: string): Promise<void> {
+		if (!serverStore.isRouterMode) return;
+
+		this.subscribe();
+
+		try {
+			await ModelsService.unload(modelId);
+			toast.info(`Load cancelled: ${this.host.toDisplayName(modelId)}`);
+		} catch (error) {
+			toast.error(`Failed to cancel load: ${this.host.toDisplayName(modelId)}`);
+
+			throw error;
+		}
+	}
+
+	constructor(private host: ModelStatusHost) {
+		// the server has no notion of a paused download, so the ids survive in
+		// localStorage; the progress snapshot is stale after a reload and stays null
+		try {
+			const raw = localStorage.getItem(PAUSED_MODEL_DOWNLOADS_LOCALSTORAGE_KEY);
+
+			for (const repoWithTag of JSON.parse(raw ?? '[]') as string[]) {
+				this.pausedDownloads.set(repoWithTag, null);
+			}
+		} catch {
+			// unreadable or corrupt: start without the paused set
+		}
+	}
+
+	/**
+	 * Trigger a model download from HuggingFace via POST /models
+	 * (ggml-org/llama.cpp#23976). The download runs in the background on the
+	 * server; the model appears in the list once the feed reports models_reload.
+	 * Progress is reported by the /models/sse feed; resuming a paused download
+	 * (same tag) continues from the partial files the pause kept on disk.
+	 */
+	async downloadModel(repoWithTag: string): Promise<void> {
+		if (!serverStore.isRouterMode) {
+			toast.error('Model downloads are only available in router mode');
+
+			return;
+		}
+
+		// the feed must be live so the resulting models_reload event refreshes the list
+		this.subscribe();
+
+		// resuming a paused download: drop the paused state, and let the server
+		// discard its stale DOWNLOADED entry (via the list fetch) before re-posting
+		if (this.deletePausedDownload(repoWithTag) || this.stopRequests.delete(repoWithTag)) {
+			await this.host.fetchRouterModels();
+		}
+
+		try {
+			const res = await ModelsService.downloadModel(repoWithTag);
+
+			if (!res.success) {
+				throw new Error(res.error?.message ?? 'Server rejected the download request');
+			}
+
+			// flip the chip to "downloading" right away; the feed refines it with real progress
+			this.downloadProgress.set(repoWithTag, { downloadedBytes: 0, files: {}, totalBytes: 0 });
+
+			toast.success(`Download started: ${this.host.toDisplayName(repoWithTag)}`);
+		} catch (error) {
+			toast.error(`Download failed: ${repoWithTag}`);
+
+			throw error;
+		}
+	}

 	async ensureLoaded(modelId: string): Promise<void> {
 		if (this.host.isModelLoaded(modelId)) return;
@@ -49,6 +229,38 @@ export class ModelStatusManager {
 		await this.load(modelId);
 	}

+	/**
+	 * All tracked downloads (in flight or paused) with their last reported
+	 * progress, for the models selector's "Download in progress" section.
+	 * Paused entries carry their frozen progress snapshot.
+	 */
+	getDownloadEntries(): {
+		isPaused: boolean;
+		progress: ModelDownloadProgress | null;
+		repoWithTag: string;
+	}[] {
+		const inFlight = Array.from(this.downloadProgress, ([repoWithTag, progress]) => ({
+			isPaused: false,
+			progress,
+			repoWithTag
+		}));
+		const paused = Array.from(this.pausedDownloads, ([repoWithTag, progress]) => ({
+			isPaused: true,
+			progress,
+			repoWithTag
+		}));
+
+		return [...inFlight, ...paused];
+	}
+
+	/**
+	 * Current download progress (bytes) for a `<repo>:<tag>` identifier, or null
+	 * when no download is being reported by the /models/sse feed.
+	 */
+	getDownloadProgress(repoWithTag: string): ModelDownloadProgress | null {
+		return this.downloadProgress.get(repoWithTag) ?? null;
+	}
+
 	/**
 	 * Current load progress for a model, or null when not loading.
 	 */
@@ -56,10 +268,57 @@ export class ModelStatusManager {
 		return this.loadProgress.get(modelId) ?? null;
 	}

+	/**
+	 * Last reported progress of a paused download, or null when no progress
+	 * event arrived before the pause.
+	 */
+	getPausedDownloadProgress(repoWithTag: string): ModelDownloadProgress | null {
+		return this.pausedDownloads.get(repoWithTag) ?? null;
+	}
+
+	/** Whether the most recent download attempt for the given entry failed. */
+	hasFailedDownload(repoWithTag: string): boolean {
+		return this.failedDownloads.has(repoWithTag);
+	}
+
+	/**
+	 * True when the feed reports an active download for the given `<repo>:<tag>`.
+	 * Cleared on download_finished / download_failed.
+	 */
+	isDownloadInProgress(repoWithTag: string): boolean {
+		return this.downloadProgress.has(repoWithTag);
+	}
+
+	/**
+	 * True when the user paused an in-flight download and it has not been resumed.
+	 */
+	isDownloadPaused(repoWithTag: string): boolean {
+		return this.pausedDownloads.has(repoWithTag);
+	}
+
+	/**
+	 * True when the given `<repo>:<tag>` is already a fully downloaded model
+	 * registered with the server (i.e. it shows up in the /v1/models list).
+	 * Both ids are normalized, see downloadIdKey().
+	 */
+	isModelDownloaded(repoWithTag: string): boolean {
+		const key = downloadIdKey(repoWithTag);
+
+		return this.host.routerModels.some((m) => downloadIdKey(m.id) === key);
+	}
+
 	isOperationInProgress(modelId: string): boolean {
 		return this.loadingStates.get(modelId) ?? false;
 	}

+	/**
+	 * True when the given sidecar file (repo-relative path) has been pulled as
+	 * the `--model-draft` or `--mmproj` of some registered model.
+	 */
+	isSidecarDownloaded(repoId: string, filePath: string): boolean {
+		return this.downloadedSidecars.has(`${repoId}/${filePath}`);
+	}
+
 	async load(modelId: string): Promise<void> {
 		if (this.host.isModelLoaded(modelId)) return;

@@ -90,6 +349,31 @@ export class ModelStatusManager {
 		}
 	}

+	/**
+	 * Pause an in-flight download (ROUTER mode only). The server stops the
+	 * download child but keeps the partial files on disk, so re-posting the
+	 * tag (downloadModel) resumes the download where it stopped. The feed
+	 * reports the stop as download_failed; a 'pause' stop request marks it as such.
+	 */
+	async pauseDownload(repoWithTag: string): Promise<void> {
+		if (!serverStore.isRouterMode) {
+			toast.error('Model downloads are only available in router mode');
+
+			return;
+		}
+
+		this.subscribe();
+
+		this.stopRequests.set(repoWithTag, ModelDownloadStopRequest.PAUSE);
+
+		try {
+			await ModelsService.unload(repoWithTag);
+		} catch {
+			this.stopRequests.delete(repoWithTag);
+			toast.error(`Failed to pause: ${repoWithTag}`);
+		}
+	}
+
 	/**
 	 * Open the /models/sse feed and keep it live with auto reconnect.
 	 * Idempotent and router mode only.
@@ -141,6 +425,87 @@ export class ModelStatusManager {
 		this.statusAbort?.abort();
 		this.statusAbort = null;
 		this.loadProgress.clear();
+		this.downloadProgress.clear();
+		this.failedDownloads.clear();
+		this.stopRequests.clear();
+	}
+
+	/**
+	 * Drop the stored progress for the model and toast the outcome.
+	 * A user pause keeps the last progress and stays resumable, a user cancel
+	 * settles silently; genuine failures are marked so the UI can offer a
+	 * delete-and-retry path.
+	 */
+	private applyDownloadFinished(event: ApiModelsSseEvent): void {
+		let request: ModelDownloadStopRequest | undefined;
+
+		if (event.event === ServerModelsSseEventType.DOWNLOAD_FAILED) {
+			request = this.stopRequests.get(event.model);
+			this.stopRequests.delete(event.model);
+		}
+
+		const progress = this.downloadProgress.get(event.model) ?? null;
+
+		this.downloadProgress.delete(event.model);
+
+		if (request === ModelDownloadStopRequest.CANCEL) {
+			// user cancel: settle silently, the feed's model_remove cleans up the entry
+			this.failedDownloads.delete(event.model);
+			this.deletePausedDownload(event.model);
+
+			return;
+		}
+
+		if (request === ModelDownloadStopRequest.PAUSE) {
+			this.setPausedDownload(event.model, progress);
+			this.failedDownloads.delete(event.model);
+
+			return;
+		}
+
+		this.deletePausedDownload(event.model);
+
+		const ok = event.event === ServerModelsSseEventType.DOWNLOAD_FINISHED;
+
+		if (ok) {
+			this.failedDownloads.delete(event.model);
+
+			// the finished download only registers in /v1/models on the next list
+			// fetch (the server reloads its model table then), so refetch to flip
+			// the quant chips to "downloaded" without waiting for a dialog reopen
+			void this.host.fetchRouterModels();
+
+			toast.success(`Download finished: ${this.host.toDisplayName(event.model)}`);
+		} else {
+			this.failedDownloads.add(event.model);
+			toast.error(`Download failed: ${this.host.toDisplayName(event.model)}`);
+		}
+	}
+
+	/**
+	 * Bucket the per-file byte counts from a `download_progress` envelope.
+	 * Total = sum of `total` across files (plan size), downloaded sum of `done`.
+	 */
+	private applyDownloadProgress(event: ApiModelsSseEvent): void {
+		const data = event.data;
+
+		if (!data || !('progress' in data)) return;
+
+		const progress = (data as ApiModelsSseDownloadProgressData).progress;
+
+		let downloaded = 0;
+		let total = 0;
+
+		for (const file of Object.values(progress)) {
+			downloaded += file?.done ?? 0;
+			total += file?.total ?? 0;
+		}
+
+		this.downloadProgress.set(event.model, {
+			downloadedBytes: downloaded,
+			files: progress,
+			totalBytes: total
+		});
 	}

 	/**
@@ -151,7 +516,7 @@ export class ModelStatusManager {
 		const model = event.model;
 		const data = event.data;

-		if (!model || !data?.status) return;
+		if (!model || !data || !('status' in data) || !data.status) return;

 		const status = data.status;

@@ -202,7 +567,33 @@ export class ModelStatusManager {

 				break;
 			case ServerModelsSseEventType.DOWNLOAD_PROGRESS:
+				this.applyDownloadProgress(event);
+
 				break;
+			case ServerModelsSseEventType.DOWNLOAD_FINISHED:
+			case ServerModelsSseEventType.DOWNLOAD_FAILED:
+				this.applyDownloadFinished(event);
+
+				break;
+		}
+	}
+
+	private deletePausedDownload(repoWithTag: string): boolean {
+		if (!this.pausedDownloads.delete(repoWithTag)) return false;
+
+		this.persistPausedDownloads();
+
+		return true;
+	}
+
+	private persistPausedDownloads(): void {
+		try {
+			localStorage.setItem(
+				PAUSED_MODEL_DOWNLOADS_LOCALSTORAGE_KEY,
+				JSON.stringify(Array.from(this.pausedDownloads.keys()))
+			);
+		} catch {
+			// storage unavailable: the pauses just do not survive a reload
 		}
 	}

@@ -226,7 +617,15 @@ export class ModelStatusManager {

 		this.host.routerModels = this.host.routerModels.filter((m) => m.id !== modelId);
 		this.loadProgress.delete(modelId);
+		this.downloadProgress.delete(modelId);
+		this.failedDownloads.delete(modelId);
+		this.deletePausedDownload(modelId);
+		this.stopRequests.delete(modelId);
 		this.rejectStatus(modelId, new Error(`Model removed: ${this.host.toDisplayName(modelId)}`));
+
+		// drop the row from the selector options too; they rebuild from the list
+		// response, which only a refetch provides
+		void this.host.fetchRouterModels();
 	}

 	/**
@@ -236,6 +635,11 @@ export class ModelStatusManager {
 		await ModelsService.watchModelEvents(signal, (event) => this.applyStatusEvent(event));
 	}

+	private setPausedDownload(repoWithTag: string, progress: ModelDownloadProgress | null): void {
+		this.pausedDownloads.set(repoWithTag, progress);
+		this.persistPausedDownloads();
+	}
+
 	/**
 	 * Update one model row status in place, reassigning to trigger reactivity.
 	 */
@@ -138,6 +138,14 @@ export interface ApiModelsSseData {
 	exit_code?: number;
 }

+/**
+ * Per-file size snapshot reported by the download_progress SSE envelope.
+ * Keys are file URLs, values are byte counters (done <= total).
+ */
+export interface ApiModelsSseDownloadProgressData {
+	progress: Record<string, { done: number; total: number }>;
+}
+
 /**
 * Event kind multiplexed on the /models/sse feed.
 * Only the status_* events carry a status payload, models_reload signals a
@@ -150,7 +158,26 @@ export interface ApiModelsSseData {
 export interface ApiModelsSseEvent {
 	model: string;
 	event: ServerModelsSseEventType;
-	data: ApiModelsSseData;
+	data?: ApiModelsSseData | ApiModelsSseDownloadProgressData;
+}
+
+/**
+ * Request body for POST /models (model download).
+ * `model` is a HuggingFace repo id, optionally suffixed with `:<tag>` to
+ * pin a quantization or sidecar file (e.g. `ggml-org/gemma-3-4b-it-GGUF:Q4_K_M`).
+ */
+export interface ApiModelsDownloadRequest {
+	model: string;
+}
+
+/**
+ * Response from POST /models and DELETE /models. The POST endpoint returns
+ * immediately; the download itself runs in the background and emits events
+ * on /models/sse.
+ */
+export interface ApiModelsDownloadResponse {
+	success: boolean;
+	error?: { code: number; message: string; type: string };
 }

 export interface ApiModelDetails {
@@ -174,12 +201,6 @@ export interface ApiModelDetails {
 	};
 }

-export interface ApiModelListResponse {
-	object: string;
-	data: ApiModelDataEntry[];
-	models?: ApiModelDetails[];
-}
-
 export interface ApiLlamaCppServerProps {
 	default_generation_settings: {
 		id: number;
@@ -448,79 +469,27 @@ export interface ApiProcessingState {
 }

 /**
- * Router model metadata - extended from ApiModelDataEntry with additional router-specific fields
- * @deprecated Use ApiModelDataEntry instead - the /models endpoint returns this structure directly
+ * Response from POST /models/load
 */
-export interface ApiRouterModelMeta {
-	/** Model identifier (e.g., "ggml-org/Qwen2.5-Omni-7B-GGUF:latest") */
-	name: string;
-	/** Path to model file or manifest */
-	path: string;
-	/** Optional path to multimodal projector */
-	path_mmproj?: string;
-	/** Whether model is in HuggingFace cache */
-	in_cache: boolean;
-	/** Port where model instance is running (0 if not loaded) */
-	port?: number;
-	/** Current status of the model */
-	status: ApiModelStatus;
-	/** Error message if status is FAILED */
-	error?: string;
-}
-
-/**
- * Request to load a model
- */
-export interface ApiRouterModelsLoadRequest {
-	model: string;
-}
-
-/**
- * Response from loading a model
- */
-export interface ApiRouterModelsLoadResponse {
+export interface ApiModelsLoadResponse {
 	success: boolean;
 	error?: string;
 }

 /**
- * Request to check model status
+ * Response with list of all models from /v1/models and /models endpoints
+ * (same structure regardless of server mode)
 */
-export interface ApiRouterModelsStatusRequest {
-	model: string;
-}
-
-/**
- * Response with model status
- */
-export interface ApiRouterModelsStatusResponse {
-	model: string;
-	status: ModelStatus;
-	port?: number;
-	error?: string;
-}
-
-/**
- * Response with list of all models from /models endpoint
- * Note: This is the same as ApiModelListResponse - the endpoint returns the same structure
- * regardless of server mode (MODEL or ROUTER)
- */
-export interface ApiRouterModelsListResponse {
+export interface ApiModelsListResponse {
 	object: string;
 	data: ApiModelDataEntry[];
+	models?: ApiModelDetails[];
 }

 /**
- * Request to unload a model
+ * Response from POST /models/unload
 */
-export interface ApiRouterModelsUnloadRequest {
-	model: string;
-}
-
-/**
- * Response from unloading a model
- */
-export interface ApiRouterModelsUnloadResponse {
+export interface ApiModelsUnloadResponse {
 	success: boolean;
 	error?: string;
 }
@@ -0,0 +1,227 @@
+/**
+ * HuggingFace Hub Model Browsing Types
+ *
+ * Types for the HuggingFace REST API (/api/models)
+ * Reference: https://huggingface.co/docs/huggingface_hub/package_reference/hf_api
+ */
+
+// Search Options
+
+export interface HfModelSearchParams {
+	/** Full-text search query */
+	search?: string;
+	/** Filter by pipeline task (e.g., "text-generation", "image-generation") */
+	pipeline_tag?: string;
+	/** Filter by library (e.g., "transformers", "diffusers", "gguf") */
+	library_name?: string;
+	/** Filter by tag (e.g., "gguf") */
+	filter?: string;
+	/** Filter by author or organization */
+	author?: string;
+	/** Sort field */
+	sort?: HfModelSort;
+	/** Results per page (1-100) */
+	limit?: number;
+	/** Pagination offset */
+	offset?: number;
+	/** Filter by model config */
+	config?: string;
+	/** Return full model info */
+	full?: boolean;
+	/**
+	 * Fields to include beyond the default set (repeated as `expand=<field>`).
+	 * The list endpoint returns only `_id`, `id`, `modelId` and the sort field
+	 * unless this is given, so callers rendering badges must ask for them.
+	 */
+	expand?: string[];
+	/** Filter by visibility */
+	private?: boolean;
+	/** Filter by gated status */
+	gated?: boolean;
+}
+
+import type { HfEntryType, HfModelSort } from '$lib/enums';
+
+// Model Info (from /api/models)
+
+export interface HfModelInfo {
+	/** Unique document ID */
+	_id: string;
+	/** Model ID (e.g., "meta-llama/Llama-3.1-8B-Instruct") */
+	id: string;
+	/** Number of likes */
+	likes: number;
+	/** Trending score */
+	trendingScore: number;
+	/** Whether the model is private */
+	private: boolean;
+	/** Number of downloads */
+	downloads: number;
+	/** Model tags */
+	tags: string[];
+	/** Pipeline task (e.g., "text-generation") */
+	pipeline_tag: string | null;
+	/** Library name (e.g., "transformers", "diffusers") */
+	library_name: string | null;
+	/** Creation timestamp */
+	createdAt: string;
+	/** Model ID (alias for id) */
+	modelId: string;
+	/** Author / organization (present when full=true) */
+	author?: string;
+	/** Last modified timestamp (present when full=true) */
+	lastModified?: string;
+	/** Repository file listing (present when full=true) */
+	siblings?: HfModelSiblingRef[];
+	/** GGUF metadata (context length, architecture, etc.) */
+	gguf?: HfModelGguf;
+}
+
+// Model Details (with full=true)
+
+export interface HfModelCardData {
+	/** License identifier */
+	license?: string;
+	/** License URL */
+	license_link?: string;
+	/** Model description */
+	description?: string;
+	/** Model library */
+	language?: string[];
+	/** Tags */
+	tags?: string[];
+	/** Original (non-GGUF) model(s) this repo was converted from, e.g. `Qwen/Qwen3.8-27B`. The API returns a single string or a list. */
+	base_model?: string | string[];
+	/** Org that produced the quant, e.g. `bartowski` */
+	quantized_by?: string;
+	[key: string]: unknown;
+}
+
+/** GGUF metadata returned by /api/models/{id}?full=true for GGUF repos. */
+export interface HfModelGguf {
+	/** Total parameter count */
+	total?: number;
+	/** Architecture, e.g. `gemma3`, `qwen3` */
+	architecture?: string;
+	/** Context length */
+	context_length?: number;
+	/** Chat template (Jinja) */
+	chat_template?: string;
+	bos_token?: string;
+	eos_token?: string;
+	/** Total size of all GGUF files in the repo, in bytes */
+	totalFileSize?: number;
+}
+
+export interface HfModelDetails {
+	/** Model ID */
+	id?: string;
+	/** SHA256 digest */
+	sha?: string;
+	/** Last modified timestamp */
+	lastModified?: string;
+	/** Downloads count */
+	downloads?: number;
+	/** Number of likes */
+	likes?: number;
+	/** Whether the model is gated */
+	gated?: boolean;
+	/** Model card data */
+	cardData?: HfModelCardData;
+	/** Tags */
+	tags?: string[];
+	/** Pipeline tag */
+	pipeline_tag?: string | null;
+	/** Library name */
+	library_name?: string | null;
+	/** Safe tensors info */
+	safetensors?: Record<string, unknown>;
+	/** Model size in bytes */
+	size?: number;
+	[key: string]: unknown;
+}
+
+export interface HfModelDetailInfo extends HfModelInfo {
+	/** Whether the model is gated (true/false/'auto') */
+	gated?: boolean | string;
+	/** Repository file listing mirrors of /api/models/{id}/tree/main */
+	siblings?: HfModelSiblingRef[];
+	/** Author / organization */
+	author?: string;
+	/** Last modified timestamp */
+	lastModified?: string;
+	/** Model card YAML data (only present when full=true) */
+	cardData?: HfModelCardData;
+	/** GGUF metadata (only present when full=true for GGUF repos) */
+	gguf?: HfModelGguf;
+	/** Model config (only present when full=true) */
+	config?: Record<string, unknown>;
+	/** Total repo storage in bytes (only present when full=true) */
+	usedStorage?: number;
+	/** Sample widget prompts */
+	widgetData?: Array<{ text?: string }>;
+	/** Related spaces */
+	spaces?: string[];
+}
+
+/** A single entry in a model repository's file tree (`/tree` responses) */
+export interface HfModelSibling {
+	/** Relative path of the file or directory within the repo */
+	path: string;
+	/** Size in bytes (omitted for directories) */
+	size?: number;
+	/** Whether this entry is a directory */
+	type?: HfEntryType;
+	/** OID/hash for the blob */
+	oid?: string;
+	[key: string]: unknown;
+}
+
+/**
+ * A single file entry in a model's `siblings` list. List (`/api/models`) and
+ * detail (`/api/models/{id}`) responses use `rfilename`, unlike `/tree`.
+ */
+export interface HfModelSiblingRef {
+	/** Relative file name within the repo */
+	rfilename: string;
+	[key: string]: unknown;
+}
+
+// API Response
+
+export interface HfModelApiResponse {
+	/** List of models */
+	data: HfModelInfo[];
+	/** Total count (if available) */
+	total?: number;
+}
+
+// llama.app model catalog (https://llama.app/v1/catalog.json)
+
+/** A single GGUF build/repo within a catalog size. */
+export interface HfCatalogBuild {
+	quant: string;
+	size: string;
+	sizeBytes: number;
+	repo: string;
+}
+
+/** A size variant (e.g. `GPT-OSS 20B`) within a catalog entry. */
+export interface HfCatalogSize {
+	name: string;
+	params: string;
+	builds: HfCatalogBuild[];
+}
+
+/** A single model family in the catalog. `featured` marks the staff picks. */
+export interface HfCatalogEntry {
+	name: string;
+	brand: string;
+	description: string;
+	details: string;
+	released: string;
+	license: string;
+	featured?: boolean;
+	maxMemGb?: number;
+	sizes: HfCatalogSize[];
+}
@@ -14,9 +14,11 @@ export type {
 	ApiModelLoadStage,
 	ApiModelsSseProgress,
 	ApiModelsSseData,
+	ApiModelsSseDownloadProgressData,
 	ApiModelsSseEvent,
+	ApiModelsDownloadRequest,
+	ApiModelsDownloadResponse,
 	ApiModelDetails,
-	ApiModelListResponse,
 	ApiLlamaCppServerProps,
 	ApiChatCompletionRequest,
 	ApiChatCompletionToolCallFunctionDelta,
@@ -26,18 +28,29 @@ export type {
 	ApiChatCompletionResponse,
 	ApiSlotData,
 	ApiProcessingState,
-	ApiRouterModelMeta,
-	ApiRouterModelsLoadRequest,
-	ApiRouterModelsLoadResponse,
-	ApiRouterModelsStatusRequest,
-	ApiRouterModelsStatusResponse,
-	ApiRouterModelsListResponse,
-	ApiRouterModelsUnloadRequest,
-	ApiRouterModelsUnloadResponse,
+	ApiModelsLoadResponse,
+	ApiModelsListResponse,
+	ApiModelsUnloadResponse,
 	AudioInputFormat,
 	ApiStreamSession
 } from './api';

+// HuggingFace types
+export type {
+	HfCatalogBuild,
+	HfCatalogEntry,
+	HfCatalogSize,
+	HfModelApiResponse,
+	HfModelCardData,
+	HfModelDetails,
+	HfModelDetailInfo,
+	HfModelGguf,
+	HfModelInfo,
+	HfModelSearchParams,
+	HfModelSibling,
+	HfModelSiblingRef
+} from './huggingface';
+
 // Chat types
 export type {
 	AttachmentMenuItem,
@@ -92,10 +105,20 @@ export type {
 	ModelCapabilities,
 	ModelModalities,
 	ModelOption,
+	ModelDownloadFileProgress,
+	ModelDownloadProgress,
 	ModelLoadProgress,
 	ModalityCapabilities
 } from './models';

+// Models discover types
+export type {
+	ModelBitDepthRow,
+	ModelDownloadEntryState,
+	ModelQuantOption,
+	ModelSelectableFile
+} from './models-discover';
+
 // Settings types
 export type {
 	SettingsConfigValue,
@@ -0,0 +1,30 @@
+import type { ModelSelectableFileKind } from '$lib/enums';
+import type { HfModelSibling } from '$lib/types/huggingface';
+
+/**
+ * Option of a quant `<select>` in the download-options command builder; the
+ * picks only compose the serve command and are not bound to the quant chips.
+ */
+export interface ModelQuantOption {
+	/** Quant token, or the file name when the file carries no quant (e.g. BF16). */
+	label: string;
+	/** Repo-relative file path the option stands for. */
+	path: string;
+}
+
+/** Download state of a single repo entry, injected by the integration layer. */
+export interface ModelDownloadEntryState {
+	/** Server identifier the entry's actions (pause / resume / cancel / retry) target. */
+	repoWithTag: string;
+	isDownloading: boolean;
+	progress: ModelDownloadProgress | null;
+	isDownloaded: boolean;
+	isPaused: boolean;
+	isFailed: boolean;
+}
+
+/** A group of GGUF files sharing one bit-depth bucket. */
+export type ModelBitDepthRow = { bitDepth: number; files: HfModelSibling[] };
+
+/** A selectable GGUF, tagged with its role: main weights, draft, or aux (mmproj). */
+export type ModelSelectableFile = HfModelSibling & { kind: ModelSelectableFileKind };
@@ -1,3 +1,4 @@
+import type { ModelSidecar } from '$lib/constants/model-id.constants';
 import type { ApiModelDataEntry, ApiModelDetails, ApiModelLoadStage } from '$lib/types/api';

 export interface ModelModalities {
@@ -8,6 +9,7 @@ export interface ModelModalities {

 export interface ModelCapabilities {
 	reasoning: boolean;
+	tools: boolean;
 }

 export interface ModelOption {
@@ -24,17 +26,27 @@ export interface ModelOption {
 	tags?: string[];
 }

-/**
- * Ephemeral UI-only load progress for one model instance.
- * Lives only while a load runs, driven by the /models/sse feed.
- * stage is absent until the feed reports its first stage.
- */
+/** UI-only load progress for one model, driven by the /models/sse feed. */
 export interface ModelLoadProgress {
 	stages: ApiModelLoadStage[];
 	current: ApiModelLoadStage;
 	value: number;
 }

+/** Per-file bytes of an in-flight download. */
+export interface ModelDownloadFileProgress {
+	done: number;
+	total: number;
+}
+
+/** Progress of an in-flight download, summed across its files. */
+export interface ModelDownloadProgress {
+	downloadedBytes: number;
+	totalBytes: number;
+	/** Per-file progress keyed by file URL. */
+	files: Record<string, ModelDownloadFileProgress>;
+}
+
 export interface ParsedModelId {
 	raw: string;
 	orgName: string | null;
@@ -42,12 +54,11 @@ export interface ParsedModelId {
 	params: string | null;
 	activatedParams: string | null;
 	quantization: string | null;
+	sidecar: ModelSidecar | null;
 	tags: string[];
 }

-/**
- * Modality capabilities for file validation
- */
+/** Modality capabilities for file validation. */
 export interface ModalityCapabilities {
 	hasVision: boolean;
 	hasAudio: boolean;
@@ -49,7 +49,7 @@ export interface ApiFetchOptions extends Omit<RequestInit, 'headers'> {
 * @example
 * ```typescript
 * // GET request
- * const models = await apiFetch<ApiModelListResponse>('/v1/models');
+ * const models = await apiFetch<ApiModelsListResponse>('/v1/models');
 *
 * // POST request
 * const result = await apiFetch<ApiResponse>('/models/load', {
@@ -137,6 +137,40 @@ export async function apiPost<T, B = unknown>(
 	});
 }

+/**
+ * Send a DELETE request to an API endpoint, optionally with query parameters.
+ *
+ * @param path - API path (query string is appended if `params` is provided)
+ * @param params - Optional record of query parameters
+ * @param options - Additional fetch options
+ * @returns Parsed JSON response
+ */
+export async function apiDelete<T>(
+	path: string,
+	params?: Record<string, string>,
+	options: ApiFetchOptions = {}
+): Promise<T> {
+	// the query is appended to the path so `apiFetch` applies its base-path prefix;
+	// `apiFetchWithParams` resolves an absolute URL and would bypass it
+	let query = '';
+
+	if (params) {
+		const search = new URLSearchParams();
+
+		for (const [key, value] of Object.entries(params)) {
+			if (value !== undefined && value !== null) {
+				search.set(key, value);
+			}
+		}
+
+		const qs = search.toString();
+
+		if (qs) query = `?${qs}`;
+	}
+
+	return apiFetch<T>(`${path}${query}`, { ...options, method: 'DELETE' });
+}
+
 /**
 * Parse error message from a failed response.
 * Tries to extract error message from JSON body, falls back to status text.
@@ -16,7 +16,7 @@ import {
 } from '$lib/constants';
 import type { ToolExecutionResult } from '$lib/types';

-function detectOs(userAgent: string): string {
+export function detectOs(userAgent: string): string {
 	for (const [pattern, os] of BROWSER_INFO_OS_UA_PATTERNS) {
 		if (pattern.test(userAgent)) return os;
 	}
@@ -0,0 +1,29 @@
+/**
+ * Detects whether a model's chat template supports tool calling.
+ *
+ * There is no server flag for tool support, so we infer it from the chat
+ * template. A template that accepts a `tools` array or emits tool-call tokens
+ * is treated as tool-capable.
+ */
+
+/** Tool-call tokens emitted by the template for assistant tool calls, matched case-insensitively. */
+const TOOL_CALL_TOKENS = [
+	'tool_call',
+	'tool_calls',
+	'function_call',
+	'tool_use',
+	'<tool',
+	'<|tool'
+];
+/** Jinja reference to the `tools` array passed in by the caller. */
+const JINJA_TOOLS_VAR = /\{\{[^{}]*\btools\b[^{}]*\}\}|\{%[^{}]*\btools\b[^{}]*%\}/i;
+
+export function detectToolUseSupport(t: string): boolean {
+	if (!t) return false;
+
+	if (JINJA_TOOLS_VAR.test(t)) return true;
+
+	const template = t.toLowerCase();
+
+	return TOOL_CALL_TOKENS.some((token) => template.includes(token));
+}
@@ -9,7 +9,7 @@

 // API utilities
 export { getAuthHeaders, getJsonHeaders, sanitizeHeaders } from './api-headers';
-export { ApiError, apiFetch, apiFetchWithParams, apiPost } from './api-fetch';
+export { ApiError, apiDelete, apiFetch, apiFetchWithParams, apiPost } from './api-fetch';
 export { validateApiKey } from './api-key-validation';

 // Attachment utilities
@@ -107,6 +107,9 @@ export {
 // Model name utilities
 export { normalizeModelName, isValidModelName } from './model-names';

+// Sidecar token utilities
+export { isAuxSidecar, isDraftSidecar, sidecarFromFileToken } from './sidecars';
+
 // Portal utilities
 export { portalToBody } from './portal-to-body';

@@ -340,7 +343,13 @@ export { buildSandboxToolDefinition, SANDBOX_TOOL_DEFINITION } from './sandbox-t
 export { executeGetDatetimeTool } from './get-datetime';

 // Browser fallback for the server's get_info tool
-export { executeBrowserInfoTool } from './browser-info';
+export { detectOs, executeBrowserInfoTool } from './browser-info';
+
+// Tool-use support detection from a chat template
+export { detectToolUseSupport } from './chat-template-tool-detector';
+
+// Model memory estimation
+export { minMemoryTierGb } from './model-compatibility';

 // Cryptography utilities

@@ -0,0 +1,37 @@
+/**
+ * Model memory estimation.
+ *
+ * Mirrors the app's compatibility check (Model+Compatibility.swift): the
+ * runtime budget is RAM x 0.75 minus a fixed overhead, and a file fits when
+ * its size with headroom stays under that budget. The result is the smallest
+ * memory tier that can run the model, so the UI presents an honest machine
+ * requirement instead of a raw file size. Context length and
+ * device-specific budgets are deliberately ignored - callers present the
+ * requirement and let the user judge.
+ */
+import {
+	MB_PER_GB,
+	MEM_TIERS,
+	MIB_BYTES,
+	QUANT_WEIGHT,
+	RAM_BUDGET_RATIO,
+	RAM_OVERHEAD_MB
+} from '$lib/constants';
+
+/**
+ * Smallest memory tier (GB) that can run a model of the given file size,
+ * or null if nothing fits even the largest tier.
+ */
+export function minMemoryTierGb(sizeBytes: number): number | null {
+	if (!sizeBytes) return null;
+
+	const weightMb = (sizeBytes / MIB_BYTES) * QUANT_WEIGHT;
+
+	for (const tier of MEM_TIERS) {
+		const budgetMb = tier * MB_PER_GB * RAM_BUDGET_RATIO - RAM_OVERHEAD_MB;
+
+		if (weightMb <= budgetMb) return tier;
+	}
+
+	return null;
+}
@@ -1,4 +1,4 @@
-import { FILE_PATH_SEPARATOR_REGEX } from '$lib/constants';
+import { FILE_PATH_SEPARATOR_REGEX, MODEL_ID } from '$lib/constants';

 /**
 * Normalizes a model name by extracting the filename from a path, but preserves Hugging Face repository format.
@@ -56,3 +56,14 @@ export function normalizeModelName(modelName: string): string {
 export function isValidModelName(modelName: string): boolean {
 	return normalizeModelName(modelName).length > 0;
 }
+
+/**
+ * Org segment of a HuggingFace repo id (`ggml-org/Qwen3-8B` -> `ggml-org`).
+ * Returns the input itself when it carries no org separator, and an empty string
+ * for a missing id, so callers can use `||` against their own fallback org.
+ */
+export function orgOf(repoId: string | null | undefined): string {
+	if (!repoId) return '';
+
+	return repoId.split(MODEL_ID.ORG_SEPARATOR)[0] || repoId;
+}
@@ -0,0 +1,19 @@
+import { type ModelSidecar, SIDECAR_TOKENS } from '$lib/constants';
+import { ModelAuxSidecar, ModelDraftSidecar } from '$lib/enums';
+
+const SIDECAR_TOKEN_SET = new Set<string>(SIDECAR_TOKENS);
+const DRAFT_SIDECAR_SET = new Set<string>(Object.values(ModelDraftSidecar));
+const AUX_SIDECAR_SET = new Set<string>(Object.values(ModelAuxSidecar));
+
+/** Map a lowercase filename token (e.g. `mtp`) to its sidecar enum value. */
+export function sidecarFromFileToken(token: string): ModelSidecar | null {
+	return SIDECAR_TOKEN_SET.has(token) ? (token as ModelSidecar) : null;
+}
+
+export function isDraftSidecar(sidecar: ModelSidecar): sidecar is ModelDraftSidecar {
+	return DRAFT_SIDECAR_SET.has(sidecar);
+}
+
+export function isAuxSidecar(sidecar: ModelSidecar): sidecar is ModelAuxSidecar {
+	return AUX_SIDECAR_SET.has(sidecar);
+}
@@ -1,3 +1,4 @@
+import { ModelAuxSidecar, ModelDraftSidecar } from '$lib/enums';
 import { ModelsService } from '$lib/services/models.service';
 import { describe, expect, it } from 'vitest';

@@ -12,6 +13,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: null,
 			raw: 'model-name-1',
+			sidecar: null,
 			tags: []
 		});

@@ -22,6 +24,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: null,
 			raw: 'org/model-name-2',
+			sidecar: null,
 			tags: []
 		});
 	});
@@ -105,6 +108,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: 'Q2_K_XL',
 			raw: 'unsloth/DeepSeek-V4-Flash-0731-GGUF:Q2_K_XL',
+			sidecar: null,
 			tags: []
 		});

@@ -115,6 +119,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: 'Q4_K_XL',
 			raw: 'unsloth/Laguna-S-2.1-GGUF:Q4_K_XL',
+			sidecar: null,
 			tags: []
 		});

@@ -125,6 +130,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: null,
 			raw: 'org/Model-Name-GGUF',
+			sidecar: null,
 			tags: []
 		});
 	});
@@ -137,6 +143,7 @@ describe('parseModelId', () => {
 			params: '8B',
 			quantization: null,
 			raw: 'meta-llama/Llama-3.1-8B',
+			sidecar: null,
 			tags: []
 		});

@@ -147,6 +154,7 @@ describe('parseModelId', () => {
 			params: '120B',
 			quantization: 'MXFP4',
 			raw: 'openai/gpt-oss-120b-MXFP4',
+			sidecar: null,
 			tags: []
 		});

@@ -157,6 +165,7 @@ describe('parseModelId', () => {
 			params: '20B',
 			quantization: 'Q4_K_M',
 			raw: 'openai/gpt-oss-20b:Q4_K_M',
+			sidecar: null,
 			tags: []
 		});

@@ -167,6 +176,7 @@ describe('parseModelId', () => {
 			params: '30B',
 			quantization: 'BF16',
 			raw: 'Qwen/Qwen3-Coder-30B-A3B-Instruct-1M-BF16',
+			sidecar: null,
 			tags: ['Instruct', '1M']
 		});
 	});
@@ -179,6 +189,7 @@ describe('parseModelId', () => {
 			params: '17B',
 			quantization: 'Q4_K_M',
 			raw: 'meta-llama/Llama-4-Scout-17B-16E-Instruct-Q4_K_M',
+			sidecar: null,
 			tags: ['16E', 'Instruct']
 		});

@@ -189,6 +200,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: 'IQ4_XS',
 			raw: 'MiniMaxAI/MiniMax-M2-IQ4_XS',
+			sidecar: null,
 			tags: []
 		});

@@ -199,6 +211,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: 'UD-Q3_K_XL',
 			raw: 'MiniMaxAI/MiniMax-M2-UD-Q3_K_XL',
+			sidecar: null,
 			tags: []
 		});

@@ -209,6 +222,7 @@ describe('parseModelId', () => {
 			params: '123B',
 			quantization: 'Q4_K_M',
 			raw: 'mistralai/Devstral-2-123B-Instruct-2512-Q4_K_M',
+			sidecar: null,
 			tags: ['Instruct', '2512']
 		});

@@ -219,6 +233,7 @@ describe('parseModelId', () => {
 			params: '24B',
 			quantization: 'Q8_0',
 			raw: 'mistralai/Devstral-Small-2-24B-Instruct-2512-Q8_0',
+			sidecar: null,
 			tags: ['Instruct', '2512']
 		});

@@ -229,6 +244,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: 'MXFP4_MOE',
 			raw: 'noctrex/GLM-4.7-Flash-MXFP4_MOE',
+			sidecar: null,
 			tags: []
 		});

@@ -239,6 +255,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: 'Q4_K_M',
 			raw: 'Qwen/Qwen3-Coder-Next-Q4_K_M',
+			sidecar: null,
 			tags: []
 		});

@@ -249,6 +266,7 @@ describe('parseModelId', () => {
 			params: '120B',
 			quantization: 'Q4_K_M',
 			raw: 'openai/gpt-oss-120b-Q4_K_M',
+			sidecar: null,
 			tags: []
 		});

@@ -259,6 +277,7 @@ describe('parseModelId', () => {
 			params: '20B',
 			quantization: 'F16',
 			raw: 'openai/gpt-oss-20b-F16',
+			sidecar: null,
 			tags: []
 		});

@@ -269,6 +288,7 @@ describe('parseModelId', () => {
 			params: null,
 			quantization: 'Q4_K_M',
 			raw: 'nomic-embed-text-v2-moe.Q4_K_M',
+			sidecar: null,
 			tags: []
 		});
 	});
@@ -304,4 +324,41 @@ describe('parseModelId', () => {
 			tags: ['it']
 		});
 	});
+
+	it('parses sidecar file tokens', () => {
+		// sidecar prefix: bare filename or multi-slash path reduces to the filename
+		expect(parseModelId('mtp-Q4_0.gguf')).toMatchObject({
+			quantization: 'Q4_0',
+			sidecar: ModelDraftSidecar.MTP
+		});
+
+		expect(parseModelId('ggml-org/Model-GGUF/mtp-Q4_0.gguf')).toMatchObject({
+			quantization: 'Q4_0',
+			sidecar: ModelDraftSidecar.MTP
+		});
+
+		expect(parseModelId('ggml-org/Model-GGUF/mmproj-F16.gguf')).toMatchObject({
+			quantization: 'F16',
+			sidecar: ModelAuxSidecar.MMPROJ
+		});
+
+		// embedded-draft suffix: -<type> only strips when preceded by a quant
+		expect(parseModelId('ggml-org/Hy3-IQ1_M-mtp')).toMatchObject({
+			modelName: 'Hy3',
+			quantization: 'IQ1_M',
+			sidecar: ModelDraftSidecar.MTP
+		});
+
+		// a model literally named MyModel-mtp is not a draft
+		expect(parseModelId('ggml-org/MyModel-mtp')).toMatchObject({
+			modelName: 'MyModel-mtp',
+			sidecar: null
+		});
+
+		// no sidecar
+		expect(parseModelId('ggml-org/model-Q4_K_M')).toMatchObject({
+			quantization: 'Q4_K_M',
+			sidecar: null
+		});
+	});
 });
@@ -0,0 +1,127 @@
+import { ModelAuxSidecar, ModelDraftSidecar, SidecarForm } from '$lib/enums';
+import { HuggingFaceService } from '$lib/services/huggingface.service';
+import { ModelsService } from '$lib/services/models.service';
+import { describe, expect, it } from 'vitest';
+
+const { buildDownloadTag, isSidecarEntry } = ModelsService;
+const { extractQuantMeta } = HuggingFaceService;
+
+// the sidecar filename grammar mirrors the server (common/download.cpp):
+// the token must be lowercase, and it can sit at the start, between name
+// segments, or at the end of the file name
+describe('extractQuantMeta', () => {
+	it('parses the prefix form', () => {
+		expect(extractQuantMeta('mtp-Model-Q4_0.gguf')).toStrictEqual({
+			quant: 'Q4_0',
+			shared: false,
+			sidecar: ModelDraftSidecar.MTP,
+			sidecarForm: SidecarForm.PREFIX
+		});
+	});
+
+	it('parses the infix form', () => {
+		expect(extractQuantMeta('Model-mtp-Q4_0.gguf')).toStrictEqual({
+			quant: 'Q4_0',
+			shared: false,
+			sidecar: ModelDraftSidecar.MTP,
+			sidecarForm: SidecarForm.INFIX
+		});
+	});
+
+	it('parses the suffix form', () => {
+		expect(extractQuantMeta('gemma-4-E2B-it-BF16-mtp.gguf')).toStrictEqual({
+			quant: 'BF16',
+			shared: false,
+			sidecar: ModelDraftSidecar.MTP,
+			sidecarForm: SidecarForm.SUFFIX
+		});
+	});
+
+	it('parses an uppercase infix token', () => {
+		expect(extractQuantMeta('gemma-4-31B-it-MTP-BF16.gguf')).toStrictEqual({
+			quant: 'BF16',
+			shared: false,
+			sidecar: ModelDraftSidecar.MTP,
+			sidecarForm: SidecarForm.INFIX
+		});
+	});
+
+	it('parses an uppercase trailing token', () => {
+		expect(extractQuantMeta('gemma-4-E2B-it-BF16-MTP.gguf')).toStrictEqual({
+			quant: 'BF16',
+			shared: false,
+			sidecar: ModelDraftSidecar.MTP,
+			sidecarForm: SidecarForm.SUFFIX
+		});
+	});
+
+	it('parses a short-form sidecar with a bare quant', () => {
+		expect(extractQuantMeta('mmproj-F16.gguf')).toStrictEqual({
+			quant: 'F16',
+			shared: false,
+			sidecar: ModelAuxSidecar.MMPROJ,
+			sidecarForm: SidecarForm.PREFIX
+		});
+	});
+
+	it('parses a bare sidecar file', () => {
+		expect(extractQuantMeta('imatrix.gguf')).toStrictEqual({
+			quant: null,
+			shared: false,
+			sidecar: ModelAuxSidecar.IMATRIX,
+			sidecarForm: SidecarForm.PREFIX
+		});
+	});
+
+	it('parses a standalone sidecar with a draft tail', () => {
+		expect(extractQuantMeta('Model-mtp-draft.gguf')).toStrictEqual({
+			quant: null,
+			shared: false,
+			sidecar: ModelDraftSidecar.MTP,
+			sidecarForm: SidecarForm.SUFFIX
+		});
+	});
+
+	it('parses a nested sidecar path by its file name', () => {
+		expect(extractQuantMeta('MTP/mtp-Model-Q4_0.gguf')).toStrictEqual({
+			quant: 'Q4_0',
+			shared: false,
+			sidecar: ModelDraftSidecar.MTP,
+			sidecarForm: SidecarForm.PREFIX
+		});
+	});
+
+	it('returns null for non-weight files', () => {
+		expect(extractQuantMeta('README.md')).toBeNull();
+	});
+});
+
+describe('buildDownloadTag', () => {
+	it('appends the quantization', () => {
+		expect(buildDownloadTag('org/repo', 'Q4_0', null)).toBe('org/repo:Q4_0');
+	});
+
+	it('appends the quantization and sidecar', () => {
+		expect(buildDownloadTag('org/repo', 'Q4_0', ModelDraftSidecar.MTP)).toBe('org/repo:Q4_0-mtp');
+	});
+
+	it('uses the sidecar alone when there is no quant', () => {
+		expect(buildDownloadTag('org/repo', null, ModelAuxSidecar.MMPROJ)).toBe('org/repo:mmproj');
+	});
+
+	it('returns the repo id untouched without a tag', () => {
+		expect(buildDownloadTag('org/repo', null, null)).toBe('org/repo');
+	});
+});
+
+describe('isSidecarEntry', () => {
+	it('detects sidecar entries by their tag', () => {
+		expect(isSidecarEntry('org/repo:Q4_0-mtp')).toBe(true);
+		expect(isSidecarEntry('org/repo:mmproj')).toBe(true);
+	});
+
+	it('leaves plain model entries loadable', () => {
+		expect(isSidecarEntry('org/repo:Q4_0')).toBe(false);
+		expect(isSidecarEntry('org/repo')).toBe(false);
+	});
+});
Author	SHA1	Message	Date
Aleksander Grygier	b774d2c807	ui : model download pipeline Track HuggingFace downloads end to end: the server download/cancel endpoints, a status manager fed by the /models/sse download progress events, and a models-discover store holding the catalog and detail state for the discover view. Downloaded and in-flight entries are excluded from the loadable model list. Assisted-by: pi:GLM-5.3-Flash	2026-09-07 16:08:16 +02:00
Aleksander Grygier	affdf585c1	ui : model memory-fit estimation Replace the raw runtime-memory estimate with the app's compatibility check: the smallest real Mac memory tier that fits a model file, budgeted as RAM x 0.75 minus fixed overhead with headroom on the file size. The constants move to lib; the unused runtime-memory estimate is dropped. browser-info's OS detection is exported for reuse. Assisted-by: pi:GLM-5.3-Flash	2026-09-07 16:08:16 +02:00
Aleksander Grygier	b9bf09af84	ui : Hugging Face Hub data layer Add HuggingFaceService and its constants/enums/types: GGUF repo search, file tree and model detail fetching, quant/sidecar filename analysis, shard-set collapsing and the llama.app catalog feed, plus an orgOf() helper on the model name utils. Assisted-by: pi:GLM-5.3-Flash	2026-09-07 16:08:16 +02:00
Aleksander Grygier	a693dd45c2	ui : model id grammar for sidecars, quants and capability parsing Extend the shared model id parser with sidecar tokens (draft variants and auxiliary imatrix/mmproj files), weight-file and custom-quant regexes, and add the tools capability to ModelCapabilities; the selector option row picks it up from the model's declared capabilities. Assisted-by: pi:GLM-5.3-Flash	2026-09-07 16:08:16 +02:00
Aleksander Grygier	8635b58aae	ui : type-safe API types, fetch helpers and download-ready models store plumbing Assisted-by: pi:GLM-5.3-Flash	2026-09-07 16:08:15 +02:00
Aleksander Grygier	296b0f8881	server : fix deadlock when removing a finished download The download monitor thread acquires the mutex on its way out, so joining it while holding the lock in server_models::remove deadlocks once the status has flipped to DOWNLOADED. Join outside the lock, same pattern as load_models(). Assisted-by: pi:zai-org/GLM-5.3	2026-09-07 16:08:15 +02:00
Aleksander Grygier	b14462c0c9	common : resolve <quant>-<sidecar> download tags and list cached sidecars A Q4_0-mtp style tag now resolves the sidecar file when no model file matches it, so a solo draft or mmproj download actually pulls the file. Cached sidecar files list as their own entries so the state survives a restart, and removing such a tag deletes only the sidecar. Assisted-by: pi:zai-org/GLM-5.3	2026-09-07 16:08:15 +02:00
Zhaolun YinandGitHub	ccc3646c63	nix : update deprecated expressions (#28145 ) * fixed warnings * fixed nixfmt warning	2026-09-07 15:59:45 +02:00
PascalandGitHub	c0b1871bc7	webgpu: format the GET_ROWS case block (#28542 ) Brace on its own line and body indented one level, matching the surrounding cases, so the webgpu clang-format check passes.	2026-09-07 15:55:14 +02:00
Xuan-Son Nguyen GitHub Pascal	160bd031b2	server: fix LRU hang on multiple requests same model (#28539 ) * server: fix LRU hang on multiple requests same model * server: keep a queued model out of the victim pool until its waiters leave A waiter that gave up while its model was still loading left the model idle with no request behind it, and nothing recounted the free slots, so a second request queued behind it stayed queued forever. tick() was only driven by requests: join, claim and the end of a proxied request. Keep the queue entry alive after a successful claim so the model coming up is never picked as a victim before its waiters use it, and recount the slots on every status change and whenever a waiter abandons the queue. The model is then evicted as soon as it comes up with nobody left to serve. --------- Co-authored-by: Pascal <admin@serveurperso.com>	2026-09-07 15:50:46 +02:00
TitaniumtownandGitHub	dbeb37548e	sycl: add a batched L2_NORM kernel (#28222 ) * sycl: add a batched L2_NORM kernel * sycl: batch consecutive L2_NORM siblings in the graph dispatch Measured on Intel Arc Pro B70 (Battlemage), Qwen3.6-27B Q4_K_M, f16 KV, npp=128 ntg=128 npl=2, GGML_SYCL profiler: L2_NORM dispatches 12480 -> 6240 L2_NORM device time 68.77 -> 39.14 ms (-43%) total device time 6782 -> 6748 ms (-0.5%) wall decode t/s flat * tests: add L2_NORM_BATCH coverage	2026-09-07 15:24:14 +02:00
Kevin Hopper GitHub Kevin Hopper	7a333e7240	vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST) (#26578 ) * vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST) CUDA has these ops from the DeepSeek-V4 merge and Metal gained them in PR 26459. Vulkan was the last major backend running the unfused primitive chain. On DeepSeek-V4-Flash the unfused Sinkhorn comb chain alone takes about 32% of decode op time on gfx1151 (Strix Halo), spread over roughly 16k dispatches per token. dsv4_hc_comb runs the full 20-iteration Sinkhorn in registers. A token's 4x4 comb matrix lives in 16 consecutive subgroup lanes, with idst in bits 0-1 and isrc in bits 2-3 to match the CPU reference layout, so subgroupShuffleXor by 1\|2 reduces rows and by 4\|8 reduces columns. One dispatch replaces about 137 strictly ordered node executions per site. The shuffle masks never cross a 16-lane boundary, so a subgroup of size 64 packs 4 independent tokens. dsv4_hc_pre and dsv4_hc_post handle the elementwise stream collapse and fan-out, with per-token coefficients staged in shared memory. GGML_VK_DISABLE_DSV4_HC disables all three ops. The _COMB, _PRE and _POST variants gate each op independently so a single kernel can be bisected against the unfused graph. Adds eval cases at the production n_iter=20 across batch sizes that cross subgroup and workgroup boundaries. * vulkan: dsv4 hc review fixes Drop the per-op env-var disables and device flags, the stride divisibility check (ggml guarantees it) and the workgroup-count fallback in supports_op. Trim the comb shader comments to the lane layout. --------- Co-authored-by: Kevin Hopper <no-reply@maestro.press>	2026-09-07 15:24:03 +02:00
ravel7524 GitHub Carl Philipp Klemm	0c963452ea	CUDA: size routed MoE MMQ N-tiles from typical expert width on RDNA3 (#24546 ) * adjust ncols_picker for routed MoE in mul_mat_q_case function * Adding CDNA, RDNA2 and RDNA4 * fix: update mmq_use_routed_moe_ncols_picker to include NVIDIA + Volta support * feat: enhance mmq configuration for various architectures with moe_ncols_min_cc support * refactor: replace moe_ncols_min_cc with use_typical_moe_ncols in mmq configuration files * HIP: mmq: enable typical moe ncols on RDNA4 --------- Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>	2026-09-07 15:22:42 +02:00
AuroraRASandGitHub	4735997382	ggml: add gfx90c HIP support (#26454 ) * ggml: add gfx90c HIP support * ggml: make gfx90c HIP support compliant with specifications	2026-09-07 15:21:42 +02:00
fairydreaming GitHub Stanisław Szymczyk	d23c47f2a9	convert : refactor Hy4-preview conversion - move HC tensor mapping to the global map (#28451 ) Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>	2026-09-07 15:20:58 +02:00