model : support MTP in GLM-4.5-Air (#26534)

* readme : update links

* readme : update maintainer PRs list

Add the new members of the `ggml-org` `maintainers` team to the
author filter of the maintainer PRs link (nikwen, marty1885,
Titaniumtown), keeping the canonical team ordering. The list now
matches the team exactly (35 members).

Assisted-by: pi:llama.cpp/Qwen3.8-27B

CONTRIBUTING.md

@@ -74,6 +74,7 @@ For more info, please refer to the [AGENTS.md](AGENTS.md) file.
 - If a PR does not warrant a new release, add `[no release]` in the squashed commit to spare CI resources
 - Be mindful of maintenance: most of the work going into a feature happens after the PR is merged. If the PR author is not committed to contribute long-term, someone else needs to take responsibility (you)
 - Add the ["merge ready"](https://github.com/ggml-org/llama.cpp/pulls?q=is%3Apr+is%3Aopen+draft%3Ano+sort%3Aupdated-desc+label%3A%22merge+ready%22+) label to a PR to indicate when a PR can be fast-merged without waiting for 2 independent reviews. [(more info)](https://github.com/ggml-org/llama.cpp/pull/26178)
+- Wait for CI results before merging
 
 Maintainers reserve the right to decline review or close pull requests for any reason, without any questions, particularly under any of the following conditions:
 - The proposed change is already mentioned in the roadmap or an existing issue, and it has been assigned to someone.

README.md

@@ -13,7 +13,7 @@
 [![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/docker.yml?label=Docker)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
 [![Winget](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/winget.yml?label=Winget)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)
 
-[manifesto](https://github.com/ggml-org/llama.cpp/discussions/205) / [ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/issues?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Abartowski1182%20OR%20author%3Ahipudding%20OR%20author%3AServeurpersoCom%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3A0cc4m%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3Aallozaur%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Alhez%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming)%20sort%3Aupdated-desc) / [compile times](https://github.com/ggml-org/llama.cpp-dev/blob/master/README-compile-times.md) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
+[ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/issues?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Abartowski1182%20OR%20author%3Anikwen%20OR%20author%3Ahipudding%20OR%20author%3AServeurpersoCom%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3Amarty1885%20OR%20author%3A0cc4m%20OR%20author%3ATitaniumtown%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3Aallozaur%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Alhez%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming)%20sort%3Aupdated-desc) / [dev stats](https://github.com/ggml-org/llama.cpp-dev) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
 
 </div>
 

ci/run.sh

@@ -300,6 +300,40 @@ function gg_sum_ctest_release {
     gg_printf '```\n'
 }
 
+# test_llama_archs_tensor_split
+
+function gg_run_test_llama_archs_tensor_split {
+    cd ${SRC}
+
+    set -e
+
+    if [ ! -z ${GG_BUILD_CUDA} ]; then
+        GGML_CUDA_DEVICES=1 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+        GGML_CUDA_DEVICES=2 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+        GGML_CUDA_DEVICES=3 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+        GGML_CUDA_DEVICES=4 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+    fi
+
+    if [ ! -z ${GG_BUILD_METAL} ]; then
+        GGML_METAL_DEVICES=1 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+        GGML_METAL_DEVICES=2 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+        GGML_METAL_DEVICES=3 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+        GGML_METAL_DEVICES=4 ./build-ci-release/bin/test-llama-archs -s 1 2>&1
+    fi
+
+    set +e
+}
+
+function gg_sum_test_llama_archs_tensor_split {
+    gg_printf '### %s\n\n' "${ci}"
+
+    gg_printf 'Runs test-llama-archs with 1 to 4 devices\n'
+    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
+    gg_printf '```\n'
+    gg_printf '%s\n' "$(cat $OUT/${ci}.log)"
+    gg_printf '```\n'
+}
+
 # test_scripts
 
 function gg_run_test_scripts {
@@ -751,6 +785,8 @@ ret=0
 test $ret -eq 0 && gg_run ctest_debug
 test $ret -eq 0 && gg_run ctest_release
 
+test $ret -eq 0 && gg_run test_llama_archs_tensor_split
+
 if [ ! -z ${GG_BUILD_HIGH_PERF} ]; then
     test $ret -eq 0 && gg_run test_backend_ops_cpu
 fi

common/common.cpp

@@ -402,10 +402,11 @@ void common_params_print_info(const common_params & params, bool print_devices)
 #endif
     COM_TRC("%s: build %d (%s) with %s for %s%s\n", __func__, llama_build_number(), llama_commit(), llama_compiler(), llama_build_target(), build_type);
 
-    COM_INF("%s: verbosity = %d (adjust with the `-lv N` CLI arg)\n", __func__, common_log_get_verbosity_thold());
+    const int verbosity = common_log_get_verbosity_thold();
+    COM_INF("%s: verbosity = %d (adjust with the `-lv N` CLI arg)\n", __func__, verbosity);
 
     // device enumeration creates a primary context on CUDA backends, skip it when the caller does not own any device
-    if (print_devices) {
+    if (print_devices && verbosity >= LOG_LEVEL_TRACE) {
         COM_TRC("%s", "device_info:\n");
         for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
             auto * dev = ggml_backend_dev_get(i);

common/json.cpp

@@ -78,19 +78,21 @@ common_json_value::common_json_value(const common_json & val) :
 common_json_value::common_json_value(common_json && val) :
     type(VAL_JSON), val_json(std::make_shared<common_json>(std::move(val))) {}
 
+// the ctors and get<T>() below are explicit specializations, giving strong symbols
+// an explicit instantiation is a weak symbol, dropped by some LTO builds (clang-cl)
 template <typename T>
-common_json_value::common_json_value(const std::set<T> & vals) : type(VAL_JSON) {
+static std::shared_ptr<common_json> set_json(const std::set<T> & vals) {
     common_json out = common_json::array();
 
     for (const auto & val : vals) {
         out.push_back(val);
     }
 
-    val_json = std::make_shared<common_json>(std::move(out));
+    return std::make_shared<common_json>(std::move(out));
 }
 
 // a set value is usable only for the types below
-#define COMMON_JSON_SET(...) template common_json_value::common_json_value(const std::set<__VA_ARGS__> &);
+#define COMMON_JSON_SET(...) template <> common_json_value::common_json_value(const std::set<__VA_ARGS__> & vals) : type(VAL_JSON), val_json(set_json(vals)) {}
 
 COMMON_JSON_SET(int)
 COMMON_JSON_SET(std::string)
@@ -98,56 +100,45 @@ COMMON_JSON_SET(std::string)
 #undef COMMON_JSON_SET
 
 template <typename T>
-common_json_value::common_json_value(const std::map<std::string, T> & vals) : type(VAL_JSON) {
+static std::shared_ptr<common_json> map_json(const T & vals) {
     common_json out = common_json::object();
 
     for (const auto & val : vals) {
         out.set({ val.first, val.second });
     }
 
-    val_json = std::make_shared<common_json>(std::move(out));
+    return std::make_shared<common_json>(std::move(out));
 }
 
 // a map value is usable only for the types below
-#define COMMON_JSON_MAP(...) template common_json_value::common_json_value(const std::map<std::string, __VA_ARGS__> &);
+#define COMMON_JSON_MAP(...) template <> common_json_value::common_json_value(const std::map<std::string, __VA_ARGS__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
 
 COMMON_JSON_MAP(bool)
 COMMON_JSON_MAP(std::string)
 
 #undef COMMON_JSON_MAP
 
-template <typename T>
-common_json_value::common_json_value(const std::unordered_map<std::string, T> & vals) : type(VAL_JSON) {
-    common_json out = common_json::object();
-
-    for (const auto & val : vals) {
-        out.set({ val.first, val.second });
-    }
-
-    val_json = std::make_shared<common_json>(std::move(out));
-}
-
 // an unordered map value is usable only for the types below
-#define COMMON_JSON_UMAP(...) template common_json_value::common_json_value(const std::unordered_map<std::string, __VA_ARGS__> &);
+#define COMMON_JSON_UMAP(...) template <> common_json_value::common_json_value(const std::unordered_map<std::string, __VA_ARGS__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
 
 COMMON_JSON_UMAP(size_t)
 
 #undef COMMON_JSON_UMAP
 
 template <typename T>
-common_json_value::common_json_value(const std::vector<T> & vals) : type(VAL_JSON) {
+static std::shared_ptr<common_json> vec_json(const std::vector<T> & vals) {
     common_json out = common_json::array();
 
     for (const auto & val : vals) {
         out.push_back(val);
     }
 
-    val_json = std::make_shared<common_json>(std::move(out));
+    return std::make_shared<common_json>(std::move(out));
 }
 
 // a vector value is usable only for the types below
 // note: std::vector<bool> is not here, its proxy reference does not convert
-#define COMMON_JSON_VEC(...) template common_json_value::common_json_value(const std::vector<__VA_ARGS__> &);
+#define COMMON_JSON_VEC(...) template <> common_json_value::common_json_value(const std::vector<__VA_ARGS__> & vals) : type(VAL_JSON), val_json(vec_json(vals)) {}
 
 COMMON_JSON_VEC(int)
 COMMON_JSON_VEC(unsigned char)
@@ -404,10 +395,6 @@ common_json::items_view common_json::items() const {
     return items_view(const_cast<common_json *>(this), size());
 }
 
-template <typename T> T common_json::get() const {
-    return guard([&] { return as_json(this).get<T>(); });
-}
-
 // the backing library cannot build a common_json, so this one is just a copy
 template <> common_json common_json::get<common_json>() const {
     return *this;
@@ -415,7 +402,7 @@ template <> common_json common_json::get<common_json>() const {
 
 // get<T>() is usable only for the types below
 
-#define COMMON_JSON_GET(...) template __VA_ARGS__ common_json::get<__VA_ARGS__>() const;
+#define COMMON_JSON_GET(...) template <> __VA_ARGS__ common_json::get<__VA_ARGS__>() const { return guard([&] { return as_json(this).get<__VA_ARGS__>(); }); }
 
 COMMON_JSON_GET(bool)
 COMMON_JSON_GET(int)
@@ -435,3 +422,12 @@ COMMON_JSON_GET(std::vector<size_t>)
 COMMON_JSON_GET(std::unordered_map<std::string, size_t>)
 
 #undef COMMON_JSON_GET
+
+// must stay below the get<std::string> specialization
+common_json::operator std::string() const {
+    return get<std::string>();
+}
+
+std::string common_json::value(const std::string & key, const char * def) const {
+    return contains(key) ? at(key).get<std::string>() : std::string(def);
+}

common/json.h

@@ -221,16 +221,14 @@ class common_json {
     // implicit get<T>() for plain values, so they can be assigned to their C++ type directly
     // note: kept to this short list on purpose, a wider one makes j["key"] ambiguous
     // note: a numeric one would make "str = json;" ambiguous, a number converts to char too
-    operator std::string() const { return get<std::string>(); }
+    operator std::string() const;
 
     template <typename T>
     T value(const std::string & key, T def) const {
         return contains(key) ? at(key).get<T>() : def;
     }
 
-    std::string value(const std::string & key, const char * def) const {
-        return contains(key) ? at(key).get<std::string>() : std::string(def);
-    }
+    std::string value(const std::string & key, const char * def) const;
 
     // a JSON default needs no get<T>(), it is already the right type
     common_json value(const std::string & key, const common_json & def) const {

conversion/glm.py

@@ -112,12 +112,36 @@ class GlmOCRModel(Glm4Model):
 @ModelBase.example("zai-org/GLM-4.5-Air")
 class Glm4MoeModel(TextModel):
     model_arch = gguf.MODEL_ARCH.GLM4_MOE
+    supports_mtp_export = True
+    _n_main_layers: int | None = None
 
     def __init__(self, *args, **kwargs):
         super().__init__(*args, **kwargs)
-        # GLM4_MOE has num_hidden_layers + 1 actual layers (including NextN layer)
-        self.block_count = self.hparams["num_hidden_layers"] + self.hparams.get("num_nextn_predict_layers", 0)
-        self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count)
+        if not self.no_mtp:
+            self.block_count += self.hparams.get("num_nextn_predict_layers", 0)
+            self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count)
+
+    def index_tensors(self, remote_hf_model_id: str | None = None):
+        type(self)._n_main_layers = self.hparams["num_hidden_layers"]
+        return super().index_tensors(remote_hf_model_id=remote_hf_model_id)
+
+    @classmethod
+    def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
+        if (titem := super().filter_tensors(item)) is None:
+            return None
+        name, gen = titem
+
+        assert cls._n_main_layers is not None
+        is_mtp = (m := re.match(r"model\.layers\.(\d+)\.", name)) is not None and int(m.group(1)) >= cls._n_main_layers
+
+        if is_mtp and cls.no_mtp:
+            return None
+        if cls.mtp_only and not is_mtp and name not in (
+            "model.embed_tokens.weight", "model.norm.weight", "lm_head.weight",
+        ):
+            return None
+
+        return name, gen
 
     def set_vocab(self):
         return self._set_vocab_glm()
@@ -153,10 +177,22 @@ class Glm4MoeModel(TextModel):
         if (norm_topk_prob := self.hparams.get("norm_topk_prob")) is not None:
             self.gguf_writer.add_expert_weights_norm(norm_topk_prob)
 
-        # NextN/MTP prediction layers
-        if (num_nextn_predict_layers := self.hparams.get("num_nextn_predict_layers")) is not None:
+        if not self.no_mtp and (num_nextn_predict_layers := self.hparams.get("num_nextn_predict_layers")) is not None:
             self.gguf_writer.add_nextn_predict_layers(num_nextn_predict_layers)
 
+    def prepare_metadata(self, vocab_only: bool):
+        from_dir = self.fname_out.is_dir()
+        super().prepare_metadata(vocab_only=vocab_only)
+
+        if not self.mtp_only or not from_dir:
+            return
+
+        output_type: str = self.ftype.name.partition("_")[2]
+        fname_default: str = gguf.naming_convention(
+            self.metadata.name, self.metadata.basename, self.metadata.finetune,
+            self.metadata.version, size_label=None, output_type=output_type, model_type=None)
+        self.fname_out = self.fname_out.parent / f"mtp-{fname_default}.gguf"
+
     _experts: list[dict[str, Tensor]] | None = None
 
     # note: unlike GLM4V non-MoE, we don't need to permute Q/K here since GLM4V_MOE uses Neox ordering already
@@ -348,6 +384,7 @@ class GlmMoeDsaModel(DeepseekV2Model):
 @ModelBase.example("upstage/Solar-Open-100B")
 class SolarOpenModel(Glm4MoeModel):
     model_arch = gguf.MODEL_ARCH.GLM4_MOE
+    supports_mtp_export = False
 
     def set_vocab(self):
         from transformers import AutoTokenizer

docs/autoparser.md

@@ -443,21 +443,21 @@ Each returned parser is wrapped by `wrap_for_generation_prompt()`, which prepend
 |                                           | `wrap_for_generation_prompt()`, string helpers                                  |
 | `common/chat-peg-parser.h/cpp`            | `common_chat_peg_builder`, `common_chat_peg_mapper`, and helpers                |
 | `common/chat.cpp`                         | Entry point: `common_chat_templates_apply_jinja()`                              |
-| `tools/parser/debug-template-parser.cpp`  | Debug tool for template analysis                                                |
-| `tools/parser/template-analysis.cpp`      | Template analysis tool                                                          |
+| `tests/test-chat-auto-parser.cpp`         | Auto-parser unit tests; also a debug tool when given a template path            |
+| `tests/test-chat-analysis.cpp`            | Template differential analysis debug tool                                       |
 
 ## Testing & Debugging
 
 ### Debug Tools
 
-**Template Debugger**: `tools/parser/debug-template-parser.cpp`
+**Template Debugger**: `tests/test-chat-auto-parser.cpp`
 
-- Usage: `./bin/llama-debug-template-parser path/to/template.jinja`
+- Usage: `./bin/test-chat-auto-parser path/to/template.jinja` (without a path, it runs the automated tests)
 - Shows detected format, markers, generated parser, and GBNF grammar
 
-**Template Analysis**: `tools/parser/template-analysis.cpp`
+**Template Analysis**: `tests/test-chat-analysis.cpp`
 
-- Usage: `./bin/llama-template-analysis path/to/template.jinja`
+- Usage: `./bin/test-chat-analysis --template-file path/to/template.jinja` (without arguments, it runs on all templates from the test suite)
 
 **Debug Logging**: Enable with `LLAMA_ARG_LOG_VERBOSITY=2`
 
@@ -519,7 +519,7 @@ The following templates have active tests in `tests/test-chat.cpp`:
 
 To support a new template format:
 
-1. **If it follows standard patterns** — The auto-parser should detect it automatically. Run `llama-debug-template-parser` to verify markers are correctly extracted.
+1. **If it follows standard patterns** — The auto-parser should detect it automatically. Run `test-chat-auto-parser <template_path>` to verify markers are correctly extracted.
 2. **If differential analysis extracts incorrect markers** — Add a workaround lambda to the `workarounds` vector in `common/chat-diff-analyzer.cpp`. Inspect the template source for a unique identifying substring.
 3. **If it needs fundamentally different handling** — Add a dedicated handler function in `chat.cpp` before the auto-parser block (as done for GPT-OSS, Functionary v3.2, and Ministral).
 

ggml/src/ggml-backend-meta.cpp

@@ -602,27 +602,40 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
             case GGML_BACKEND_SPLIT_AXIS_1:
             case GGML_BACKEND_SPLIT_AXIS_2:
             case GGML_BACKEND_SPLIT_AXIS_3: {
-                GGML_ASSERT(src_ss[0].n_segments == 1);
-                if (src_ss[0].axis == ggml_n_dims(tensor->src[0]) - 1 && src_ss[0].nr[0] == 1) {
-                    return {ggml_backend_meta_split_axis(ggml_n_dims(tensor) - 1), {0}, {1}, 1};
-                }
-                int64_t base_ne_in = tensor->src[0]->ne[0];
-                for (int dim = 1; dim <= src_ss[0].axis; dim++) {
+                int64_t base_ne_in = 1;
+                for (int dim = 0; dim <= src_ss[0].axis; dim++) {
                     base_ne_in *= tensor->src[0]->ne[dim];
                 }
-                base_ne_in /= src_ss[0].nr[0];
+                if (src_ss[0].n_segments == 1) {
+                    base_ne_in /= src_ss[0].nr[0];
+                    if (src_ss[0].axis == ggml_n_dims(tensor->src[0]) - 1 && src_ss[0].nr[0] == 1) {
+                        return {ggml_backend_meta_split_axis(ggml_n_dims(tensor) - 1), {0}, {1}, 1};
+                    }
+                    if (src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_0 && tensor->ne[0] == tensor->src[0]->ne[0] &&
+                            tensor->ne[1] == 1 && src_ss[0].nr[0] == 1) {
+                        bool complete_rows = true;
+                        for (size_t j = 0; j < n_bufs; j++) {
+                            const int64_t ne = src_ss[0].ne[j];
+                            complete_rows = complete_rows && (ne == 0 || ne == tensor->src[0]->ne[0]);
+                        }
+                        if (complete_rows) {
+                            // Move a complete dim-0 split to the following singleton dimension.
+                            return {GGML_BACKEND_SPLIT_AXIS_1, {0}, {1}, 1};
+                        }
+                    }
+                }
+                // Reshape outputs use one segment; split-state propagation merges source segments.
                 int64_t base_ne_out = 1;
                 for (int dim = 0; dim < GGML_MAX_DIMS; dim++) {
-                    const int64_t base_ne_out_next = base_ne_out *= tensor->ne[dim];
-                    if (base_ne_out_next % base_ne_in == 0) {
-                        return {ggml_backend_meta_split_axis(dim), {0}, {uint32_t(base_ne_out_next/base_ne_in)}, 1};
+                    base_ne_out *= tensor->ne[dim];
+                    if (base_ne_out % base_ne_in == 0) {
+                        return {ggml_backend_meta_split_axis(dim), {0}, {uint32_t(base_ne_out/base_ne_in)}, 1};
                     }
-                    if (base_ne_out_next > base_ne_in) {
+                    if (base_ne_out > base_ne_in) {
                         GGML_ASSERT(src_ss[0].n_segments == 1);
                         GGML_ASSERT(src_ss[0].nr[0]      == 1);
                         return {ggml_backend_meta_split_axis(dim), {0}, {1}, 1};
                     }
-                    base_ne_out = base_ne_out_next;
                 }
                 GGML_ABORT("shape mismatch for %s", ggml_op_name(tensor->op));
             }
@@ -792,7 +805,7 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
             ggml_backend_dev_t dev = ggml_backend_buft_get_device(ggml_backend_buffer_get_type(tensor->buffer));
             const ggml_backend_meta_device_context * dev_ctx = (const ggml_backend_meta_device_context *) dev->context;
             ggml_backend_meta_split_state ret = dev_ctx->get_split_state(tensor, dev_ctx->get_split_state_ud);
-            if (ret.axis >= 0 && ret.axis <= GGML_MAX_DIMS) {
+            if (ret.axis >= 0 && ret.axis < GGML_MAX_DIMS) {
                 const int64_t granularity = ret.axis == GGML_BACKEND_SPLIT_AXIS_0 ? ggml_blck_size(tensor->type) : 1;
                 int64_t ne_sum = 0;
                 for (size_t s = 0; s < ret.n_segments; s++) {
@@ -802,6 +815,9 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
                     }
                 }
                 GGML_ASSERT(ne_sum == tensor->ne[ret.axis]);
+            } else if (ret.axis == GGML_BACKEND_SPLIT_AXIS_PARTIAL) {
+                GGML_ASSERT(ret.n_segments == 1);
+                GGML_ASSERT(ret.nr[0] == 1);
             }
             return ret;
         }
@@ -1352,15 +1368,29 @@ static void ggml_backend_meta_buffer_set_tensor(ggml_backend_buffer_t buffer, gg
         } break;
         case GGML_BACKEND_SPLIT_AXIS_PARTIAL: {
             GGML_ASSERT(tensor->type == GGML_TYPE_F32);
-            const int64_t ne = ggml_nelements(tensor);
-            std::vector<float> tmp;
-            tmp.reserve(ne);
-            for (int64_t i = 0; i < ne; i++) {
-                tmp.push_back(((const float *) data)[i] / n_bufs);
+            GGML_ASSERT(offset % sizeof(float) == 0);
+            GGML_ASSERT(size   % sizeof(float) == 0);
+            const size_t n_values = size / sizeof(float);
+            size_t n_contributors = 0;
+            for (size_t j = 0; j < n_bufs; j++) {
+                n_contributors += split_state.ne[j] != 0;
+            }
+            const bool has_contributor_mask = n_contributors != 0;
+            if (!has_contributor_mask) {
+                n_contributors = n_bufs;
+            }
+            std::vector<float> tmp(n_values);
+            for (size_t i = 0; i < n_values; i++) {
+                tmp[i] = ((const float *) data)[i] / n_contributors;
+            }
+            std::vector<float> zero;
+            if (has_contributor_mask) {
+                zero.resize(n_values, 0.0f);
             }
             for (size_t j = 0; j < n_bufs; j++) {
                 ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j);
-                ggml_backend_tensor_set(simple_tensor, tmp.data(), offset, size);
+                const float * partial = has_contributor_mask && split_state.ne[j] == 0 ? zero.data() : tmp.data();
+                ggml_backend_tensor_set(simple_tensor, partial, offset, size);
             }
         } break;
         default: {

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

@@ -38,6 +38,7 @@
 #include "ggml-cuda/out-prod.cuh"
 #include "ggml-cuda/pad.cuh"
 #include "ggml-cuda/pool2d.cuh"
+#include "ggml-cuda/pool1d.cuh"
 #include "ggml-cuda/quantize.cuh"
 #include "ggml-cuda/rope.cuh"
 #include "ggml-cuda/roll.cuh"
@@ -2326,6 +2327,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
         case GGML_OP_POOL_2D:
             ggml_cuda_op_pool2d(ctx, dst);
             break;
+        case GGML_OP_POOL_1D:
+            ggml_cuda_op_pool1d(ctx, dst);
+            break;
         case GGML_OP_SUM:
             ggml_cuda_op_sum(ctx, dst);
             break;
@@ -5245,6 +5249,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
         case GGML_OP_CONV_2D_DW:
             return op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_CONV_TRANSPOSE_2D:
+        case GGML_OP_POOL_1D:
         case GGML_OP_POOL_2D:
             return true;
         case GGML_OP_ACC:

ggml/src/ggml-cuda/pool1d.cu

@@ -0,0 +1,85 @@
+#include "pool1d.cuh"
+
+static __global__ void pool1d_nchw_kernel(
+        const int iw, const int ow,
+        const int kw, const int sw, const int pw,
+        const int parallel_elements,
+        const float * src, float * dst, const enum ggml_op_pool op) {
+    const int idx = threadIdx.x + blockIdx.x * blockDim.x;
+    if (idx >= parallel_elements) {
+        return;
+    }
+
+    const int nc = idx / ow;
+    const int cur_ow = idx % ow;
+
+    const float * i_ptr = src + nc * iw;
+    float * o_ptr = dst + nc * ow;
+
+    const int start = cur_ow * sw - pw;
+    const int b = max(0, start);
+    const int e = min(iw, start + kw);
+
+    float res;
+    switch (op) {
+        case GGML_OP_POOL_AVG: res = 0.0f;     break;
+        case GGML_OP_POOL_MAX: res = -FLT_MAX; break;
+        default: return;
+    }
+
+    int count = 0;
+    for (int i = b; i < e; i++) {
+#if __CUDA_ARCH__ >= 350
+        float cur = __ldg(i_ptr + i);
+#else
+        float cur = i_ptr[i];
+#endif
+        switch (op) {
+            case GGML_OP_POOL_AVG: res += cur;                break;
+            case GGML_OP_POOL_MAX: res = max(res, cur);       break;
+            default: break;
+        }
+        count++;
+    }
+
+    if (op == GGML_OP_POOL_AVG) {
+        res = (count > 0) ? (res / count) : 0.0f;
+    }
+
+    o_ptr[cur_ow] = res;
+}
+
+static void pool1d_nchw_kernel_f32_f32_cuda(
+        const int iw, const int ow,
+        const int kw, const int sw, const int pw,
+        const int parallel_elements,
+        const float * src, float * dst, const enum ggml_op_pool op,
+        cudaStream_t stream) {
+    const int num_blocks = (parallel_elements + CUDA_POOL1D_BLOCK_SIZE - 1) / CUDA_POOL1D_BLOCK_SIZE;
+    dim3 block_nums(num_blocks);
+    pool1d_nchw_kernel<<<block_nums, CUDA_POOL1D_BLOCK_SIZE, 0, stream>>>(iw, ow, kw, sw, pw, parallel_elements, src, dst, op);
+}
+
+void ggml_cuda_op_pool1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+    const ggml_tensor * src0 = dst->src[0];
+    const float * src0_d = (const float *)src0->data;
+    float * dst_d = (float *)dst->data;
+    cudaStream_t stream = ctx.stream();
+
+    GGML_ASSERT(src0->type == GGML_TYPE_F32);
+    GGML_ASSERT( dst->type == GGML_TYPE_F32);
+
+    const int32_t * opts = (const int32_t *)dst->op_params;
+    enum ggml_op_pool op = static_cast<ggml_op_pool>(opts[0]);
+    const int k0 = opts[1];
+    const int s0 = opts[2];
+    const int p0 = opts[3];
+
+    const int64_t IW = src0->ne[0];
+    const int64_t OW = dst->ne[0];
+    const int64_t nr = ggml_nrows(src0);
+
+    const int parallel_elements = (int)(nr * OW);
+
+    pool1d_nchw_kernel_f32_f32_cuda(IW, OW, k0, s0, p0, parallel_elements, src0_d, dst_d, op, stream);
+}

ggml/src/ggml-cuda/pool1d.cuh

@@ -0,0 +1,5 @@
+#include "common.cuh"
+
+#define CUDA_POOL1D_BLOCK_SIZE 256
+
+void ggml_cuda_op_pool1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst);

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

@@ -955,6 +955,7 @@ struct vk_device_struct {
     vk_pipeline pipeline_diag[2];
     vk_pipeline pipeline_clamp[2];
     vk_pipeline pipeline_pad_f32;
+    vk_pipeline pipeline_pad_reflect_1d_f32;
     vk_pipeline pipeline_roll_f32;
     vk_pipeline pipeline_repeat_i32, pipeline_repeat_back_f32;
     vk_pipeline pipeline_repeat_i16;
@@ -5630,6 +5631,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
     ggml_vk_create_pipeline(device, device->pipeline_diag[1], "diag_f16", diag_f16_len, diag_f16_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1);
 
     ggml_vk_create_pipeline(device, device->pipeline_pad_f32, "pad_f32", pad_f32_len, pad_f32_data, "main", 2, sizeof(vk_op_pad_push_constants), {512, 1, 1}, {}, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_pad_reflect_1d_f32, "pad_reflect_1d_f32", pad_reflect_1d_f32_len, pad_reflect_1d_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1);
 
     ggml_vk_create_pipeline(device, device->pipeline_roll_f32, "roll_f32", roll_f32_len, roll_f32_data, "main", 2, sizeof(vk_op_unary_push_constants), {512, 1, 1}, {}, 1);
 
@@ -11336,6 +11338,11 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const
             return ctx->device->pipeline_pad_f32;
         }
         return nullptr;
+    case GGML_OP_PAD_REFLECT_1D:
+        if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
+            return ctx->device->pipeline_pad_reflect_1d_f32;
+        }
+        return nullptr;
     case GGML_OP_ROLL:
         if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
             return ctx->device->pipeline_roll_f32;
@@ -12239,6 +12246,7 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co
     case GGML_OP_CLAMP:
     case GGML_OP_LEAKY_RELU:
     case GGML_OP_PAD:
+    case GGML_OP_PAD_REFLECT_1D:
     case GGML_OP_ROLL:
     case GGML_OP_REPEAT:
     case GGML_OP_REPEAT_BACK:
@@ -13111,6 +13119,17 @@ static void ggml_vk_pad(ggml_backend_vk_context * ctx, vk_context& subctx, const
     ggml_vk_op_f32(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_PAD, std::move(p));
 }
 
+static void ggml_vk_pad_reflect_1d(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
+    const uint32_t p0 = (uint32_t)dst->op_params[0];
+    const uint32_t p1 = (uint32_t)dst->op_params[1];
+
+    vk_op_unary_push_constants p = vk_op_unary_push_constants_init(src0, dst, ggml_nelements(dst));
+    memcpy(&p.param1, &p0, sizeof(float));
+    memcpy(&p.param2, &p1, sizeof(float));
+
+    ggml_vk_op_f32(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_PAD_REFLECT_1D, std::move(p));
+}
+
 static void ggml_vk_roll(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
     const int32_t s0 = ggml_get_op_params_i32(dst, 0);
     const int32_t s1 = ggml_get_op_params_i32(dst, 1);
@@ -15520,6 +15539,10 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr
     case GGML_OP_PAD:
         ggml_vk_pad(ctx, compute_ctx, src0, node);
 
+        break;
+    case GGML_OP_PAD_REFLECT_1D:
+        ggml_vk_pad_reflect_1d(ctx, compute_ctx, src0, node);
+
         break;
     case GGML_OP_ROLL:
         ggml_vk_roll(ctx, compute_ctx, src0, node);
@@ -18446,6 +18469,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
         case GGML_OP_SCALE:
             return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_PAD:
+        case GGML_OP_PAD_REFLECT_1D:
         case GGML_OP_ROLL:
             return op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_DIAG_MASK_INF:
@@ -19228,6 +19252,8 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph *
         } else if (tensor->op == GGML_OP_PAD) {
             tensor_clone = ggml_pad_ext(ggml_ctx, src_clone[0], tensor->op_params[0], tensor->op_params[1], tensor->op_params[2], tensor->op_params[3],
                                                                 tensor->op_params[4], tensor->op_params[5], tensor->op_params[6], tensor->op_params[7]);
+        } else if (tensor->op == GGML_OP_PAD_REFLECT_1D) {
+            tensor_clone = ggml_pad_reflect_1d(ggml_ctx, src_clone[0], tensor->op_params[0], tensor->op_params[1]);
         } else if (tensor->op == GGML_OP_REPEAT) {
             tensor_clone = ggml_repeat(ggml_ctx, src_clone[0], tensor);
         } else if (tensor->op == GGML_OP_REPEAT_BACK) {

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

@@ -0,0 +1,43 @@
+#version 450
+
+#include "types.glsl"
+#include "generic_unary_head.glsl" // included to use functions like fastdiv etc.
+
+layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in;
+
+void main() {
+
+    const uint idx = get_idx();
+
+    if (idx >= p.ne) {
+        return;
+    }
+
+    const uint p0 = floatBitsToUint(p.param1);
+    const uint p1 = floatBitsToUint(p.param2);
+
+    const uint i3 = fastdiv(idx, p.ne1_012mp, fastdiv_L(p.ne1_Ls, 0));
+    const uint i3_offset = i3 * p.ne12 * p.ne11 * p.ne10;
+
+    const uint i2 = fastdiv(idx - i3_offset, p.ne1_01mp, fastdiv_L(p.ne1_Ls, 1));
+    const uint i2_offset = i2 * p.ne11 * p.ne10;
+
+    const uint i1 = fastdiv(idx - i3_offset - i2_offset, p.ne1_0mp, fastdiv_L(p.ne1_Ls, 2));
+    const uint i0 = idx - i3_offset - i2_offset - i1 * p.ne10;
+
+    uint src_col;
+
+    if (i0 < p0) {
+        src_col = p0 - i0;    // left pad area
+    } else if (i0 < p0 + p.ne00) {
+        src_col = i0 - p0;    // center area
+    } else {
+        src_col = 2u * p.ne00 - 2u - (i0 - p0);  // right pad area
+    }
+
+    const uint src_idx = i3 * p.nb03 + i2 * p.nb02 + i1 * p.nb01 + src_col * p.nb00;
+    const uint d_idx = i3 * p.nb13 + i2 * p.nb12 + i1 * p.nb11 + i0 * p.nb10;
+
+    // copy the computed value to the destination tensor
+    data_d[get_doffset() + d_idx] = D_TYPE(data_a[get_aoffset() + src_idx]);
+}

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

@@ -896,6 +896,7 @@ void process_shaders() {
     string_to_spv("scale_f32", "scale.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}});
 
     string_to_spv("pad_f32", "pad.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}});
+    string_to_spv("pad_reflect_1d_f32", "pad_reflect_1d.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}});
 
     string_to_spv("concat_i8", "concat.comp", {{"A_TYPE", "uint8_t"}, {"B_TYPE", "uint8_t"}, {"D_TYPE", "uint8_t"}});
     string_to_spv("concat_i16", "concat.comp", {{"A_TYPE", "uint16_t"}, {"B_TYPE", "uint16_t"}, {"D_TYPE", "uint16_t"}});

gguf-py/gguf/constants.py

@@ -3822,7 +3822,7 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
         MODEL_TENSOR.FFN_DOWN_SHEXP,
         MODEL_TENSOR.FFN_UP_SHEXP,
         MODEL_TENSOR.FFN_EXP_PROBS_B,
-        # NextN/MTP tensors - preserved but unused
+        # NextN/MTP tensors
         MODEL_TENSOR.NEXTN_EH_PROJ,
         MODEL_TENSOR.NEXTN_EMBED_TOKENS,
         MODEL_TENSOR.NEXTN_ENORM,

include/llama.h

@@ -733,7 +733,7 @@ extern "C" {
 
     // Removes all tokens that belong to the specified sequence and have positions in [p0, p1)
     // Returns false if a partial sequence cannot be removed. Removing a whole sequence never fails
-    // seq_id < 0 : match any sequence
+    // seq_id < 0 : match any sequence [TAG_LLAMA_SEQ_ID_NEG]
     // p0 < 0     : [0,  p1]
     // p1 < 0     : [p0, inf)
     LLAMA_API bool llama_memory_seq_rm(

scripts/sync_vendor.py

@@ -27,7 +27,7 @@ vendor = {
     f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/split.py":  "split.py",
     f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/LICENSE":   "vendor/cpp-httplib/LICENSE",
 
-    "https://raw.githubusercontent.com/sheredom/subprocess.h/9ce0d701b6fb10f8f8c4445edd31e7c60a1237e3/subprocess.h": "vendor/sheredom/subprocess.h",
+    "https://raw.githubusercontent.com/sheredom/subprocess.h/0dccaa9aa176dd6d7ef8afeca3c18d6e80a32795/subprocess.h": "vendor/sheredom/subprocess.h",
 
     f"https://raw.githubusercontent.com/Cyan4973/xxHash/{XXHASH_COMMIT}/xxhash.c":      "vendor/hash/xxhash/xxhash.c",
     f"https://raw.githubusercontent.com/Cyan4973/xxHash/{XXHASH_COMMIT}/xxhash.h":      "vendor/hash/xxhash/xxhash.h",

skills/add-new-model/SKILL.md

@@ -66,7 +66,7 @@ These recur often enough in review comments on past add-model PRs that they're w
 - Optional hparams that are genuinely absent from some configs (e.g. a shared-expert count) should be read with an explicit optional/fallback accessor, not assumed present.
 - Hparams that are actually load-bearing (the model produces wrong output or crashes without them, e.g. `sliding_window_pattern`, norm-eps) must hard-error if missing, not silently fall back to a default.
 - Don't bake a default chat template into the C++ binary - inject it into the GGUF at conversion time instead, since one `llm_arch` can be reused by multiple fine-tunes with different templates, and a baked-in C++ default fails silently for those.
-- Before writing a dedicated tool-call/output parser, check whether the existing autoparser already handles the template (`llama-debug-template-parser <jinja>` shows what it detects).
+- Before writing a dedicated tool-call/output parser, check whether the existing autoparser already handles the template (`test-chat-auto-parser <jinja>` shows what it detects).
 - Marking a custom EOS/closing-tag token as `eot` at conversion time isn't always sufficient - in long/agentic generations a model can emit the closing sequence as literal text instead of the token, so generation never stops on EOG and raw text leaks past the parser. Verify this case, not just the token path.
 - If reusing or aliasing an existing pre-tokenizer for convenience, justify and test that choice explicitly - silent reuse is an easy source of subtle tokenizer bugs.
 - Watch for excessive graph splits caused by building per-layer view/index tensors inside the layer loop - hoist tensors that don't vary per layer out of the loop (relevant if you hit `GGML_SCHED_MAX_SPLIT_INPUTS`).

src/llama-context.cpp

@@ -3218,8 +3218,6 @@ size_t llama_context::state_read_data(llama_io_read_i & io) {
 }
 
 size_t llama_context::state_seq_write_data(llama_io_write_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) {
-    GGML_UNUSED(seq_id);
-
     if (memory) {
         memory->state_write(io, seq_id, flags);
     }
@@ -3228,8 +3226,6 @@ size_t llama_context::state_seq_write_data(llama_io_write_i & io, llama_seq_id s
 }
 
 size_t llama_context::state_seq_read_data(llama_io_read_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) {
-    GGML_UNUSED(seq_id);
-
     if (memory) {
         memory->state_read(io, seq_id, flags);
     }

src/llama-kv-cache-dsv4.cpp

@@ -599,6 +599,33 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
         }
     }
 
+    if (ratio == DSV4_HCA_RATIO && !plan.state_pos.empty() && plan.state_write_idxs.empty()) {
+        assert(kv_size > 0);
+        // the last slot must not be live, or the dummy write would corrupt it;
+        // a full stream implies a completed block, which implies real writes
+        assert(plan.n_kv < (int64_t) kv_size);
+
+        // Keep the compress/write ops in the graph when no HCA block completes
+        // in this ubatch. The dummy block writes to the last cache slot and is
+        // masked out.
+        uint32_t i = 0;
+        while (i < ubatch.n_tokens && ubatch.pos[i] < 0) {
+            ++i;
+        }
+        assert(i < ubatch.n_tokens);
+
+        const llama_seq_id seq_id = ubatch.seq_id[i][0];
+        const int64_t cache_off = dsv4_stream_offset(n_stream, seq_id, kv_size);
+        const int32_t source_idx = state_source_idx(seq_id, ubatch.pos[i]);
+
+        plan.state_write_idxs.push_back(cache_off + kv_size - 1);
+        plan.state_write_pos .push_back(0);
+
+        for (uint32_t j = 0; j < ratio; ++j) {
+            plan.state_read_idxs.push_back(source_idx);
+        }
+    }
+
     if (overlap) {
         // [ all blocks' prev-window indices | all blocks' cur-window indices ]
         plan.state_read_idxs.reserve(overlap_prev_reads.size() + overlap_cur_reads.size());
@@ -608,7 +635,10 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
                 overlap_cur_reads.begin(), overlap_cur_reads.end());
     }
 
-    plan.n_kv = GGML_PAD(plan.n_kv, 256u);
+    // Keep the mask (and with it the compressed-attention branch) present even
+    // before the first block is visible, so the graph topology never changes.
+    // Padded slots are masked out; comp cache buffers are zero-initialized.
+    plan.n_kv = std::max<int64_t>(GGML_PAD(plan.n_kv, 256u), 256);
 
     std::sort(persist_rows.begin(), persist_rows.end(),
             [](const persist_row & a, const persist_row & b) {
@@ -620,16 +650,26 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
         plan.state_persist_dst_idxs.push_back(row.dst);
     }
 
-
     if (n_rs_seq > 0) {
-        for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
-            const llama_seq_id seq_id = ubatch.seq_id_unq[s];
-            if (seq_id < 0 || (uint32_t) seq_id >= n_stream) {
-                continue;
+        // Emit restore/snapshot entries for all layout streams so that the
+        // graph tensor sizes do not depend on the ubatch's sequence count.
+        // Streams not present in the ubatch get no-op entries.
+        for (uint32_t stream = 0; stream < n_stream; ++stream) {
+            llama_seq_id seq_id = -1;
+            if (n_stream == 1) {
+                // a unified stream serves any single sequence
+                seq_id = ubatch.n_seqs_unq > 0 ? ubatch.seq_id_unq[0] : -1;
+            } else {
+                for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
+                    if (ubatch.seq_id_unq[s] == (llama_seq_id) stream) {
+                        seq_id = ubatch.seq_id_unq[s];
+                        break;
+                    }
+                }
             }
 
-            const int64_t stream_off = dsv4_stream_offset(n_stream, seq_id, state_size);
-            const uint32_t rollback = (uint32_t) seq_id < rs_idx.size() ? rs_idx[seq_id] : 0;
+            const int64_t stream_off = (int64_t) stream*state_size;
+            const uint32_t rollback = seq_id >= 0 && (uint32_t) seq_id < rs_idx.size() ? rs_idx[seq_id] : 0;
             // Keep the restore graph fixed-width when no rollback is pending.
             const int64_t src_plane = rollback > 0 && rollback <= n_rs_seq ? (int64_t) rollback*state_rows : 0;
             for (uint32_t r = 0; r < state_size; ++r) {
@@ -639,35 +679,33 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
 
             std::vector<uint32_t> token_idxs;
             token_idxs.reserve(ubatch.n_tokens);
-            for (uint32_t i = 0; i < ubatch.n_tokens; ++i) {
-                if (dsv4_token_has_seq(ubatch, i, seq_id)) {
-                    token_idxs.push_back(i);
+            if (seq_id >= 0) {
+                for (uint32_t i = 0; i < ubatch.n_tokens; ++i) {
+                    if (dsv4_token_has_seq(ubatch, i, seq_id)) {
+                        token_idxs.push_back(i);
+                    }
                 }
             }
-            if (token_idxs.empty()) {
-                continue;
-            }
 
             const uint32_t n_seq_tokens = (uint32_t) token_idxs.size();
             const int64_t scratch_off = (int64_t) state_rows*(1 + n_rs_seq);
             for (uint32_t d = 1; d <= n_rs_seq; ++d) {
                 const int64_t dst_plane = (int64_t) d*state_rows;
+                const uint32_t prefix = d <= n_seq_tokens ? n_seq_tokens - d : 0;
 
                 for (uint32_t r = 0; r < state_size; ++r) {
-                    int32_t src;
-                    if (d <= n_seq_tokens) {
-                        const uint32_t prefix = n_seq_tokens - d;
-                        src = (int32_t) (stream_off + r);
+                    int32_t src = (int32_t) (stream_off + r);
 
-                        for (uint32_t j = 0; j < prefix; ++j) {
-                            const uint32_t i_tok = token_idxs[j];
-                            if (ubatch.pos[i_tok] >= 0 && (uint32_t) (ubatch.pos[i_tok]%state_size) == r) {
-                                src = (int32_t) (scratch_off + i_tok);
-                            }
+                    for (uint32_t j = 0; j < prefix; ++j) {
+                        const uint32_t i_tok = token_idxs[j];
+                        if (ubatch.pos[i_tok] >= 0 && (uint32_t) (ubatch.pos[i_tok]%state_size) == r) {
+                            src = (int32_t) (scratch_off + i_tok);
                         }
-                    } else {
-                        const int64_t src_plane = (int64_t) (d - n_seq_tokens)*state_rows;
-                        src = (int32_t) (src_plane + stream_off + r);
+                    }
+
+                    if (n_seq_tokens == 0) {
+                        // no-op: copy the snapshot plane onto itself
+                        src = (int32_t) (dst_plane + stream_off + r);
                     }
 
                     plan.state_snapshot_src_idxs.push_back(src);
@@ -683,10 +721,16 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
     }();
 
     if (debug) {
-        LLAMA_LOG_INFO("%s: ratio=%u, n_tokens=%u, state_persist_dst=%s, state_write_pos=%s\n",
-                __func__, ratio, ubatch.n_tokens,
+        LLAMA_LOG_DEBUG("%s: ratio=%u, n_tokens=%u, n_seqs_unq=%u, state_persist_dst=%s, state_write_pos=%s\n",
+                __func__, ratio, ubatch.n_tokens, ubatch.n_seqs_unq,
                 dsv4_plan_positions(plan.state_persist_dst_idxs).c_str(),
                 dsv4_plan_positions(plan.state_write_pos).c_str());
+        for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
+            const llama_seq_id seq_id = ubatch.seq_id_unq[s];
+            const uint32_t rollback = seq_id >= 0 && (uint32_t) seq_id < rs_idx.size() ? rs_idx[seq_id] : 0;
+            LLAMA_LOG_DEBUG("%s:   seq %d pos [%d, %d] rollback=%u\n", __func__, seq_id,
+                    ubatch.pos[0], ubatch.pos[ubatch.n_tokens - 1], rollback);
+        }
     }
 
     return plan;
@@ -704,8 +748,17 @@ static std::vector<llama_kv_cache_dsv4_context::comp_plan> dsv4_build_comp_plans
     std::vector<llama_kv_cache_dsv4_context::comp_plan> plans;
     plans.reserve(ubatches.size());
 
+    // the first ubatch touching a seq consumes its rollback restore
+    std::vector<uint32_t> rs(rs_idx);
     for (const llama_ubatch & ubatch : ubatches) {
-        plans.push_back(dsv4_build_comp_plan(ubatch, ratio, overlap, state_size, kv_size, n_stream, n_rs_seq, rs_idx));
+        plans.push_back(dsv4_build_comp_plan(ubatch, ratio, overlap, state_size, kv_size, n_stream, n_rs_seq, rs));
+
+        for (uint32_t s = 0; s < ubatch.n_seqs_unq; ++s) {
+            const llama_seq_id seq_id = ubatch.seq_id_unq[s];
+            if (seq_id >= 0 && (size_t) seq_id < rs.size()) {
+                rs[seq_id] = 0;
+            }
+        }
     }
 
     return plans;
@@ -803,16 +856,15 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_reserve_comp_plan(
         return plan;
     }
 
-    const uint32_t n_seqs       = std::max<uint32_t>(1, ubatch.n_seqs);
-    const uint32_t n_seq_tokens = std::max<uint32_t>(1, ubatch.n_seq_tokens);
-    const uint64_t n_blocks_u64 = (uint64_t) n_seqs*((n_seq_tokens + ratio - 1)/ratio);
-    const size_t n_blocks = (size_t) std::max<uint64_t>(1, n_blocks_u64);
-    GGML_ASSERT((uint64_t) n_blocks == std::max<uint64_t>(1, n_blocks_u64));
+    // worst case over every seq split: sum of per-seq ceil(tokens/ratio) is at
+    // most floor(n_tokens/ratio) + n_seqs
+    const uint32_t n_seqs = std::max<uint32_t>(1, ubatch.n_seqs);
+    const size_t n_blocks = (size_t) ubatch.n_tokens/ratio + n_seqs;
 
     const uint64_t state_rows = (uint64_t) state_size*n_stream;
     const size_t n_persist = (size_t) std::min<uint64_t>(ubatch.n_tokens, state_rows);
-    const size_t n_restore = n_rs_seq > 0 ? (size_t) state_size*std::max<uint32_t>(1, ubatch.n_seqs_unq) : 0;
-    const size_t n_snapshot = (size_t) n_rs_seq*state_size*std::max<uint32_t>(1, ubatch.n_seqs_unq);
+    const size_t n_restore = n_rs_seq > 0 ? (size_t) state_size*n_stream : 0;
+    const size_t n_snapshot = (size_t) n_rs_seq*state_size*n_stream;
 
     plan.state_pos .resize(ubatch.n_tokens);
     plan.state_persist_src_idxs.resize(n_persist);
@@ -1356,7 +1408,9 @@ llama_memory_context_ptr llama_kv_cache_dsv4::init_batch(
             if (has_coupled) {
                 ubatch = balloc.split_seq(n_ubatch);
             } else {
-                ubatch = balloc.split_equal(n_ubatch, raw_per_seq || comp_per_seq, 0);
+                // [TAG_RECURRENT_ROLLBACK_SPLITS]
+                // the trailing (1 + n_rs_seq) tokens of each seq must stay in the same ubatch
+                ubatch = balloc.split_equal(n_ubatch, raw_per_seq || comp_per_seq, n_rs_seq > 0 ? n_rs_seq + 1 : 0);
             }
 
             if (ubatch.n_tokens == 0) {
@@ -1433,6 +1487,11 @@ bool llama_kv_cache_dsv4::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1
             return false;
         }
 
+        // pending rollback is single-use: stacked partial removals don't compose
+        if (rs_idx[seq_id] != 0) {
+            return false;
+        }
+
         const bool res = kv_raw->seq_rm(seq_id, p0, p1);
         if (res) {
             rs_idx[seq_id] = (uint32_t) rollback;
@@ -1594,9 +1653,7 @@ void llama_kv_cache_dsv4::state_read(llama_io_read_i & io, llama_seq_id seq_id,
     kv_raw->state_read(io, seq_id, flags);
 
     if (!partial_only) {
-        kv_csa->clear(true);
-        kv_hca->clear(true);
-        kv_lid->clear(true);
+        clear_compressed(seq_id, true);
 
         dsv4_state_read_k_cache(io, kv_csa.get(), seq_id, flags);
         dsv4_state_read_k_cache(io, kv_hca.get(), seq_id, flags);

src/llama-kv-cache.cpp

@@ -383,6 +383,7 @@ bool llama_kv_cache::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1) {
         return true;
     }
 
+    // TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
     GGML_ASSERT(seq_id == -1 || (seq_id >= 0 && (size_t) seq_id < seq_to_stream.size()));
 
     if (p0 < 0) {
@@ -2043,6 +2044,7 @@ void llama_kv_cache::state_read(llama_io_read_i & io, llama_seq_id seq_id, llama
 
     GGML_UNUSED(flags);
 
+    // TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
     GGML_ASSERT(seq_id == -1 || (seq_id >= 0 && (size_t) seq_id < seq_to_stream.size()));
 
     uint32_t n_stream_cur;

src/llama-memory-recurrent.cpp

@@ -158,13 +158,14 @@ bool llama_memory_recurrent::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos
         p1 = std::numeric_limits<llama_pos>::max();
     }
 
+    if ((uint32_t) seq_id >= this->n_seq_max) {
+        LLAMA_LOG_ERROR("%s: invalid seq_id (%d) - larger than n_seq_max (%d)\n", __func__, seq_id, this->n_seq_max);
+        return false;
+    }
+
     const bool rm_all = p0 == 0 && p1 == std::numeric_limits<llama_pos>::max();
     if (rm_all) {
-        if (seq_id >= 0) {
-            set_rs_idx(seq_id, 0);
-        } else {
-            std::fill(rs_idx.begin(), rs_idx.end(), 0);
-        }
+        set_rs_idx(seq_id, 0);
     }
 
     // models like Mamba or RWKV can't have a state partially erased at the end
@@ -181,7 +182,9 @@ bool llama_memory_recurrent::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos
             // partial rollback via per-token snapshot index (bounded by n_rs_seq)
             if (0 < p0 && p0 <= cell.pos && p1 > cell.pos) {
                 const llama_pos rollback = cell.pos - (p0 - 1);
-                if (rollback >= 1 && rollback <= (llama_pos) n_rs_seq) {
+                // pending rollback is single-use
+                const bool pending = rs_idx[seq_id] != 0;
+                if (!pending && rollback >= 1 && rollback <= (llama_pos) n_rs_seq) {
                     set_rs_idx(seq_id, (uint32_t) rollback);
                     cell.pos = p0 - 1;
                     return true;
@@ -390,10 +393,17 @@ llama_pos llama_memory_recurrent::seq_pos_max(llama_seq_id seq_id) const {
 }
 
 void llama_memory_recurrent::set_rs_idx(llama_seq_id seq_id, uint32_t idx) {
-    if (seq_id < 0 || (size_t) seq_id >= rs_idx.size()) {
+    if (seq_id < 0) {
+        std::fill(rs_idx.begin(), rs_idx.end(), 0);
         return;
     }
-    rs_idx[seq_id] = (idx > n_rs_seq) ? n_rs_seq : idx;
+
+    assert(n_seq_max == rs_idx.size());
+
+    GGML_ASSERT((uint32_t) seq_id < n_seq_max);
+    GGML_ASSERT(idx <= n_rs_seq);
+
+    rs_idx[seq_id] = idx;
 }
 
 std::map<ggml_backend_buffer_type_t, size_t> llama_memory_recurrent::memory_breakdown() const {
@@ -742,6 +752,7 @@ void llama_memory_recurrent::state_write(llama_io_write_i & io, llama_seq_id seq
     uint32_t cell_range_begin = size;
     for (uint32_t i = 0; i < size; ++i) {
         const auto & cell = cells[i];
+        // TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
         if ((seq_id == -1 && !cell.is_empty()) || cell.has_seq_id(seq_id)) {
             ++cell_count;
             uint32_t rs_idx_cur = 0;
@@ -827,6 +838,7 @@ void llama_memory_recurrent::state_read(llama_io_read_i & io, llama_seq_id seq_i
     }
 
     if (!res) {
+        // TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
         if (seq_id == -1) {
             clear(true);
         } else {
@@ -836,11 +848,7 @@ void llama_memory_recurrent::state_read(llama_io_read_i & io, llama_seq_id seq_i
     }
 
     if (n_rs_seq != 0) {
-        if (seq_id == -1) {
-            std::fill(rs_idx.begin(), rs_idx.end(), 0);
-        } else {
-            set_rs_idx(seq_id, 0);
-        }
+        set_rs_idx(seq_id, 0);
     }
 }
 

src/llama-model-saver.cpp

@@ -293,6 +293,14 @@ void llama_model_saver::add_kv_from_model() {
     add_kv(LLM_KV_ATTENTION_INDEXER_LOCAL_BLOCKS,    hparams.indexer_local_blocks);
     add_kv(LLM_KV_ATTENTION_INDEXER_TYPES,           hparams.is_indexer_full_impl, true);
     add_kv(LLM_KV_ATTENTION_RECURRENT_LAYERS,        hparams.is_recr_impl, true);
+    add_kv(LLM_KV_ATTENTION_OUTPUT_GROUP_COUNT,      hparams.dsv4_o_group_count);
+    add_kv(LLM_KV_ATTENTION_OUTPUT_LORA_RANK,        hparams.dsv4_o_lora_rank);
+    add_kv(LLM_KV_ATTENTION_COMPRESS_ROPE_FREQ_BASE, hparams.dsv4_compress_rope_base);
+    add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS,         hparams.dsv4_compress_ratios, true);
+    add_kv(LLM_KV_HYPER_CONNECTION_COUNT,            hparams.dsv4_hc_mult);
+    add_kv(LLM_KV_HYPER_CONNECTION_SINKHORN_ITERATIONS, hparams.dsv4_hc_sinkhorn_iters);
+    add_kv(LLM_KV_HYPER_CONNECTION_EPSILON,          hparams.dsv4_hc_eps);
+    add_kv(LLM_KV_HASH_LAYER_COUNT,                  hparams.dsv4_hash_layer_count);
 
     const float rope_scaling_factor = hparams.rope_freq_scale_train == 1.0f ? 0.0f : 1.0f/hparams.rope_freq_scale_train;
 
@@ -422,6 +430,9 @@ void llama_model_saver::add_tensors_from_model() {
     add_tensor(model->cls_out);
     add_tensor(model->cls_out_b);
     add_tensor(model->cls_norm);
+    add_tensor(model->hc_head_fn);
+    add_tensor(model->hc_head_base);
+    add_tensor(model->hc_head_scale);
 
     for (const struct llama_layer & layer : model->layers) {
         for (size_t i = 0; i < sizeof(layer)/sizeof(struct ggml_tensor *); ++i) {

src/llama-model.cpp

@@ -520,7 +520,7 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
             return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_MIRRORED);
         }
         if (std::regex_match(tensor_name, pattern_ffn_down_exps_bias)) {
-            return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_PARTIAL);
+            return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_PARTIAL, "ffn_down_exps.weight");
         }
 
         // output
@@ -554,6 +554,9 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
                     GGML_ASSERT(tensor->ne[axis] == 2*key_dim + value_dim);
                     return {{key_dim, 2}, {value_dim, 1}};
                 }
+                if (std::regex_match(tensor_name, pattern_r_cache)) {
+                    return {{key_dim * (hparams.ssm_d_conv - 1), 2}, {value_dim * (hparams.ssm_d_conv - 1), 1}};
+                }
             } else {
                 const int64_t head_ratio = n_v_heads / n_k_heads;
                 if (std::regex_match(tensor_name, pattern_qkv_weight) || std::regex_match(tensor_name, pattern_ssm_conv1d)) {
@@ -642,12 +645,12 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
                 blck_size_perf *= 2;
             }
 
+            const int64_t granularity_q    = std::lcm(n_embd_q, blck_size_perf);
+            const int64_t granularity_head = granularity_q / hparams.n_embd_head_k(il); // for tensors with one value per head
             if (std::regex_match(tensor_name, pattern_attn_sinks)) {
                 GGML_ASSERT(segments.size() == 1);
-                return {std::lcm(n_embd_q, blck_size_perf)/n_embd_q * n_gqa};
+                return {granularity_head};
             }
-
-            const int64_t granularity_q = std::lcm(n_embd_q, blck_size_perf);
             if (std::regex_match(tensor_name, pattern_q_weight) || std::regex_match(tensor_name, pattern_q_bias)) {
                 GGML_ASSERT(segments.size() == 1);
                 // some models have Q gate tensors, for those cases the granularity needs to be doubled:
@@ -660,6 +663,13 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
                 GGML_ASSERT(segments.size() == 1);
                 return {granularity_q};
             }
+            if (std::regex_match(tensor_name, pattern_attn_gate_weight)) {
+                GGML_ASSERT(segments.size() == 1);
+                if (tensor->ne[1] == hparams.n_head(il)) {
+                    return {granularity_head};
+                }
+                return {granularity_q};
+            }
 
             const int64_t granularity_kv = granularity_q / n_gqa;
             if (std::regex_match(tensor_name, pattern_kv_weight) ||
@@ -728,6 +738,16 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
         memset(split_state.ne, 0, sizeof(split_state.ne));
         split_state.nr[0] = 1;
         split_state.n_segments = 1;
+        if (split_state.axis == GGML_BACKEND_SPLIT_AXIS_PARTIAL) {
+            GGML_ASSERT(tc.tensor_axis_0 != tensor);
+            const ggml_backend_meta_split_state source_split_state = llama_meta_device_get_split_state(tc.tensor_axis_0, userdata);
+            GGML_ASSERT(source_split_state.axis >= 0 && source_split_state.axis < GGML_MAX_DIMS);
+            for (size_t j = 0; j < ud->n_devices; j++) {
+                for (size_t is = 0; is < source_split_state.n_segments; is++) {
+                    split_state.ne[j] += source_split_state.ne[is*ud->n_devices + j] * source_split_state.nr[is];
+                }
+            }
+        }
     }
     return split_state;
     GGML_UNUSED(userdata);

src/models/deepseek4.cpp

@@ -1,3 +1,4 @@
+#include "llama-hparams.h"
 #include "models.h"
 
 #include "llama-kv-cache-dsv4.h"
@@ -58,6 +59,7 @@ void llama_model_deepseek4::load_arch_hparams(llama_model_loader & ml) {
     if (n_compress_ratios < hparams.n_layer_all) {
         throw std::runtime_error("DeepSeek-V4 compress_ratios is shorter than block_count");
     }
+    GGML_ASSERT(n_compress_ratios <= LLAMA_MAX_LAYERS);
     ml.get_arr(LLM_KV_ATTENTION_COMPRESS_RATIOS, hparams.dsv4_compress_ratios);
 
     ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func);

src/models/glm4-moe.cpp

@@ -29,10 +29,19 @@ void llama_model_glm4_moe::load_arch_hparams(llama_model_loader & ml) {
     }
 }
 
-void llama_model_glm4_moe::load_arch_tensors(llama_model_loader &) {
+void llama_model_glm4_moe::load_arch_tensors(llama_model_loader & ml) {
     LLAMA_LOAD_LOCALS;
     const int64_t n_expert_shared = hparams.n_expert_shared;
 
+    const bool mtp_only = (hparams.n_layer_nextn > 0) && (ml.get_weight("blk.0.attn_norm.weight") == nullptr);
+    const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight";
+    const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr);
+    const int trunk_flags = mtp_only  ? TENSOR_NOT_REQUIRED : 0;
+    int       mtp_flags   = trunk_only ? TENSOR_NOT_REQUIRED : 0;
+
+    if (!ml.load_mtp) {
+        mtp_flags |= TENSOR_SKIP;
+    }
 
     GGML_ASSERT(hparams.n_expert > 0 && "n_expert must be > 0 for GLM4_MOE MoE layers");
     GGML_ASSERT(hparams.n_expert_used > 0 && "n_expert_used must be > 0 for GLM4_MOE MoE layers");
@@ -47,16 +56,9 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader &) {
         output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, TENSOR_DUPLICATED);
     }
 
-    // Load ALL tensors including NextN layer to satisfy total tensor count
-    // but only PROCESS up to last layer (skipping final NextN layer) in forward pass
     for (int i = 0; i < n_layer_all; ++i) {
-        int flags = 0;
-        if (i >= n_layer) {
-            // skip all tensors in the NextN layers
-            flags |= TENSOR_SKIP;
-        }
-
         auto & layer = layers[i];
+        const int flags = i < n_layer ? trunk_flags : mtp_flags;
 
         layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), { n_embd }, flags);
 
@@ -110,24 +112,186 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader &) {
             layer.ffn_up   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), { n_embd, n_ff }, flags);
         }
 
-        // NextN/MTP tensors (preserved but unused) - conditionally load for last nextn_predict_layers
+        // NextN/MTP tensors
         if (i >= n_layer) {
             layer.nextn.eh_proj          = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", i), { 2 * n_embd, n_embd }, flags);
             layer.nextn.enorm            = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), { n_embd }, flags);
             layer.nextn.hnorm            = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM, "weight", i), { n_embd }, flags);
 
             // Optional tensors
-            layer.nextn.embed_tokens     = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS, "weight", i), { n_embd, n_vocab }, flags | TENSOR_NOT_REQUIRED);
-            layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", i), { n_embd, n_vocab }, flags | TENSOR_NOT_REQUIRED);
-            layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", i), { n_embd }, flags | TENSOR_NOT_REQUIRED);
+            layer.nextn.embed_tokens     = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS, "weight", i), { n_embd, n_vocab }, TENSOR_NOT_REQUIRED | flags);
+            layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", i), { n_embd, n_vocab }, TENSOR_NOT_REQUIRED | flags);
+            layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", i), { n_embd }, TENSOR_NOT_REQUIRED | flags);
         }
     }
 }
 
 std::unique_ptr<llm_graph_context> llama_model_glm4_moe::build_arch_graph(const llm_graph_params & params) const {
+    if (params.gtype == LLM_GRAPH_TYPE_DECODER_MTP) {
+        return std::make_unique<graph_mtp>(*this, params);
+    }
     return std::make_unique<graph>(*this, params);
 }
 
+llama_model_glm4_moe::graph_mtp::graph_mtp(const llama_model & model, const llm_graph_params & params)
+    : llm_graph_context(params) {
+    GGML_ASSERT(hparams.n_layer_nextn > 0 && "GLM4_MOE MTP requires n_layer_nextn > 0");
+    GGML_ASSERT(hparams.n_layer_nextn == 1 && "GLM4_MOE MTP currently only supports a single MTP block");
+
+    const int64_t n_embd_head = hparams.n_embd_head_v();
+    GGML_ASSERT(n_embd_head == hparams.n_embd_head_k());
+
+    const int il = hparams.n_layer() + cparams.nextn_layer_offset;
+    GGML_ASSERT(cparams.nextn_layer_offset >= 0 &&
+                cparams.nextn_layer_offset < (int) hparams.n_layer_nextn &&
+                "nextn_layer_offset out of range [0, n_layer_nextn)");
+
+    const auto & layer = model.layers[il];
+
+    GGML_ASSERT(layer.nextn.eh_proj && "MTP block missing nextn.eh_proj");
+    GGML_ASSERT(layer.nextn.enorm   && "MTP block missing nextn.enorm");
+    GGML_ASSERT(layer.nextn.hnorm   && "MTP block missing nextn.hnorm");
+    GGML_ASSERT(layer.ffn_gate_inp  && "MTP block missing ffn_gate_inp");
+
+    auto inp = std::make_unique<llm_graph_input_embd_h>(hparams.n_embd);
+
+    inp->tokens = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_tokens);
+    ggml_set_input(inp->tokens);
+
+    inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, hparams.n_embd_inp(), n_tokens);
+    ggml_set_input(inp->embd);
+
+    ggml_tensor * tok_embd;
+    if (ubatch.token) {
+        ggml_tensor * tok_embd_w = layer.nextn.embed_tokens ? layer.nextn.embed_tokens : model.tok_embd;
+        tok_embd = ggml_get_rows(ctx0, tok_embd_w, inp->tokens);
+    } else {
+        tok_embd = inp->embd;
+    }
+    cb(tok_embd, "mtp_tok_embd", il);
+
+    inp->h = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, hparams.n_embd, n_tokens);
+    ggml_set_input(inp->h);
+    ggml_set_name(inp->h, "mtp_h_input");
+
+    ggml_tensor * h_embd = inp->h;
+
+    res->add_input(std::move(inp));
+
+    ggml_tensor * inp_pos     = build_inp_pos();
+    ggml_tensor * inp_out_ids = build_inp_out_ids();
+
+    auto * inp_attn = build_attn_inp_kv();
+
+    ggml_tensor * h_norm = build_norm(h_embd, layer.nextn.hnorm, nullptr, LLM_NORM_RMS, il);
+    cb(h_norm, "mtp_hnorm", il);
+
+    ggml_tensor * e_norm = build_norm(tok_embd, layer.nextn.enorm, nullptr, LLM_NORM_RMS, il);
+    cb(e_norm, "mtp_enorm", il);
+
+    ggml_tensor * concat = ggml_concat(ctx0, e_norm, h_norm, 0);
+    cb(concat, "mtp_concat", il);
+
+    ggml_tensor * cur = build_lora_mm(layer.nextn.eh_proj, concat, layer.nextn.eh_proj_s);
+    cb(cur, "mtp_eh_proj", il);
+
+    ggml_tensor * inpSA = cur;
+
+    cur = build_norm(cur, layer.attn_norm, nullptr, LLM_NORM_RMS, il);
+    cb(cur, "mtp_attn_norm", il);
+
+    auto [Qcur, Kcur, Vcur] = build_qkv(layer, cur,
+            n_embd_head, n_head, n_head_kv, il);
+
+    if (layer.attn_q_norm) {
+        Qcur = build_norm(Qcur, layer.attn_q_norm, nullptr, LLM_NORM_RMS, il);
+        cb(Qcur, "mtp_Qcur_normed", il);
+    }
+    if (layer.attn_k_norm) {
+        Kcur = build_norm(Kcur, layer.attn_k_norm, nullptr, LLM_NORM_RMS, il);
+        cb(Kcur, "mtp_Kcur_normed", il);
+    }
+
+    Qcur = ggml_rope_ext(ctx0, Qcur, inp_pos, nullptr, n_rot,
+            rope_type, n_ctx_orig, freq_base, freq_scale,
+            ext_factor, attn_factor, beta_fast, beta_slow);
+
+    Kcur = ggml_rope_ext(ctx0, Kcur, inp_pos, nullptr, n_rot,
+            rope_type, n_ctx_orig, freq_base, freq_scale,
+            ext_factor, attn_factor, beta_fast, beta_slow);
+
+    cb(Qcur, "mtp_Qcur", il);
+    cb(Kcur, "mtp_Kcur", il);
+    cb(Vcur, "mtp_Vcur", il);
+
+    cur = build_attn(inp_attn,
+            layer.wo, nullptr, layer.wo_s,
+            Qcur, Kcur, Vcur, nullptr, nullptr, nullptr,
+            1.0f / sqrtf(float(n_embd_head)), il);
+    cb(cur, "mtp_attn_out", il);
+
+    ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
+    cb(ffn_inp, "mtp_ffn_inp", il);
+
+    cur = build_norm(ffn_inp, layer.attn_post_norm, nullptr, LLM_NORM_RMS, il);
+    cb(cur, "mtp_post_attn_norm", il);
+
+    ggml_tensor * routed_out = build_moe_ffn(cur,
+            layer.ffn_gate_inp,
+            layer.ffn_up_exps,
+            layer.ffn_gate_exps,
+            layer.ffn_down_exps,
+            layer.ffn_exp_probs_b,
+            n_expert, n_expert_used,
+            LLM_FFN_SILU, hparams.expert_weights_norm,
+            hparams.expert_weights_scale,
+            (llama_expert_gating_func_type) hparams.expert_gating_func,
+            il);
+    cb(routed_out, "mtp_ffn_moe_out", il);
+
+    ggml_tensor * shared_out = build_ffn(cur,
+            layer.ffn_up_shexp,   nullptr, nullptr,
+            layer.ffn_gate_shexp, nullptr, nullptr,
+            layer.ffn_down_shexp, nullptr, nullptr,
+            nullptr,
+            LLM_FFN_SILU, LLM_FFN_PAR, il);
+    cb(shared_out, "mtp_ffn_shexp_out", il);
+
+    cur = ggml_add(ctx0, routed_out, shared_out);
+    cb(cur, "mtp_ffn_out", il);
+
+    cur = ggml_add(ctx0, cur, ffn_inp);
+    cb(cur, "mtp_post_ffn", il);
+
+    ggml_tensor * head_norm_w = layer.nextn.shared_head_norm
+            ? layer.nextn.shared_head_norm
+            : model.output_norm;
+    GGML_ASSERT(head_norm_w && "GLM4_MOE MTP: missing both nextn.shared_head_norm and output_norm");
+
+    cur = build_norm(cur, head_norm_w, nullptr, LLM_NORM_RMS, -1);
+    cb(cur, "h_nextn", -1);
+    res->t_h_nextn = cur;
+
+    if (inp_out_ids) {
+        cur = ggml_get_rows(ctx0, cur, inp_out_ids);
+    }
+    cb(cur, "mtp_shared_head_norm", -1);
+
+    ggml_tensor * head_w = layer.nextn.shared_head_head
+            ? layer.nextn.shared_head_head
+            : model.output;
+    ggml_tensor * head_s = layer.nextn.shared_head_head
+            ? layer.nextn.shared_head_head_s
+            : model.output_s;
+    GGML_ASSERT(head_w && "GLM4_MOE MTP: missing LM head (nextn.shared_head_head or model.output)");
+
+    cur = build_lora_mm(head_w, cur, head_s);
+    cb(cur, "result_output", -1);
+
+    res->t_logits = cur;
+    ggml_build_forward_expand(gf, cur);
+}
+
 llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) {
     const int64_t n_embd_head = hparams.n_embd_head_v();
 
@@ -154,8 +318,7 @@ llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_pa
 
     ggml_tensor * inp_out_ids = build_inp_out_ids();
 
-    // Only process up to last layer (skip final NextN layer)
-    // Final layer tensors are loaded but not processed in forward pass
+    // NextN layers are processed by graph_mtp.
     for (int il = 0; il < n_layer; ++il) {
         ggml_tensor * inpSA = inpL;
 
@@ -205,7 +368,7 @@ llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_pa
                     model.layers[il].wo, NULL, model.layers[il].wo_s,
                     Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f/sqrtf(float(n_embd_head)), il);
         }
-        if (il == n_layer - 1 && inp_out_ids) {
+        if (il == n_layer - 1 && inp_out_ids && (!cparams.embeddings_nextn || cparams.embeddings_nextn_masked)) {
             cur   = ggml_get_rows(ctx0, cur, inp_out_ids);
             inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
         }
@@ -265,6 +428,13 @@ llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_pa
     cur = inpL;
     cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1);
 
+    cb(cur, "h_nextn", -1);
+    res->t_h_nextn = cur;
+
+    if (cparams.embeddings_nextn && !cparams.embeddings_nextn_masked && inp_out_ids) {
+        cur = ggml_get_rows(ctx0, cur, inp_out_ids);
+    }
+
     cb(cur, "result_norm", -1);
     res->t_embd = cur;
 

src/models/models.h

@@ -1412,6 +1412,10 @@ struct llama_model_glm4_moe : public llama_model_base {
         graph(const llama_model & model, const llm_graph_params & params);
     };
 
+    struct graph_mtp : public llm_graph_context {
+        graph_mtp(const llama_model & model, const llm_graph_params & params);
+    };
+
     std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
 };
 

tests/CMakeLists.txt

@@ -228,6 +228,15 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
     set_tests_properties(test-recurrent-state-rollback-nemotron-h PROPERTIES
         FIXTURES_REQUIRED generate-models
     )
+    llama_test(
+        test-recurrent-state-rollback
+        NAME test-recurrent-state-rollback-dsv4
+        LABEL main
+        ARGS -m "${MODEL_DIR}/deepseek4-moe.gguf"
+    )
+    set_tests_properties(test-recurrent-state-rollback-dsv4 PROPERTIES
+        FIXTURES_REQUIRED generate-models
+    )
 endif()
 
 llama_build_and_test(test-chat-peg-parser.cpp peg-parser/simple-tokenize.cpp)
@@ -235,6 +244,8 @@ llama_build_and_test(test-jinja.cpp)
 llama_test(test-jinja NAME test-jinja-py ARGS -py LABEL python)
 llama_build_and_test(test-chat-auto-parser.cpp WORKING_DIRECTORY ${PROJECT_SOURCE_DIR})
 llama_build_and_test(test-chat-template.cpp)
+# debug tool for chat template differential analysis (not registered as a test, run it manually)
+llama_build(test-chat-analysis.cpp)
 llama_build_and_test(test-log.cpp)
 llama_build_and_test(
     test-peg-parser.cpp

tests/test-chat-analysis.cpp

@@ -84,11 +84,12 @@ static std::string read_file(const std::string & path) {
 }
 
 static void print_usage(const char * program_name) {
-    LOG_ERR("Usage: %s [options]\n", program_name);
+    LOG_ERR("Debug the auto-parser's differential analysis: render a template with/without tools, reasoning, etc. and show the diffs.\n");
+    LOG_ERR("\nUsage: %s [options]\n", program_name);
     LOG_ERR("\nOptions:\n");
     LOG_ERR("  --template <name>       Analyze specific template from test suite (e.g., 'deepseek' or 'DeepSeek-V3.1')\n");
     LOG_ERR("  --template-file <path>  Analyze custom template file\n");
-    LOG_ERR("  --all                   Analyze all templates from test suite\n");
+    LOG_ERR("  --all                   Analyze all templates from test suite (default when no arguments are given)\n");
     LOG_ERR("\nExamples:\n");
     LOG_ERR("  %s --all\n", program_name);
     LOG_ERR("  %s --template deepseek\n", program_name);
@@ -97,14 +98,17 @@ static void print_usage(const char * program_name) {
 
 static bool parse_options(int argc, char ** argv, analysis_options & opts) {
     if (argc < 2) {
-        print_usage(argv[0]);
-        return false;
+        // default mode: analyze all templates from the test suite
+        opts.analyze_all = true;
     }
 
     for (int i = 1; i < argc; ++i) {
         std::string arg = argv[i];
 
-        if (arg == "--all") {
+        if (arg == "-h" || arg == "--help") {
+            print_usage(argv[0]);
+            return false;
+        } else if (arg == "--all") {
             opts.analyze_all = true;
         } else if (arg == "--template") {
             if (i + 1 >= argc) {

tests/test-chat-auto-parser.cpp

@@ -2,11 +2,18 @@
 #include "chat-auto-parser.h"
 #include "chat-peg-parser.h"
 #include "chat.h"
+#include "gguf.h"
+#include "jinja/runtime.h"
+#include "log.h"
 #include "peg-parser.h"
 #include "testing.h"
 
+#include <cstdlib>
+#include <filesystem>
 #include <fstream>
 #include <iostream>
+#include <iterator>
+#include <optional>
 #include <sstream>
 #include <string>
 
@@ -94,11 +101,447 @@ static void test_bailing_v3_tool_format(testing & t);
 
 static void test_role_markers_all_templates(testing & t);
 
+static json build_tools_definition();
+
+//
+// debug mode: analyze a single template and dump the generated parser and grammar
+//
+
+enum class output_mode {
+    ANALYSIS,  // Only output analysis results (default)
+    TEMPLATE,  // Only output rendered template
+    BOTH       // Output both
+};
+
+enum class input_message_type {
+    NONE,                    // Don't render any message scenarios (only analysis)
+    CONTENT_ONLY,            // Simple assistant message with content
+    REASONING_CONTENT,       // Message with reasoning_content + content
+    TOOL_CALL_ONLY,          // Message with tool_calls only
+    CONTENT_TOOL_CALL,       // Message with content + tool_calls
+    REASONING_TOOL_CALL,     // Message with reasoning_content + tool_calls
+    CONTENT_FAKE_TOOL_CALL,  // Message with content but no actual tool_calls (for testing)
+    ALL                      // Render all scenarios
+};
+
+struct debug_options {
+    std::string        template_path;
+    bool               with_tools        = true;
+    bool               generation_prompt = true;
+    bool               enable_reasoning  = true;
+    bool               debug_jinja       = false;
+    bool               force_tool_call   = false;
+    bool               parallel_tool_calls = true;
+    output_mode        mode              = output_mode::BOTH;
+    input_message_type input_message     = input_message_type::NONE;
+};
+
+static std::string read_file(const std::string & path) {
+    std::ifstream fin(path, std::ios::binary);
+    if (!fin.is_open()) {
+        throw std::runtime_error("Could not open file: " + path);
+    }
+    std::ostringstream buf;
+    buf << fin.rdbuf();
+    return buf.str();
+}
+
+static std::string read_gguf_chat_template(const std::string & path) {
+    struct gguf_init_params params = { /*no_alloc =*/true,  // We only need metadata, not tensor data
+                                       /*ctx=*/nullptr };
+
+    struct gguf_context * ctx = gguf_init_from_file(path.c_str(), params);
+    if (ctx == nullptr) {
+        throw std::runtime_error("Could not open GGUF file: " + path);
+    }
+
+    const char * key    = "tokenizer.chat_template";
+    int64_t      key_id = gguf_find_key(ctx, key);
+
+    if (key_id == -1) {
+        gguf_free(ctx);
+        throw std::runtime_error("GGUF file does not contain chat template key: " + std::string(key));
+    }
+
+    const char * template_str = gguf_get_val_str(ctx, key_id);
+    if (template_str == nullptr) {
+        gguf_free(ctx);
+        throw std::runtime_error("GGUF file contains chat template key but value is null");
+    }
+
+    std::string result = template_str;
+    gguf_free(ctx);
+    return result;
+}
+
+static void print_usage(const char * program_name) {
+    LOG_ERR("Test the chat template auto-parser; also usable as a debug tool that shows the generated PEG parser, GBNF grammar and triggers for a given template.\n");
+    LOG_ERR("\nUsage: %s [filter_regex]                       run the automated tests (default)\n", program_name);
+    LOG_ERR("       %s <template_or_gguf_path> [options]     debug a single template\n", program_name);
+    LOG_ERR("\nDebug mode options:\n");
+    LOG_ERR("  --no-tools              Disable tool definitions\n");
+    LOG_ERR("  --force-tool-call       Set tool calls to forced\n");
+    LOG_ERR("  --parallel-tool-calls=0|1 Set parallel_tool_calls (default: 1)\n");
+    LOG_ERR("  --generation-prompt=0|1 Set add_generation_prompt (default: 1)\n");
+    LOG_ERR("  --enable-reasoning=0|1  Enable reasoning parsing (default: 1)\n");
+    LOG_ERR("  --output=MODE           Output mode: analysis, template, both (default: both)\n");
+    LOG_ERR("  --debug-jinja           Enable Jinja fine-grained debug\n");
+    LOG_ERR("  --input-message=TYPE    Message type to render:\n");
+    LOG_ERR("                          content_only, reasoning_content, tool_call_only,\n");
+    LOG_ERR("                          content_tool_call, reasoning_tool_call,\n");
+    LOG_ERR("                          content_fake_tool_call, all\n");
+    LOG_ERR("\nExamples:\n");
+    LOG_ERR("  %s template.jinja --input-message=all --generation-prompt=1\n", program_name);
+    LOG_ERR("  %s template.jinja --output=template --input-message=tool_call_only\n", program_name);
+}
+
+static bool parse_bool_option(const std::string & value) {
+    return value == "1" || value == "true" || value == "yes";
+}
+
+static bool parse_debug_options(int argc, char ** argv, debug_options & opts) {
+    opts.template_path = argv[1];
+
+    for (int i = 2; i < argc; ++i) {
+        std::string arg = argv[i];
+
+        if (arg == "--force-tool-call") {
+            opts.force_tool_call = true;
+        } else if (arg == "--debug-jinja") {
+            opts.debug_jinja = true;
+        } else if (arg == "--no-tools") {
+            opts.with_tools = false;
+        } else if (arg.rfind("--parallel-tool-calls=", 0) == 0) {
+            opts.parallel_tool_calls = parse_bool_option(arg.substr(22));
+        } else if (arg.rfind("--generation-prompt=", 0) == 0) {
+            opts.generation_prompt = parse_bool_option(arg.substr(20));
+        } else if (arg.rfind("--enable-reasoning=", 0) == 0) {
+            opts.enable_reasoning = parse_bool_option(arg.substr(19));
+        } else if (arg.rfind("--output=", 0) == 0) {
+            std::string mode = arg.substr(9);
+            if (mode == "analysis") {
+                opts.mode = output_mode::ANALYSIS;
+            } else if (mode == "template") {
+                opts.mode = output_mode::TEMPLATE;
+            } else if (mode == "both") {
+                opts.mode = output_mode::BOTH;
+            } else {
+                LOG_ERR("Unknown output mode: %s\n", mode.c_str());
+                return false;
+            }
+        } else if (arg.rfind("--input-message=", 0) == 0) {
+            std::string type = arg.substr(16);
+            if (type == "content_only") {
+                opts.input_message = input_message_type::CONTENT_ONLY;
+            } else if (type == "reasoning_content") {
+                opts.input_message = input_message_type::REASONING_CONTENT;
+            } else if (type == "tool_call_only") {
+                opts.input_message = input_message_type::TOOL_CALL_ONLY;
+            } else if (type == "content_tool_call") {
+                opts.input_message = input_message_type::CONTENT_TOOL_CALL;
+            } else if (type == "reasoning_tool_call") {
+                opts.input_message = input_message_type::REASONING_TOOL_CALL;
+            } else if (type == "content_fake_tool_call") {
+                opts.input_message = input_message_type::CONTENT_FAKE_TOOL_CALL;
+            } else if (type == "all") {
+                opts.input_message = input_message_type::ALL;
+            } else {
+                LOG_ERR("Unknown input message type: %s\n", type.c_str());
+                return false;
+            }
+        } else {
+            LOG_ERR("Unknown option: %s\n", arg.c_str());
+            print_usage(argv[0]);
+            return false;
+        }
+    }
+
+    return true;
+}
+
+static json build_debug_user_message() {
+    return json{
+        { "role",    "user"                               },
+        { "content", "Hello, please help me with a task." }
+    };
+}
+
+static json build_content_only_message() {
+    return json{
+        { "role",    "assistant"                                   },
+        { "content", "Hello! I'm here to help you with your task." }
+    };
+}
+
+static json build_reasoning_content_message() {
+    return json{
+        { "role",              "assistant"                                                               },
+        { "content",           "Hello! I'm here to help you with your task."                             },
+        { "reasoning_content", "The user is greeting me and asking for help. I should respond politely." }
+    };
+}
+
+static json build_tool_call_only_message() {
+    return json{
+        { "role",       "assistant"      },
+        { "content",    nullptr          },
+        { "tool_calls",
+         json::array({ json{
+              { "type", "function" },
+              { "function", json{ { "name", "test_function_name" },
+                                  { "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } },
+              { "id", "123456789" } } }) }
+    };
+}
+
+static json build_content_tool_call_message() {
+    return json{
+        { "role",       "assistant"                                                                              },
+        { "content",    "I'll help you by calling a function."                                                   },
+        { "tool_calls",
+         json::array({ json{
+              { "type", "function" },
+              { "function",
+                json{ { "name", "test_function_name" },
+                      { "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } } } }) }
+    };
+}
+
+static json build_reasoning_tool_call_message() {
+    return json{
+        { "role",              "assistant"                                                                       },
+        { "content",           nullptr                                                                           },
+        { "reasoning_content", "I need to call a function to help with this task."                               },
+        { "tool_calls",
+         json::array({ json{
+              { "type", "function" },
+              { "function",
+                json{ { "name", "test_function_name" },
+                      { "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } } } }) }
+    };
+}
+
+static json build_content_fake_tool_call_message() {
+    // This message has content but NO tool_calls field
+    // It's used to test if a template renders tool definitions but not tool calls
+    return json{
+        { "role",    "assistant"                            },
+        { "content", "I'll help you by calling a function." }
+    };
+}
+
+static void render_scenario(const common_chat_template & tmpl,
+                            const std::string &          scenario_name,
+                            const json &                 messages,
+                            const json &                 tools,
+                            bool                         add_generation_prompt,
+                            bool                         enable_thinking) {
+    LOG_ERR("\n=== Scenario: %s ===\n", scenario_name.c_str());
+    LOG_ERR("add_generation_prompt: %s, enable_thinking: %s\n", add_generation_prompt ? "true" : "false",
+            enable_thinking ? "true" : "false");
+
+    // When add_generation_prompt is true, add a trailing user message to trigger the prompt
+    json final_messages = messages;
+    if (add_generation_prompt && !messages.empty() && messages.back().value("role", "") == "assistant") {
+        final_messages.push_back(json{
+            { "role",    "user"                                       },
+            { "content", "Now please continue with another response." }
+        });
+    }
+
+    LOG_ERR("Messages:\n%s\n", final_messages.dump(2).c_str());
+
+    try {
+        generation_params inputs;
+        inputs.messages                         = final_messages;
+        inputs.add_generation_prompt            = add_generation_prompt;
+        inputs.extra_context["enable_thinking"] = enable_thinking;
+
+        if (!tools.is_null() && tools.is_array() && !tools.empty()) {
+            inputs.tools = tools;
+        }
+
+        std::string output = common_chat_template_direct_apply(tmpl, inputs);
+
+        LOG_ERR("\n--- Rendered Output ---\n");
+        LOG_ERR("%s\n", output.c_str());
+        LOG_ERR("--- End Output (length: %zu) ---\n", output.length());
+    } catch (const std::exception & e) {
+        LOG_ERR("Rendering failed: %s\n", e.what());
+    }
+}
+
+static void render_all_scenarios(const common_chat_template & tmpl,
+                                 const json &                 tools,
+                                 bool                         add_generation_prompt,
+                                 bool                         enable_thinking,
+                                 input_message_type           message_type) {
+    json user_msg = build_debug_user_message();
+
+    auto render_if = [&](input_message_type type, const std::string & name, const json & assistant_msg) {
+        if (message_type == input_message_type::ALL || message_type == type) {
+            json messages = json::array({ user_msg, assistant_msg });
+            render_scenario(tmpl, name, messages, tools, add_generation_prompt, enable_thinking);
+        }
+    };
+
+    render_if(input_message_type::CONTENT_ONLY, "content_only", build_content_only_message());
+    render_if(input_message_type::REASONING_CONTENT, "reasoning_content", build_reasoning_content_message());
+    render_if(input_message_type::TOOL_CALL_ONLY, "tool_call_only", build_tool_call_only_message());
+    render_if(input_message_type::CONTENT_TOOL_CALL, "content_tool_call", build_content_tool_call_message());
+    render_if(input_message_type::REASONING_TOOL_CALL, "reasoning_tool_call", build_reasoning_tool_call_message());
+    render_if(input_message_type::CONTENT_FAKE_TOOL_CALL, "content_fake_tool_call",
+              build_content_fake_tool_call_message());
+
+    // Also render with add_generation_prompt=true to show the prompt ending
+    if (message_type == input_message_type::ALL) {
+        LOG_ERR("\n\n=== Generation Prompt Scenarios (add_generation_prompt=true) ===\n");
+
+        json prompt_messages = json::array({ user_msg });
+        render_scenario(tmpl, "generation_prompt_only", prompt_messages, tools, true, enable_thinking);
+
+        // With enable_thinking toggled
+        render_scenario(tmpl, "generation_prompt_thinking_disabled", prompt_messages, tools, true, false);
+    }
+}
+
+static generation_params prepare_debug_params(const debug_options & opts, const json & tools) {
+    generation_params params;
+    params.messages         = json::array({ build_debug_user_message() });
+    params.reasoning_format = opts.enable_reasoning ? COMMON_REASONING_FORMAT_DEEPSEEK : COMMON_REASONING_FORMAT_NONE;
+    params.enable_thinking  = opts.enable_reasoning;
+    params.add_generation_prompt = opts.generation_prompt;
+
+    if (opts.with_tools) {
+        params.tools       = tools;
+        params.tool_choice = opts.force_tool_call ? COMMON_CHAT_TOOL_CHOICE_REQUIRED : COMMON_CHAT_TOOL_CHOICE_AUTO;
+    } else {
+        params.tools       = json();
+        params.tool_choice = COMMON_CHAT_TOOL_CHOICE_NONE;
+    }
+    params.parallel_tool_calls = opts.parallel_tool_calls;
+    return params;
+}
+
+static int debug_single_template(const debug_options & opts) {
+    std::string template_source;
+    try {
+        // Check if the file is a GGUF file
+        if (opts.template_path.size() >= 5 &&
+            opts.template_path.compare(opts.template_path.size() - 5, 5, ".gguf") == 0) {
+            template_source = read_gguf_chat_template(opts.template_path);
+        } else {
+            template_source = read_file(opts.template_path);
+        }
+    } catch (const std::exception & e) {
+        LOG_ERR("Error reading template: %s\n", e.what());
+        return 1;
+    }
+
+    LOG_ERR("Analyzing template: %s\n", opts.template_path.c_str());
+    LOG_ERR("Options: with_tools=%s, generation_prompt=%s, enable_reasoning=%s\n", opts.with_tools ? "true" : "false",
+            opts.generation_prompt ? "true" : "false", opts.enable_reasoning ? "true" : "false");
+
+    try {
+        common_chat_template chat_template(template_source, "", "");
+
+        json tools = opts.with_tools ? build_tools_definition() : json();
+
+        generation_params params = prepare_debug_params(opts, tools);
+        common_chat_params            parser_data;
+        if (std::optional<common_chat_params> spec_tmpl =
+                common_chat_try_specialized_template(chat_template, template_source, params)) {
+            LOG_ERR("\n");
+            LOG_ERR("This template uses a specialized parser, analysis results will not be available.\n");
+            parser_data = *spec_tmpl;
+        } else {
+            // Render template scenarios if requested
+            if (opts.input_message != input_message_type::NONE &&
+                (opts.mode == output_mode::TEMPLATE || opts.mode == output_mode::BOTH)) {
+                LOG_ERR("\n");
+                LOG_ERR("================================================================================\n");
+                LOG_ERR("                         TEMPLATE RENDERING OUTPUT\n");
+                LOG_ERR("================================================================================\n");
+
+                render_all_scenarios(chat_template, tools, opts.generation_prompt, opts.enable_reasoning,
+                                     opts.input_message);
+            }
+
+            // Output analysis if requested
+            if (opts.mode == output_mode::ANALYSIS || opts.mode == output_mode::BOTH) {
+                LOG_ERR("\n");
+                LOG_ERR("================================================================================\n");
+                LOG_ERR("                           TEMPLATE ANALYSIS\n");
+                LOG_ERR("================================================================================\n");
+
+                struct autoparser analysis;
+                analysis.analyze_template(chat_template);
+
+                // Generate Parser
+                parser_data = peg_generator::generate_parser(chat_template, params, analysis);
+            }
+        }
+
+        if (!std::empty(parser_data.parser)) {
+            LOG_ERR("\n=== Generated Parser ===\n");
+            common_peg_arena arena;
+            arena.load(parser_data.parser);
+            LOG_ERR("%s\n", arena.dump(arena.root()).c_str());
+
+            LOG_ERR("\n=== Generated Grammar ===\n");
+            LOG_ERR("%s\n", parser_data.grammar.c_str());
+
+            LOG_ERR("\n=== Generated Lazy Grammar ===\n");
+            LOG_ERR("%d\n", parser_data.grammar_lazy);
+
+            LOG_ERR("\n=== Generated Grammar Triggers ===\n");
+            for (const common_grammar_trigger & cgt : parser_data.grammar_triggers) {
+                LOG_ERR("Token: %d | Type: %d | Value: %s\n", cgt.token, cgt.type, cgt.value.c_str());
+            }
+
+            LOG_ERR("\n=== Preserved Tokens ===\n");
+            for (const std::string & token : parser_data.preserved_tokens) {
+                LOG_ERR("  '%s'\n", token.c_str());
+            }
+        }
+    } catch (const std::exception & e) {
+        LOG_ERR("Analysis failed: %s\n", e.what());
+        return 1;
+    }
+
+    return 0;
+}
+
 int main(int argc, char * argv[]) {
+    if (argc > 1) {
+        std::string arg = argv[1];
+        if (arg == "-h" || arg == "--help") {
+            common_log_set_verbosity_thold(99);
+            print_usage(argv[0]);
+            return 0;
+        }
+
+        // debug mode: if the first argument is an existing file, analyze that template instead of running the automated tests
+        if (std::filesystem::is_regular_file(arg)) {
+            common_log_set_verbosity_thold(99);
+
+            debug_options opts;
+            if (!parse_debug_options(argc, argv, opts)) {
+                return 1;
+            }
+
+            if (opts.debug_jinja || std::getenv("LLAMA_DEBUG_JINJA") != nullptr) {
+                jinja::enable_debug(true);
+            }
+
+            return debug_single_template(opts);
+        }
+    }
+
     testing t(std::cout);
     t.verbose = true;
 
-    // usage: test-chat-auto-parser-helpers [filter_regex]
+    // usage: test-chat-auto-parser [filter_regex]
 
     if (argc > 1) {
         t.set_filter(argv[1]);

tests/test-chat-template.cpp

@@ -28,6 +28,8 @@ static void run_multiple(const std::string& dir_path, bool stop_on_first_failure
 static void run_single(const std::string& contents, json input, bool use_common = false, bool dump_prog = false, const std::string & output_path = "");
 
 static std::string HELP = R"(
+Test the Jinja engine by rendering chat templates and comparing the output against expected results.
+
 Usage: test-chat-template [OPTIONS] PATH_TO_TEMPLATE
 Options:
   -h, --help               Show this help message and exit.

tests/test-llama-archs.cpp

@@ -101,6 +101,15 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
         n_head = 1;
         n_ff   = 96;
         n_layer = 22; // hparams.n_layer_kv_from_start = 20 is hardcoded
+    } else if (arch == LLM_ARCH_DEEPSEEK4) {
+        n_embd  = 128;
+        n_head  = 1;
+        n_ff    = 192;
+        n_layer = 3; // uncompressed + csa + hca, one layer of each ratio kind
+    } else if (arch == LLM_ARCH_STEP35 || arch == LLM_ARCH_LAGUNA) {
+        n_embd = 160; // exercise per-head tensor split granularity with head size 80
+    } else if (arch == LLM_ARCH_QWEN3 || arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_AFMOE) {
+        n_head = 4;
     } else if (arch == LLM_ARCH_DEEPSEEK2
             || arch == LLM_ARCH_DEEPSEEK32
             || arch == LLM_ARCH_GLM_DSA
@@ -120,6 +129,12 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
         n_vocab = 4096; // must be >= the hard-coded codec head size (3072)
     }
 
+    uint32_t n_head_kv = n_head;
+    if (arch == LLM_ARCH_QWEN3) {
+        n_head_kv = 1; // MQA coverage
+    } else if (arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_AFMOE) {
+        n_head_kv = 2; // GQA coverage
+    }
     const uint32_t n_embd_head = n_embd / n_head;
 
     ms.add_kv(LLM_KV_GENERAL_ARCHITECTURE,      llm_arch_name(arch));
@@ -160,7 +175,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
         ms.add_kv(LLM_KV_ATTENTION_HEAD_COUNT_KV, n_head_per_layer);
     } else {
         ms.add_kv(LLM_KV_ATTENTION_HEAD_COUNT, n_head);
-        ms.add_kv(LLM_KV_ATTENTION_HEAD_COUNT_KV, n_head);
+        ms.add_kv(LLM_KV_ATTENTION_HEAD_COUNT_KV, n_head_kv);
     }
 
     ms.add_kv(LLM_KV_ATTENTION_MAX_ALIBI_BIAS, 8.0f);
@@ -193,6 +208,10 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
             }
             ms.add_kv(LLM_KV_ATTENTION_INDEXER_TYPES, indexer_types);
         }
+    } else if (arch == LLM_ARCH_DEEPSEEK4) {
+        ms.add_kv(LLM_KV_ATTENTION_KEY_LENGTH,       uint32_t(128));
+        ms.add_kv(LLM_KV_ATTENTION_VALUE_LENGTH,     uint32_t(128));
+        ms.add_kv(LLM_KV_ROPE_DIMENSION_COUNT,       uint32_t(64));
     } else if (arch == LLM_ARCH_MINIMAX_M3) {
         // partial rotary: n_rot must not exceed the indexer key length (64)
         ms.add_kv(LLM_KV_ROPE_DIMENSION_COUNT,       uint32_t(64));
@@ -229,6 +248,20 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
 
     // MSA requires one indexer head per GQA (KV) head, unlike the DSA archs where the
     // indexer head count is independent of the main attention head count.
+    if (arch == LLM_ARCH_DEEPSEEK4) {
+        ms.add_kv(LLM_KV_EXPERT_WEIGHTS_SCALE,                 2.5f);
+        ms.add_kv(LLM_KV_EXPERT_WEIGHTS_NORM,                  true);
+        ms.add_kv(LLM_KV_SWIGLU_CLAMP_EXP,                     7.0f);
+        ms.add_kv(LLM_KV_ATTENTION_OUTPUT_GROUP_COUNT,         uint32_t(1));
+        ms.add_kv(LLM_KV_ATTENTION_OUTPUT_LORA_RANK,           uint32_t(64));
+        ms.add_kv(LLM_KV_ATTENTION_COMPRESS_ROPE_FREQ_BASE,    10000.0f);
+        ms.add_kv(LLM_KV_HYPER_CONNECTION_COUNT,               uint32_t(4));
+        ms.add_kv(LLM_KV_HYPER_CONNECTION_SINKHORN_ITERATIONS, uint32_t(4));
+        ms.add_kv(LLM_KV_HYPER_CONNECTION_EPSILON,             1e-6f);
+        ms.add_kv(LLM_KV_HASH_LAYER_COUNT,                     uint32_t(0));
+        ms.add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS,            std::vector<uint32_t>({0, 4, 128}));
+    }
+
     ms.add_kv(LLM_KV_ATTENTION_INDEXER_HEAD_COUNT,   arch == LLM_ARCH_MINIMAX_M3 ? n_head : uint32_t(1));
     ms.add_kv(LLM_KV_ATTENTION_INDEXER_KEY_LENGTH,   uint32_t(64));
     ms.add_kv(LLM_KV_ATTENTION_INDEXER_TOP_K,        uint32_t(8));
@@ -247,7 +280,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
         ms.add_kv(LLM_KV_EXPERT_COUNT,               uint32_t(2));
         ms.add_kv(LLM_KV_EXPERT_USED_COUNT,          uint32_t(1));
         ms.add_kv(LLM_KV_EXPERT_SHARED_COUNT,        uint32_t(1));
-        ms.add_kv(LLM_KV_EXPERT_GATING_FUNC,         uint32_t(2)); // sigmoid
+        ms.add_kv(LLM_KV_EXPERT_GATING_FUNC,         arch == LLM_ARCH_DEEPSEEK4 ? uint32_t(4) : uint32_t(2)); // sqrtsoftplus : sigmoid
         ms.add_kv(LLM_KV_EXPERT_GROUP_SCALE,         1.0f);
         ms.add_kv(LLM_KV_EXPERTS_PER_GROUP,          uint32_t(1));
     }
@@ -385,6 +418,7 @@ static bool moe_mandatory(const llm_arch arch) {
         case LLM_ARCH_DEEPSEEK2:
         case LLM_ARCH_DEEPSEEK32:
         case LLM_ARCH_DOTS3NOTE:
+        case LLM_ARCH_DEEPSEEK4:
         case LLM_ARCH_GLM4_MOE:
         case LLM_ARCH_GLM_DSA:
         case LLM_ARCH_EXAONE_MOE:
@@ -470,9 +504,6 @@ static bool arch_supported(const llm_arch arch) {
     if (arch == LLM_ARCH_DEEPSEEK2OCR) {
         return false;
     }
-    if (arch == LLM_ARCH_DEEPSEEK4) {
-        return false;
-    }
 
     // FIXME: these hit scheduler/view-backed-output issues with WebGPU on CI.
 #ifdef GGML_USE_WEBGPU

tests/test-recurrent-state-rollback.cpp

@@ -35,6 +35,178 @@ static bool decode_one(llama_context * ctx, llama_token tok, llama_pos pos) {
     return ok;
 }
 
+// Roll back multiple sequences, then replay them in a single batch whose
+// per-seq token count exceeds n_ubatch: each seq's replay spans several
+// ubatches while its rollback restore is still pending. Compared against a
+// reference context that never advanced past the rollback point and decodes
+// the identical replay batch.
+static bool test_multi_seq_split_replay(const common_params & params, llama_model * model, const int n_vocab) {
+    constexpr uint32_t  n_seqs     = 2;
+    constexpr uint32_t  n_ubatch   = 16;
+    constexpr uint32_t  n_prompt   = 19;
+    constexpr uint32_t  n_rollback = 3;
+    constexpr uint32_t  n_replay   = 40; // > n_ubatch so each seq spans multiple ubatches
+    constexpr llama_pos p0         = n_prompt - n_rollback;
+
+    const auto make_ctx_multi = [&]() {
+        auto cparams = common_context_params_to_llama(params);
+        cparams.n_seq_max  = n_seqs;
+        cparams.n_rs_seq   = 8;
+        cparams.n_ctx      = 256;
+        cparams.n_batch    = 256;
+        cparams.n_ubatch   = n_ubatch;
+        cparams.kv_unified = false;
+        return llama_init_from_model(model, cparams);
+    };
+
+    llama_context * ctx_roll = make_ctx_multi();
+    llama_context * ctx_ref  = make_ctx_multi();
+    if (ctx_roll == nullptr || ctx_ref == nullptr) {
+        fprintf(stderr, "%s : failed to init multi-seq contexts\n", __func__);
+        return false;
+    }
+
+    const auto cleanup = [&]() {
+        llama_free(ctx_roll);
+        llama_free(ctx_ref);
+    };
+
+    if (llama_n_rs_seq(ctx_roll) < n_rollback) {
+        fprintf(stderr, "%s : skipping because n_rs_seq is too small\n", __func__);
+        cleanup();
+        return true;
+    }
+
+    const auto tok = [&](uint32_t seq, llama_pos pos) {
+        return (llama_token) ((7*(uint32_t) pos + 31*seq + 1) % (uint32_t) n_vocab);
+    };
+
+    bool ok = true;
+
+    // both contexts decode the identical [0, p0) prefill; only ctx_roll decodes
+    // the tail, which is then rolled back so its restore is pending at replay
+    for (uint32_t s = 0; s < n_seqs && ok; ++s) {
+        llama_batch batch = llama_batch_init(n_prompt, 0, 1);
+        for (llama_pos pos = 0; pos < (llama_pos) p0; ++pos) {
+            common_batch_add(batch, tok(s, pos), pos, { (llama_seq_id) s }, false);
+        }
+        ok = ok && llama_decode(ctx_roll, batch) == 0;
+        ok = ok && llama_decode(ctx_ref,  batch) == 0;
+
+        common_batch_clear(batch);
+        for (llama_pos pos = p0; pos < (llama_pos) n_prompt; ++pos) {
+            common_batch_add(batch, tok(s, pos), pos, { (llama_seq_id) s }, false);
+        }
+        ok = ok && llama_decode(ctx_roll, batch) == 0;
+        llama_batch_free(batch);
+
+        ok = ok && llama_memory_seq_rm(llama_get_memory(ctx_roll), (llama_seq_id) s, p0, -1);
+
+        // a second partial removal while one is pending must be refused
+        ok = ok && !llama_memory_seq_rm(llama_get_memory(ctx_roll), (llama_seq_id) s, p0 - 1, -1);
+    }
+    if (!ok) {
+        fprintf(stderr, "%s : multi-seq prefill/rollback failed\n", __func__);
+        cleanup();
+        return false;
+    }
+
+    llama_batch batch = llama_batch_init(n_seqs*n_replay, 0, 1);
+    for (uint32_t s = 0; s < n_seqs; ++s) {
+        for (uint32_t i = 0; i < n_replay; ++i) {
+            const llama_pos pos = p0 + (llama_pos) i;
+            common_batch_add(batch, tok(s, pos), pos, { (llama_seq_id) s }, true);
+        }
+    }
+    ok = llama_decode(ctx_roll, batch) == 0;
+    ok = ok && llama_decode(ctx_ref, batch) == 0;
+    llama_batch_free(batch);
+    if (!ok) {
+        fprintf(stderr, "%s : multi-seq replay decode failed\n", __func__);
+        cleanup();
+        return false;
+    }
+
+    // identical ubatch shapes from bit-exact states: a correct implementation
+    // matches bitwise, so eps only allows backend scheduling noise
+    constexpr float eps = 1e-7f;
+
+    float    diff_max  = 0.0f;
+    uint32_t seq_first = 0;
+    int32_t  pos_first = -1;
+    for (uint32_t i = 0; i < n_seqs*n_replay; ++i) {
+        const float * l_roll = llama_get_logits_ith(ctx_roll, i);
+        const float * l_ref  = llama_get_logits_ith(ctx_ref,  i);
+        if (l_roll == nullptr || l_ref == nullptr) {
+            fprintf(stderr, "%s : missing multi-seq logits at index %u\n", __func__, i);
+            cleanup();
+            return false;
+        }
+        for (int t = 0; t < n_vocab; ++t) {
+            const float diff = std::fabs(l_roll[t] - l_ref[t]);
+            if (diff > eps && pos_first < 0) {
+                seq_first = i/n_replay;
+                pos_first = p0 + (int32_t) (i%n_replay);
+            }
+            diff_max = std::max(diff_max, diff);
+        }
+    }
+
+    if (diff_max > eps) {
+        fprintf(stderr, "%s : multi-seq split replay logits mismatch (max diff %g, first at seq %u pos %d)\n",
+                __func__, (double) diff_max, seq_first, pos_first);
+        cleanup();
+        return false;
+    }
+
+    fprintf(stderr, "%s : multi-seq split replay matched (max diff %g)\n", __func__, (double) diff_max);
+
+    // seq-1-only decodes must be independent of seq 0's content: diverge seq 0
+    // in ctx_ref only, then compare identical seq-1-only continuations bitwise
+    constexpr uint32_t n_tail = 4;
+
+    {
+        llama_batch batch_tail = llama_batch_init(n_tail, 0, 1);
+        for (uint32_t i = 0; i < n_tail; ++i) {
+            const llama_pos pos = p0 + (llama_pos) (n_replay + i);
+            common_batch_add(batch_tail, tok(0, pos + 7), pos, { 0 }, false);
+        }
+        ok = llama_decode(ctx_ref, batch_tail) == 0;
+        llama_batch_free(batch_tail);
+    }
+
+    float diff_tail = 0.0f;
+    for (uint32_t i = 0; i < n_tail && ok; ++i) {
+        const llama_pos pos = p0 + (llama_pos) (n_replay + i);
+        llama_batch batch_one = llama_batch_init(1, 0, 1);
+        common_batch_add(batch_one, tok(1, pos), pos, { 1 }, true);
+        ok = llama_decode(ctx_roll, batch_one) == 0;
+        ok = ok && llama_decode(ctx_ref, batch_one) == 0;
+        llama_batch_free(batch_one);
+        if (!ok) {
+            break;
+        }
+
+        const float * l_roll = llama_get_logits_ith(ctx_roll, 0);
+        const float * l_ref  = llama_get_logits_ith(ctx_ref,  0);
+        ok = l_roll != nullptr && l_ref != nullptr;
+        for (int t = 0; ok && t < n_vocab; ++t) {
+            diff_tail = std::max(diff_tail, std::fabs(l_roll[t] - l_ref[t]));
+        }
+    }
+
+    if (!ok || diff_tail > eps) {
+        fprintf(stderr, "%s : seq-1-only decode leaked seq 0 state (ok=%d, max diff %g)\n",
+                __func__, ok ? 1 : 0, (double) diff_tail);
+        cleanup();
+        return false;
+    }
+
+    fprintf(stderr, "%s : seq-1-only decode independent of seq 0 (max diff %g)\n", __func__, (double) diff_tail);
+    cleanup();
+    return true;
+}
+
 int main(int argc, char ** argv) {
     std::setlocale(LC_NUMERIC, "C");
 
@@ -220,5 +392,10 @@ int main(int argc, char ** argv) {
     llama_free(ctx_src);
     llama_free(ctx_dst);
     llama_free(ctx_dirty);
+
+    if (!test_multi_seq_split_replay(params, model, n_vocab)) {
+        return 1;
+    }
+
     return 0;
 }

tools/CMakeLists.txt

@@ -27,7 +27,6 @@ else()
         add_subdirectory(server)
     endif()
     add_subdirectory(tokenize)
-    add_subdirectory(parser)
     add_subdirectory(tts)
     add_subdirectory(mtmd)
     if (GGML_RPC)

tools/mtmd/clip-model.h

@@ -29,10 +29,10 @@ enum patch_merge_type {
     PATCH_MERGE_SPATIAL_UNPAD,
 };
 
+// all algos are Pillow-compatible (matching PIL.Image.resize output)
 enum resize_algo {
-    RESIZE_ALGO_BILINEAR, // stretch to target resolution
-    RESIZE_ALGO_BICUBIC, // center-crop when aspect ratio doesn't match
-    RESIZE_ALGO_BICUBIC_PILLOW,
+    RESIZE_ALGO_BILINEAR,
+    RESIZE_ALGO_BICUBIC,
     RESIZE_ALGO_LANCZOS,
 };
 
@@ -73,7 +73,7 @@ struct clip_hparams {
     int32_t preproc_max_tiles = 0;
     int32_t preproc_tile_size = 0; // local tile size (deepseek-ocr)
     resize_algo image_resize_algo_rf = RESIZE_ALGO_BICUBIC;
-    resize_algo image_resize_algo_ov = RESIZE_ALGO_BILINEAR;
+    resize_algo image_resize_algo_ov = RESIZE_ALGO_BICUBIC;
     pad_style image_pad_rf = PAD_CEIL;  // padding style for the refined image (e.g. llava-1.6)
     pad_style image_pad_ov = PAD_NONE;  // padding style for the overview image (e.g. llava-1.6)
     std::array<uint8_t, 3> image_pad_color_rf = {0, 0, 0}; // padding color for refined image

tools/mtmd/clip.cpp

@@ -1420,20 +1420,18 @@ struct clip_model_loader {
                         hparams.image_pad_color     = {122, 116, 104};
                         if (!hparams.image_res_candidates.empty()) {
                             hparams.image_resize_pad  = PAD_CEIL;
-                            hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                            hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         } else {
                             // llava-1.6 default params
                             hparams.image_pad_ov         = PAD_NONE;
                             hparams.image_pad_rf         = PAD_CEIL;
                             hparams.image_pad_color_rf   = {122, 116, 104};
-                            hparams.image_resize_algo_rf = RESIZE_ALGO_BICUBIC;
-                            hparams.image_resize_algo_ov = RESIZE_ALGO_BILINEAR;
                         }
                     } break;
                 case PROJECTOR_TYPE_GLM_EDGE:
                     {
                         hparams.image_resize_pad  = PAD_CEIL;
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                     } break;
                 case PROJECTOR_TYPE_MINICPMV:
                     {
@@ -1490,6 +1488,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_IDEFICS3:
                     {
                         // use default llava-uhd preprocessing params
+                        hparams.image_resize_algo = RESIZE_ALGO_LANCZOS;
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
                         get_u32(KEY_PREPROC_IMAGE_SIZE, hparams.image_longest_edge, false);
                         hparams.set_limit_image_tokens();
@@ -1516,7 +1515,7 @@ struct clip_model_loader {
                         // ref: https://huggingface.co/mistral-community/pixtral-12b/blob/main/preprocessor_config.json
                         // TODO: verify the image_min_tokens
                         hparams.n_merge = 1; // the original pixtral does not use patch merging
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         hparams.rope_theta = 10000.0f;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         hparams.set_limit_image_tokens(8, 1024);
@@ -1544,7 +1543,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_DOTS3NOTE_V:
                     {
                         hparams.rope_theta = 10000.0f;
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge);
                         get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);
                         get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels);
@@ -1562,7 +1561,7 @@ struct clip_model_loader {
                     } break;
                 case PROJECTOR_TYPE_KIMIVL:
                     {
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         hparams.rope_theta = 10000.0f;
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
                         // TODO: check kimivl preprocessor for exact values
@@ -1601,7 +1600,7 @@ struct clip_model_loader {
                     {
                         hparams.rope_theta = 100.0f;
                         hparams.n_merge = 3; // pooling_kernel_size
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
                         if (model.proj_type == PROJECTOR_TYPE_GEMMA4UV) {
                             // for "unified" variant, we directly use a bigger patch size, because the "token merging" is done directly on conv layer
@@ -1618,6 +1617,7 @@ struct clip_model_loader {
                         // Gemma3n uses MobileNetV5 which produces 256 tokens (16x16)
                         // Similar configuration to Gemma3
                         hparams.n_merge = 1;  // MobileNetV5 handles resizing internally
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
                     } break;
                 case PROJECTOR_TYPE_QWEN2VL:
@@ -1625,7 +1625,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_QWEN3VL:
                     {
                         hparams.n_merge = 2; // default value for Qwen 2 and 2.5
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         get_u32(KEY_WIN_ATTN_PATTERN, hparams.n_wa_pattern, model.proj_type == PROJECTOR_TYPE_QWEN25VL); // only 2.5 requires it
                         // ref: https://huggingface.co/Qwen/Qwen2.5-VL-7B-Instruct/blob/main/preprocessor_config.json
@@ -1641,7 +1641,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_MINIMAX_M3:
                     {
                         hparams.n_merge = 2; // spatial_merge_size
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         hparams.image_resize_pad  = PAD_NONE;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         // n_merge is used as a divisor in clip_image_batch_encode
@@ -1666,7 +1666,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_MIMOVL:
                     {
                         hparams.n_merge = 2; // spatial_merge_size
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         get_u32(string_format(KEY_N_HEAD_KV, "vision"), hparams.n_head_kv);
                         // 1D banded sliding-window radius (visual_token_window_size); required
@@ -1713,15 +1713,15 @@ struct clip_model_loader {
                         log_ffn_op = "gelu_erf";
                         hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
 
-                        // reka model performs better when using resize_bicubic, which stretches
-                        // the image to fit fixed square size
+                        // reka model performs better when the image is stretched to fit
+                        // fixed square size (no padding)
                         hparams.image_resize_pad = PAD_NONE;
                     } break;
                 case PROJECTOR_TYPE_GLM4V:
                     {
                         hparams.rope_theta = 10000.0f;
                         hparams.n_merge = 2; // default value for GLM4-V
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         hparams.set_limit_image_tokens(8, 4096);
                         hparams.set_warmup_n_tokens(46*46); // avoid OOM on warmup
@@ -1729,6 +1729,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_LLAMA4:
                     {
                         hparams.rope_theta = 10000.0f;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
                         set_llava_uhd_res_candidates(model, 3);
                     } break;
@@ -1840,7 +1841,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_PADDLEOCR:
                     {
                         hparams.n_merge = 2;
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);
                         get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels);
 
@@ -1852,7 +1853,7 @@ struct clip_model_loader {
                         hparams.patch_size = 16;
                         hparams.image_size = 1024;
                         hparams.warmup_image_size = 1024;
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         hparams.image_pad_color = {127, 127, 127};
 
                         get_u32(KEY_SAM_N_BLOCK, hparams.sam_n_layer, true);
@@ -1882,7 +1883,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_HUNYUANVL:
                     {
                         hparams.n_merge = 2;
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
+                        hparams.image_resize_algo = RESIZE_ALGO_LANCZOS;
                         hparams.image_resize_pad = PAD_NONE;
                         hparams.ffn_op = FFN_GELU;
                         hparams.set_limit_image_tokens(256, 16384);
@@ -1955,12 +1956,12 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_JANUS_PRO:
                     {
                         hparams.image_pad_color   = {127, 127, 127};
-                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                     } break;
                 case PROJECTOR_TYPE_GRANITE4_VISION:
                     {
                         // SigLIP tower.
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
                         hparams.image_resize_pad = PAD_CEIL;
 
                         // NOTE: feature_layers loaded in common path as optional

tools/mtmd/mtmd-image.cpp

@@ -58,22 +58,7 @@ struct img_tool {
 
         if (padding == PAD_NONE) {
             // direct resize
-            switch (algo) {
-                case RESIZE_ALGO_BILINEAR:
-                    resize_bilinear(src, dst, target_resolution.width, target_resolution.height);
-                    break;
-                case RESIZE_ALGO_BICUBIC:
-                    resize_bicubic(src, dst, target_resolution.width, target_resolution.height);
-                    break;
-                case RESIZE_ALGO_BICUBIC_PILLOW:
-                    resize_bicubic_pillow(src, dst, target_resolution.width, target_resolution.height);
-                    break;
-                case RESIZE_ALGO_LANCZOS:
-                    resize_lanczos_pillow(src, dst, target_resolution.width, target_resolution.height);
-                    break;
-                default:
-                    throw std::runtime_error("Unsupported resize algorithm");
-            }
+            resize_pillow(src, dst, target_resolution.width, target_resolution.height, algo);
         } else {
             // resize with padding
             clip_image_u8 resized_image;
@@ -90,22 +75,7 @@ struct img_tool {
                 new_height = std::min(static_cast<int>(std::ceil(src.get_size().height * scale)), target_resolution.height);
             }
 
-            switch (algo) {
-                case RESIZE_ALGO_BILINEAR:
-                    resize_bilinear(src, resized_image, new_width, new_height);
-                    break;
-                case RESIZE_ALGO_BICUBIC:
-                    resize_bicubic(src, resized_image, new_width, new_height);
-                    break;
-                case RESIZE_ALGO_BICUBIC_PILLOW:
-                    resize_bicubic_pillow(src, resized_image, new_width, new_height);
-                    break;
-                case RESIZE_ALGO_LANCZOS:
-                    resize_lanczos_pillow(src, resized_image, new_width, new_height);
-                    break;
-                default:
-                    throw std::runtime_error("Unsupported resize algorithm");
-            }
+            resize_pillow(src, resized_image, new_width, new_height, algo);
 
             // fill dst with pad_color
             fill(dst, pad_color);
@@ -224,152 +194,37 @@ struct img_tool {
     }
 
 private:
-    // Bilinear resize function
-    static void resize_bilinear(const clip_image_u8 & src, clip_image_u8 & dst, int target_width, int target_height) {
-        const auto src_size = src.get_size();
-        if (src_size.width == 0 || src_size.height == 0) { dst.set_size({0, 0}, false); return; }
-        if (target_width  <= 0) target_width  = 1;
-        if (target_height <= 0) target_height = 1;
-
-        dst.set_size({target_width, target_height}, false);
-
-        if (src.is_placeholder()) {
-            // no-op for placeholder image, just set the size and return
-            return;
-        }
-
-        float x_ratio = target_width  > 1 ? static_cast<float>(src_size.width  - 1) / (target_width  - 1) : 0.0f;
-        float y_ratio = target_height > 1 ? static_cast<float>(src_size.height - 1) / (target_height - 1) : 0.0f;
-
-        for (int y = 0; y < target_height; ++y) {
-            for (int x = 0; x < target_width; ++x) {
-                float px = x * x_ratio;
-                float py = y * y_ratio;
-
-                int x0 = std::min(static_cast<int>(px), src_size.width  - 1);
-                int y0 = std::min(static_cast<int>(py), src_size.height - 1);
-                int x1 = std::min(x0 + 1, src_size.width  - 1);
-                int y1 = std::min(y0 + 1, src_size.height - 1);
-
-                float xf = px - x0;
-                float yf = py - y0;
-
-                const auto p00 = src.get_pixel(x0, y0);
-                const auto p10 = src.get_pixel(x1, y0);
-                const auto p01 = src.get_pixel(x0, y1);
-                const auto p11 = src.get_pixel(x1, y1);
-
-                std::array<uint8_t, 3> pixel;
-                for (int c = 0; c < 3; ++c) {
-                    float top    = lerp(static_cast<float>(p00[c]), static_cast<float>(p10[c]), xf);
-                    float bottom = lerp(static_cast<float>(p01[c]), static_cast<float>(p11[c]), xf);
-                    pixel[c] = static_cast<uint8_t>(lerp(top, bottom, yf));
-                }
-                dst.set_pixel(x, y, pixel);
-            }
-        }
-    }
-
-    // Bicubic resize function
-    // part of image will be cropped if the aspect ratio is different
-    static void resize_bicubic(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) {
-        const auto img_size = img.get_size();
-        const int nx = img_size.width;
-        const int ny = img_size.height;
-
-        dst.set_size({target_width, target_height}, false);
-
-        if (img.is_placeholder()) {
-            // no-op for placeholder image, just set the size and return
-            return;
-        }
-
-        float Cc;
-        float C[5] = {};
-        float d0, d2, d3, a0, a1, a2, a3;
-        int i, j, k, jj;
-        int x, y;
-        float dx, dy;
-        float tx, ty;
-
-        tx = (float)nx / (float)target_width;
-        ty = (float)ny / (float)target_height;
-
-        // Bicubic interpolation; adapted from ViT.cpp, inspired from :
-        //    -> https://github.com/yglukhov/bicubic-interpolation-image-processing/blob/master/libimage.c#L36
-        //    -> https://en.wikipedia.org/wiki/Bicubic_interpolation
-
-        for (i = 0; i < target_height; i++) {
-            for (j = 0; j < target_width; j++) {
-                x = (int)(tx * j);
-                y = (int)(ty * i);
-
-                dx = tx * j - x;
-                dy = ty * i - y;
-
-                std::array<uint8_t, 3> pixel;
-                for (k = 0; k < 3; k++) {
-                    for (jj = 0; jj <= 3; jj++) {
-                        d0 = img.get_pixel(clip(x - 1, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k] - img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
-                        d2 = img.get_pixel(clip(x + 1, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k] - img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
-                        d3 = img.get_pixel(clip(x + 2, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k] - img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
-                        a0 = img.get_pixel(clip(x, 0, nx - 1), clip(y - 1 + jj, 0, ny - 1))[k];
-
-                        a1 = -1.0 / 3 * d0 + d2 - 1.0 / 6 * d3;
-                        a2 =  1.0 / 2 * d0 +      1.0 / 2 * d2;
-                        a3 = -1.0 / 6 * d0 -      1.0 / 2 * d2 + 1.0 / 6 * d3;
-
-                        C[jj] = a0 + a1 * dx + a2 * dx * dx + a3 * dx * dx * dx;
-
-                        d0 = C[0] - C[1];
-                        d2 = C[2] - C[1];
-                        d3 = C[3] - C[1];
-                        a0 = C[1];
-                        a1 = -1.0 / 3 * d0 + d2 - 1.0 / 6 * d3;
-                        a2 =  1.0 / 2 * d0 +      1.0 / 2 * d2;
-                        a3 = -1.0 / 6 * d0 -      1.0 / 2 * d2 + 1.0 / 6 * d3;
-                        Cc = a0 + a1 * dy + a2 * dy * dy + a3 * dy * dy * dy;
-
-                        const uint8_t Cc2 = std::min(std::max(std::round(Cc), 0.0f), 255.0f);
-                        pixel[k] = Cc2;
-                    }
-                }
-                dst.set_pixel(j, i, pixel);
-            }
-        }
-    }
-
-    // Pillow-compatible separable resampling (Bicubic and Lanczos)
+    // Pillow-compatible separable resampling (Bilinear, Bicubic and Lanczos)
     // Adapted from https://github.com/python-pillow/Pillow/blob/main/src/libImaging/Resample.c
     //
     // Key properties:
     // 1. Separable filtering: horizontal pass followed by vertical pass
     // 2. Pre-computes normalized filter coefficients for each output pixel
     // 3. Fixed-point integer arithmetic (22 fractional bits) for speed and determinism
-    static bool resize_bicubic_pillow(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) {
-        return resize_pillow(img, dst, target_width, target_height, /*use_lanczos=*/false);
-    }
-
-    // Lanczos-3 (support radius 3), matches Pillow's Image.LANCZOS
-    static bool resize_lanczos_pillow(const clip_image_u8 & img, clip_image_u8 & dst, int target_width, int target_height) {
-        return resize_pillow(img, dst, target_width, target_height, /*use_lanczos=*/true);
-    }
-
     static bool resize_pillow(
             const clip_image_u8 & img,
             clip_image_u8 & dst,
             int target_width,
             int target_height,
-            bool use_lanczos) {
+            resize_algo algo) {
         // Fixed-point precision: 22 bits = 32 (int32_t) - 8 (uint8_t pixels) - 2 (headroom for accumulation)
         // This allows encoding fractional weights as integers: weight * 2^22
         const int PRECISION_BITS = 32 - 8 - 2;
 
-        // Resample filter: Lanczos-3 (support [-3, 3]) or bicubic with a = -0.5 (support [-2, 2])
-        // Note: GGML/PyTorch bicubic uses a = -0.75, Pillow uses a = -0.5
+        // Filter support radius
+        double filter_support;
+        switch (algo) {
+            case RESIZE_ALGO_BILINEAR: filter_support = 1.0; break;
+            case RESIZE_ALGO_BICUBIC:  filter_support = 2.0; break;
+            case RESIZE_ALGO_LANCZOS:  filter_support = 3.0; break;
+            default:
+                throw std::runtime_error("Unsupported resize algorithm");
+        }
+
         // Returns filter weight for distance x from pixel center
-        auto resample_filter = [use_lanczos](double x) -> double {
-            if (use_lanczos) {
+        // Note: for bicubic, Pillow uses a = -0.5 while GGML/PyTorch use a = -0.75
+        auto resample_filter = [algo](double x) -> double {
+            if (algo == RESIZE_ALGO_LANCZOS) {
                 if (-3.0 <= x && x < 3.0) {
                     auto sinc = [](double v) {
                         if (v == 0.0) {
@@ -383,10 +238,15 @@ private:
                 return 0.0;
             }
 
-            constexpr double a = -0.5;
             if (x < 0.0) {
                 x = -x;
             }
+
+            if (algo == RESIZE_ALGO_BILINEAR) {
+                return x < 1.0 ? 1.0 - x : 0.0;
+            }
+
+            constexpr double a = -0.5;
             if (x < 1.0) {
                 return ((a + 2.0) * x - (a + 3.0)) * x * x + 1;
             }
@@ -396,9 +256,6 @@ private:
             return 0.0;  // Zero outside [-2, 2]
         };
 
-        // Filter support radius: 2 for bicubic, 3 for lanczos
-        const double filter_support = use_lanczos ? 3.0 : 2.0;
-
         // Clipping function for 8-bit values
         auto clip8 = [](int val) -> uint8_t {
             if (val < 0) return 0;
@@ -493,100 +350,92 @@ private:
             const double fxp_scale = std::ldexp(1.0, PRECISION_BITS); // 1.0 * 2^PRECISION_BITS
 
             for (int i = 0; i < outSize * ksize; i++) {
-                if (use_lanczos) {
-                    // Pillow adds +/- 0.5 then truncates toward zero; std::round would round twice
-                    const double rounded = pre_weights[i] * fxp_scale + (pre_weights[i] < 0 ? -0.5 : 0.5);
-                    weights[i] = static_cast<int32_t>(rounded);
-                    continue;
-                }
-                double tmp_val = pre_weights[i] * fxp_scale;
-                if (pre_weights[i] < 0) {
-                    tmp_val -= 0.5;
-                } else {
-                    tmp_val += 0.5;
-                }
-                tmp_val = std::round(tmp_val);
-                tmp_val = std::clamp(tmp_val,
-                                     static_cast<double>(std::numeric_limits<int32_t>::min()),
-                                     static_cast<double>(std::numeric_limits<int32_t>::max()));
-                weights[i] = static_cast<int32_t>(tmp_val);
+                // Pillow adds +/- 0.5 then truncates toward zero; std::round would round twice
+                const double rounded = pre_weights[i] * fxp_scale + (pre_weights[i] < 0 ? -0.5 : 0.5);
+                weights[i] = static_cast<int32_t>(rounded);
             }
 
             return ksize;
         };
 
         // Horizontal resampling pass
-        // Resizes width from imIn to out_nx, preserving height
-        auto resample_horizontal = [&](const clip_image_u8 & imIn, clip_image_u8 & imOut,
+        // Resizes width from src to out_nx, preserving height
+        auto resample_horizontal = [&](const uint8_t * src, int in_nx, int in_ny,
                                        int out_nx,
                                        int ksize, const std::vector<int> & bounds, const std::vector<int32_t> & weights) {
-            const int in_ny = imIn.get_size().height;
-            imOut.set_size({out_nx, in_ny}, false);
+            std::vector<uint8_t> out((size_t) out_nx * in_ny * 3);
 
             // Process each row independently
             for (int yy = 0; yy < in_ny; yy++) {
+                const uint8_t * src_row = src + (size_t) yy * in_nx * 3;
+                uint8_t * dst_row = out.data() + (size_t) yy * out_nx * 3;
+
                 // For each output pixel in this row
                 for (int xx = 0; xx < out_nx; xx++) {
-                    // Get the range of input pixels and filter coefficients
-                    int xmin = bounds[xx * 2 + 0];  // First input pixel index
-                    int xcnt = bounds[xx * 2 + 1];  // Number of input pixels
+                    const int xmin = bounds[xx * 2 + 0];  // First input pixel index
+                    const int xcnt = bounds[xx * 2 + 1];  // Number of input pixels
+                    const int32_t * k = &weights[xx * ksize];
+                    const uint8_t * p = src_row + (size_t) xmin * 3;
 
-                    // Initialize accumulators for RGB channels with rounding bias (0.5 in fixed-point)
+                    // Accumulators for RGB channels, with rounding bias (0.5 in fixed-point)
                     int32_t ss0 = 1 << (PRECISION_BITS - 1);
                     int32_t ss1 = 1 << (PRECISION_BITS - 1);
                     int32_t ss2 = 1 << (PRECISION_BITS - 1);
 
                     // Convolve: sum weighted input pixels
                     for (int x = 0; x < xcnt; x++) {
-                        const auto src_px = imIn.get_pixel(x + xmin, yy);
-                        ss0 += src_px[0] * weights[xx * ksize + x];  // R channel
-                        ss1 += src_px[1] * weights[xx * ksize + x];  // G channel
-                        ss2 += src_px[2] * weights[xx * ksize + x];  // B channel
+                        ss0 += p[0] * k[x];
+                        ss1 += p[1] * k[x];
+                        ss2 += p[2] * k[x];
+                        p += 3;
                     }
 
                     // Convert back from fixed-point (divide by 2^PRECISION_BITS) and clamp to [0,255]
-                    imOut.set_pixel(xx, yy, {clip8(ss0 >> PRECISION_BITS),
-                                             clip8(ss1 >> PRECISION_BITS),
-                                             clip8(ss2 >> PRECISION_BITS)});
+                    dst_row[xx * 3 + 0] = clip8(ss0 >> PRECISION_BITS);
+                    dst_row[xx * 3 + 1] = clip8(ss1 >> PRECISION_BITS);
+                    dst_row[xx * 3 + 2] = clip8(ss2 >> PRECISION_BITS);
                 }
             }
+
+            return out;
         };
 
         // Vertical resampling pass
-        // Resizes height from imIn to out_ny, preserving width
-        auto resample_vertical = [&](const clip_image_u8 & imIn, clip_image_u8 & imOut,
+        // Resizes height from src to out_ny, preserving width
+        // Accumulates whole rows at once (contiguous access, auto-vectorizes well)
+        auto resample_vertical = [&](const uint8_t * src, int in_nx,
                                      int out_ny,
                                      int ksize, const std::vector<int> & bounds, const std::vector<int32_t> & weight) {
-            const int in_nx = imIn.get_size().width;
-            imOut.set_size({in_nx, out_ny}, false);
+            const size_t row_elems = (size_t) in_nx * 3;
+            std::vector<uint8_t> out(row_elems * out_ny);
+            std::vector<int32_t> acc(row_elems);
 
             // For each output row
             for (int yy = 0; yy < out_ny; yy++) {
-                // Get the range of input rows and filter coefficients
-                int ymin = bounds[yy * 2 + 0];  // First input row index
-                int ycnt = bounds[yy * 2 + 1];  // Number of input rows
+                const int ymin = bounds[yy * 2 + 0];  // First input row index
+                const int ycnt = bounds[yy * 2 + 1];  // Number of input rows
+                const int32_t * k = &weight[yy * ksize];
 
-                // Process each column in this output row
-                for (int xx = 0; xx < in_nx; xx++) {
-                    // Initialize accumulators for RGB channels with rounding bias
-                    int32_t ss0 = 1 << (PRECISION_BITS - 1);
-                    int32_t ss1 = 1 << (PRECISION_BITS - 1);
-                    int32_t ss2 = 1 << (PRECISION_BITS - 1);
+                // Rounding bias (0.5 in fixed-point)
+                std::fill(acc.begin(), acc.end(), 1 << (PRECISION_BITS - 1));
 
-                    // Convolve: sum weighted input pixels vertically
-                    for (int y = 0; y < ycnt; y++) {
-                        const auto src_px = imIn.get_pixel(xx, y + ymin);
-                        ss0 += src_px[0] * weight[yy * ksize + y];  // R channel
-                        ss1 += src_px[1] * weight[yy * ksize + y];  // G channel
-                        ss2 += src_px[2] * weight[yy * ksize + y];  // B channel
+                // Convolve: accumulate each weighted input row
+                for (int y = 0; y < ycnt; y++) {
+                    const uint8_t * src_row = src + (size_t) (ymin + y) * row_elems;
+                    const int32_t w = k[y];
+                    for (size_t i = 0; i < row_elems; i++) {
+                        acc[i] += src_row[i] * w;
                     }
+                }
 
-                    // Convert back from fixed-point and clamp to [0,255]
-                    imOut.set_pixel(xx, yy, {clip8(ss0 >> PRECISION_BITS),
-                                             clip8(ss1 >> PRECISION_BITS),
-                                             clip8(ss2 >> PRECISION_BITS)});
+                // Convert back from fixed-point and clamp to [0,255]
+                uint8_t * dst_row = out.data() + (size_t) yy * row_elems;
+                for (size_t i = 0; i < row_elems; i++) {
+                    dst_row[i] = clip8(acc[i] >> PRECISION_BITS);
                 }
             }
+
+            return out;
         };
 
         // Main resampling logic using separable two-pass approach
@@ -610,36 +459,25 @@ private:
         }
 
         // Perform two-pass resampling
+        const uint8_t * src = img.get_ro_buf().data();
         if (need_horizontal && need_vertical) {
-            // Both horizontal and vertical
-            clip_image_u8 temp;
-            resample_horizontal(img, temp, target_width, ksize_horiz, bounds_horiz, weights_horiz);
-            resample_vertical(temp, dst, target_height, ksize_vert, bounds_vert, weights_vert);
+            auto temp = resample_horizontal(src, src_width, src_height, target_width, ksize_horiz, bounds_horiz, weights_horiz);
+            dst.set_size({target_width, target_height}, false);
+            dst.cpy_buf(resample_vertical(temp.data(), target_width, target_height, ksize_vert, bounds_vert, weights_vert));
         } else if (need_horizontal) {
-            // Only horizontal
-            resample_horizontal(img, dst, target_width, ksize_horiz, bounds_horiz, weights_horiz);
+            dst.set_size({target_width, src_height}, false);
+            dst.cpy_buf(resample_horizontal(src, src_width, src_height, target_width, ksize_horiz, bounds_horiz, weights_horiz));
         } else if (need_vertical) {
-            // Only vertical
-            resample_vertical(img, dst, target_height, ksize_vert, bounds_vert, weights_vert);
+            dst.set_size({src_width, target_height}, false);
+            dst.cpy_buf(resample_vertical(src, src_width, target_height, ksize_vert, bounds_vert, weights_vert));
         } else {
             // No resizing needed - direct copy
-            dst.set_size(img.get_size(), img.is_placeholder());
-            if (!img.is_placeholder()) {
-                dst.cpy_buf(img.get_ro_buf());
-            }
+            dst.set_size(img.get_size(), false);
+            dst.cpy_buf(img.get_ro_buf());
         }
 
         return true;
     }
-
-    static inline int clip(int x, int lower, int upper) {
-        return std::max(lower, std::min(x, upper));
-    }
-
-    // Linear interpolation between two points
-    static inline float lerp(float s, float e, float t) {
-        return s + (e - s) * t;
-    }
 };
 
 
@@ -1264,7 +1102,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_deepseekocr::preprocess(const cli
     clip_image_u8 padded;
     img_tool::resize(img, padded,
                     { base_size, base_size },
-                    RESIZE_ALGO_BICUBIC_PILLOW,
+                    RESIZE_ALGO_BICUBIC,
                     PAD_NEAREST,
                     hparams.image_pad_color);
     output.append_overview(hparams, padded, true);
@@ -1280,7 +1118,7 @@ mtmd_image_preproc_out mtmd_image_preprocessor_deepseekocr::preprocess(const cli
         grid_h = grid.height;
 
         clip_image_u8 refined;
-        img_tool::resize(img, refined, { tile_size * grid_w, tile_size * grid_h }, RESIZE_ALGO_BICUBIC_PILLOW,
+        img_tool::resize(img, refined, { tile_size * grid_w, tile_size * grid_h }, RESIZE_ALGO_BICUBIC,
                          PAD_NONE);
 
         for (int row = 0; row < grid_h; row++) {

tools/parser/CMakeLists.txt

@@ -1,20 +0,0 @@
-if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
-    # this tool is disabled on Windows when building with shared libraries because it uses internal functions not exported with LLAMA_API
-    set(TARGET llama-debug-template-parser)
-    add_executable(${TARGET} debug-template-parser.cpp)
-    target_link_libraries(${TARGET} PRIVATE llama-common llama ${CMAKE_THREAD_LIBS_INIT})
-    target_compile_features(${TARGET} PRIVATE cxx_std_17)
-
-    if(LLAMA_TOOLS_INSTALL)
-        install(TARGETS ${TARGET} RUNTIME)
-    endif()
-endif()
-
-set(TARGET llama-template-analysis)
-add_executable(${TARGET} template-analysis.cpp)
-target_link_libraries(${TARGET} PRIVATE llama-common llama ${CMAKE_THREAD_LIBS_INIT})
-target_compile_features(${TARGET} PRIVATE cxx_std_17)
-
-if(LLAMA_TOOLS_INSTALL)
-    install(TARGETS ${TARGET} RUNTIME)
-endif()

tools/parser/debug-template-parser.cpp

@@ -1,469 +0,0 @@
-#include "../src/llama-grammar.h"
-#include "chat-auto-parser.h"
-#include "chat.h"
-#include "common.h"
-#include "gguf.h"
-#include "jinja/runtime.h"
-#include "log.h"
-#include "json.h"
-#include "peg-parser.h"
-
-#include <fstream>
-#include <iterator>
-#include <numeric>
-#include <optional>
-#include <sstream>
-#include <string>
-
-using json = common_json;
-
-enum class output_mode {
-    ANALYSIS,  // Only output analysis results (default)
-    TEMPLATE,  // Only output rendered template
-    BOTH       // Output both
-};
-
-enum class input_message_type {
-    NONE,                    // Don't render any message scenarios (only analysis)
-    CONTENT_ONLY,            // Simple assistant message with content
-    REASONING_CONTENT,       // Message with reasoning_content + content
-    TOOL_CALL_ONLY,          // Message with tool_calls only
-    CONTENT_TOOL_CALL,       // Message with content + tool_calls
-    REASONING_TOOL_CALL,     // Message with reasoning_content + tool_calls
-    CONTENT_FAKE_TOOL_CALL,  // Message with content but no actual tool_calls (for testing)
-    ALL                      // Render all scenarios
-};
-
-struct debug_options {
-    std::string        template_path;
-    bool               with_tools        = true;
-    bool               generation_prompt = true;
-    bool               enable_reasoning  = true;
-    bool               debug_jinja       = false;
-    bool               force_tool_call   = false;
-    bool               parallel_tool_calls = true;
-    output_mode        mode              = output_mode::BOTH;
-    input_message_type input_message     = input_message_type::NONE;
-};
-
-static std::string read_file(const std::string & path) {
-    std::ifstream fin(path, std::ios::binary);
-    if (!fin.is_open()) {
-        throw std::runtime_error("Could not open file: " + path);
-    }
-    std::ostringstream buf;
-    buf << fin.rdbuf();
-    return buf.str();
-}
-
-static std::string read_gguf_chat_template(const std::string & path) {
-    struct gguf_init_params params = { /*no_alloc =*/true,  // We only need metadata, not tensor data
-                                       /*ctx=*/nullptr };
-
-    struct gguf_context * ctx = gguf_init_from_file(path.c_str(), params);
-    if (ctx == nullptr) {
-        throw std::runtime_error("Could not open GGUF file: " + path);
-    }
-
-    const char * key    = "tokenizer.chat_template";
-    int64_t      key_id = gguf_find_key(ctx, key);
-
-    if (key_id == -1) {
-        gguf_free(ctx);
-        throw std::runtime_error("GGUF file does not contain chat template key: " + std::string(key));
-    }
-
-    const char * template_str = gguf_get_val_str(ctx, key_id);
-    if (template_str == nullptr) {
-        gguf_free(ctx);
-        throw std::runtime_error("GGUF file contains chat template key but value is null");
-    }
-
-    std::string result = template_str;
-    gguf_free(ctx);
-    return result;
-}
-
-static void print_usage(const char * program_name) {
-    LOG_ERR("Usage: %s <template_or_gguf_path> [options]\n", program_name);
-    LOG_ERR("\nOptions:\n");
-    LOG_ERR("  --no-tools              Disable tool definitions\n");
-    LOG_ERR("  --force-tool-call       Set tool calls to forced\n");
-    LOG_ERR("  --parallel-tool-calls=0|1 Set parallel_tool_calls (default: 1)\n");
-    LOG_ERR("  --generation-prompt=0|1 Set add_generation_prompt (default: 1)\n");
-    LOG_ERR("  --enable-reasoning=0|1  Enable reasoning parsing (default: 1)\n");
-    LOG_ERR("  --output=MODE           Output mode: analysis, template, both (default: both)\n");
-    LOG_ERR("  --debug-jinja           Enable Jinja fine-grained debug\n");
-    LOG_ERR("  --input-message=TYPE    Message type to render:\n");
-    LOG_ERR("                          content_only, reasoning_content, tool_call_only,\n");
-    LOG_ERR("                          content_tool_call, reasoning_tool_call,\n");
-    LOG_ERR("                          content_fake_tool_call, all\n");
-    LOG_ERR("\nExamples:\n");
-    LOG_ERR("  %s template.jinja --input-message=all --generation-prompt=1\n", program_name);
-    LOG_ERR("  %s template.jinja --output=template --input-message=tool_call_only\n", program_name);
-}
-
-static bool parse_bool_option(const std::string & value) {
-    return value == "1" || value == "true" || value == "yes";
-}
-
-static bool parse_options(int argc, char ** argv, debug_options & opts) {
-    if (argc < 2) {
-        print_usage(argv[0]);
-        return false;
-    }
-
-    opts.template_path = argv[1];
-
-    for (int i = 2; i < argc; ++i) {
-        std::string arg = argv[i];
-
-        if (arg == "--force-tool-call") {
-            opts.force_tool_call = true;
-        } else if (arg == "--debug-jinja") {
-            opts.debug_jinja = true;
-        } else if (arg == "--no-tools") {
-            opts.with_tools = false;
-        } else if (arg.rfind("--parallel-tool-calls=", 0) == 0) {
-            opts.parallel_tool_calls = parse_bool_option(arg.substr(22));
-        } else if (arg.rfind("--generation-prompt=", 0) == 0) {
-            opts.generation_prompt = parse_bool_option(arg.substr(20));
-        } else if (arg.rfind("--enable-reasoning=", 0) == 0) {
-            opts.enable_reasoning = parse_bool_option(arg.substr(19));
-        } else if (arg.rfind("--output=", 0) == 0) {
-            std::string mode = arg.substr(9);
-            if (mode == "analysis") {
-                opts.mode = output_mode::ANALYSIS;
-            } else if (mode == "template") {
-                opts.mode = output_mode::TEMPLATE;
-            } else if (mode == "both") {
-                opts.mode = output_mode::BOTH;
-            } else {
-                LOG_ERR("Unknown output mode: %s\n", mode.c_str());
-                return false;
-            }
-        } else if (arg.rfind("--input-message=", 0) == 0) {
-            std::string type = arg.substr(16);
-            if (type == "content_only") {
-                opts.input_message = input_message_type::CONTENT_ONLY;
-            } else if (type == "reasoning_content") {
-                opts.input_message = input_message_type::REASONING_CONTENT;
-            } else if (type == "tool_call_only") {
-                opts.input_message = input_message_type::TOOL_CALL_ONLY;
-            } else if (type == "content_tool_call") {
-                opts.input_message = input_message_type::CONTENT_TOOL_CALL;
-            } else if (type == "reasoning_tool_call") {
-                opts.input_message = input_message_type::REASONING_TOOL_CALL;
-            } else if (type == "content_fake_tool_call") {
-                opts.input_message = input_message_type::CONTENT_FAKE_TOOL_CALL;
-            } else if (type == "all") {
-                opts.input_message = input_message_type::ALL;
-            } else {
-                LOG_ERR("Unknown input message type: %s\n", type.c_str());
-                return false;
-            }
-        } else {
-            LOG_ERR("Unknown option: %s\n", arg.c_str());
-            print_usage(argv[0]);
-            return false;
-        }
-    }
-
-    return true;
-}
-
-static json build_user_message() {
-    return json{
-        { "role",    "user"                               },
-        { "content", "Hello, please help me with a task." }
-    };
-}
-
-static json build_content_only_message() {
-    return json{
-        { "role",    "assistant"                                   },
-        { "content", "Hello! I'm here to help you with your task." }
-    };
-}
-
-static json build_reasoning_content_message() {
-    return json{
-        { "role",              "assistant"                                                               },
-        { "content",           "Hello! I'm here to help you with your task."                             },
-        { "reasoning_content", "The user is greeting me and asking for help. I should respond politely." }
-    };
-}
-
-static json build_tool_call_only_message() {
-    return json{
-        { "role",       "assistant"      },
-        { "content",    nullptr          },
-        { "tool_calls",
-         json::array({ json{
-              { "type", "function" },
-              { "function", json{ { "name", "test_function_name" },
-                                  { "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } },
-              { "id", "123456789" } } }) }
-    };
-}
-
-static json build_content_tool_call_message() {
-    return json{
-        { "role",       "assistant"                                                                              },
-        { "content",    "I'll help you by calling a function."                                                   },
-        { "tool_calls",
-         json::array({ json{
-              { "type", "function" },
-              { "function",
-                json{ { "name", "test_function_name" },
-                      { "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } } } }) }
-    };
-}
-
-static json build_reasoning_tool_call_message() {
-    return json{
-        { "role",              "assistant"                                                                       },
-        { "content",           nullptr                                                                           },
-        { "reasoning_content", "I need to call a function to help with this task."                               },
-        { "tool_calls",
-         json::array({ json{
-              { "type", "function" },
-              { "function",
-                json{ { "name", "test_function_name" },
-                      { "arguments", json::object({ { "param1", "value1" }, { "param2", "value2" } }) } } } } }) }
-    };
-}
-
-static json build_content_fake_tool_call_message() {
-    // This message has content but NO tool_calls field
-    // It's used to test if a template renders tool definitions but not tool calls
-    return json{
-        { "role",    "assistant"                            },
-        { "content", "I'll help you by calling a function." }
-    };
-}
-
-static json build_tools_definition() {
-    json parameters_schema                    = json::object();
-    parameters_schema["type"]                 = "object";
-    parameters_schema["properties"]           = json::object();
-    parameters_schema["properties"]["param1"] = json::object({
-        { "type",        "string"          },
-        { "description", "First parameter" }
-    });
-    parameters_schema["properties"]["param2"] = json::object({
-        { "type",        "string"           },
-        { "description", "Second parameter" }
-    });
-    parameters_schema["required"]             = json::array({ "param1" });
-
-    return json::array({
-        json{ { "type", "function" },
-             { "function", json{ { "name", "test_function_name" },
-                                  { "description", "A test function for debugging" },
-                                  { "parameters", parameters_schema } } } }
-    });
-}
-
-static void render_scenario(const common_chat_template & tmpl,
-                            const std::string &          scenario_name,
-                            const json &                 messages,
-                            const json &                 tools,
-                            bool                         add_generation_prompt,
-                            bool                         enable_thinking) {
-    LOG_ERR("\n=== Scenario: %s ===\n", scenario_name.c_str());
-    LOG_ERR("add_generation_prompt: %s, enable_thinking: %s\n", add_generation_prompt ? "true" : "false",
-            enable_thinking ? "true" : "false");
-
-    // When add_generation_prompt is true, add a trailing user message to trigger the prompt
-    json final_messages = messages;
-    if (add_generation_prompt && !messages.empty() && messages.back().value("role", "") == "assistant") {
-        final_messages.push_back(json{
-            { "role",    "user"                                       },
-            { "content", "Now please continue with another response." }
-        });
-    }
-
-    LOG_ERR("Messages:\n%s\n", final_messages.dump(2).c_str());
-
-    try {
-        autoparser::generation_params inputs;
-        inputs.messages                         = final_messages;
-        inputs.add_generation_prompt            = add_generation_prompt;
-        inputs.extra_context["enable_thinking"] = enable_thinking;
-
-        if (!tools.is_null() && tools.is_array() && !tools.empty()) {
-            inputs.tools = tools;
-        }
-
-        std::string output = common_chat_template_direct_apply(tmpl, inputs);
-
-        LOG_ERR("\n--- Rendered Output ---\n");
-        LOG_ERR("%s\n", output.c_str());
-        LOG_ERR("--- End Output (length: %zu) ---\n", output.length());
-    } catch (const std::exception & e) {
-        LOG_ERR("Rendering failed: %s\n", e.what());
-    }
-}
-
-static void render_all_scenarios(const common_chat_template & tmpl,
-                                 const json &                 tools,
-                                 bool                         add_generation_prompt,
-                                 bool                         enable_thinking,
-                                 input_message_type           message_type) {
-    json user_msg = build_user_message();
-
-    auto render_if = [&](input_message_type type, const std::string & name, const json & assistant_msg) {
-        if (message_type == input_message_type::ALL || message_type == type) {
-            json messages = json::array({ user_msg, assistant_msg });
-            render_scenario(tmpl, name, messages, tools, add_generation_prompt, enable_thinking);
-        }
-    };
-
-    render_if(input_message_type::CONTENT_ONLY, "content_only", build_content_only_message());
-    render_if(input_message_type::REASONING_CONTENT, "reasoning_content", build_reasoning_content_message());
-    render_if(input_message_type::TOOL_CALL_ONLY, "tool_call_only", build_tool_call_only_message());
-    render_if(input_message_type::CONTENT_TOOL_CALL, "content_tool_call", build_content_tool_call_message());
-    render_if(input_message_type::REASONING_TOOL_CALL, "reasoning_tool_call", build_reasoning_tool_call_message());
-    render_if(input_message_type::CONTENT_FAKE_TOOL_CALL, "content_fake_tool_call",
-              build_content_fake_tool_call_message());
-
-    // Also render with add_generation_prompt=true to show the prompt ending
-    if (message_type == input_message_type::ALL) {
-        LOG_ERR("\n\n=== Generation Prompt Scenarios (add_generation_prompt=true) ===\n");
-
-        json prompt_messages = json::array({ user_msg });
-        render_scenario(tmpl, "generation_prompt_only", prompt_messages, tools, true, enable_thinking);
-
-        // With enable_thinking toggled
-        render_scenario(tmpl, "generation_prompt_thinking_disabled", prompt_messages, tools, true, false);
-    }
-}
-
-static autoparser::generation_params prepare_params(const debug_options & opts, const json & tools) {
-    autoparser::generation_params params;
-    params.messages         = json::array({ build_user_message() });
-    params.reasoning_format = opts.enable_reasoning ? COMMON_REASONING_FORMAT_DEEPSEEK : COMMON_REASONING_FORMAT_NONE;
-    params.enable_thinking  = opts.enable_reasoning;
-    params.add_generation_prompt = opts.generation_prompt;
-
-    if (opts.with_tools) {
-        params.tools       = tools;
-        params.tool_choice = opts.force_tool_call ? COMMON_CHAT_TOOL_CHOICE_REQUIRED : COMMON_CHAT_TOOL_CHOICE_AUTO;
-    } else {
-        params.tools       = json();
-        params.tool_choice = COMMON_CHAT_TOOL_CHOICE_NONE;
-    }
-    params.parallel_tool_calls = opts.parallel_tool_calls;
-    return params;
-}
-
-int main(int argc, char ** argv) {
-    // Set log level to most verbose to capture all debug output
-    common_log_set_verbosity_thold(99);
-
-    debug_options opts;
-    if (!parse_options(argc, argv, opts)) {
-        return 1;
-    }
-
-    if (opts.debug_jinja || std::getenv("LLAMA_DEBUG_JINJA") != nullptr) {
-        jinja::enable_debug(true);
-    }
-
-    std::string template_source;
-    try {
-        // Check if the file is a GGUF file
-        if (opts.template_path.size() >= 5 &&
-            opts.template_path.compare(opts.template_path.size() - 5, 5, ".gguf") == 0) {
-            template_source = read_gguf_chat_template(opts.template_path);
-        } else {
-            template_source = read_file(opts.template_path);
-        }
-    } catch (const std::exception & e) {
-        LOG_ERR("Error reading template: %s\n", e.what());
-        return 1;
-    }
-
-    LOG_ERR("Analyzing template: %s\n", opts.template_path.c_str());
-    LOG_ERR("Options: with_tools=%s, generation_prompt=%s, enable_reasoning=%s\n", opts.with_tools ? "true" : "false",
-            opts.generation_prompt ? "true" : "false", opts.enable_reasoning ? "true" : "false");
-
-    try {
-        common_chat_template chat_template(template_source, "", "");
-
-        json tools = opts.with_tools ? build_tools_definition() : json();
-
-        autoparser::generation_params params = prepare_params(opts, tools);
-        common_chat_params            parser_data;
-        if (std::optional<common_chat_params> spec_tmpl =
-                common_chat_try_specialized_template(chat_template, template_source, params)) {
-            LOG_ERR("\n");
-            LOG_ERR("This template uses a specialized parser, analysis results will not be available.\n");
-            parser_data = *spec_tmpl;
-        } else {
-            // Render template scenarios if requested
-            if (opts.input_message != input_message_type::NONE &&
-                (opts.mode == output_mode::TEMPLATE || opts.mode == output_mode::BOTH)) {
-                LOG_ERR("\n");
-                LOG_ERR("================================================================================\n");
-                LOG_ERR("                         TEMPLATE RENDERING OUTPUT\n");
-                LOG_ERR("================================================================================\n");
-
-                render_all_scenarios(chat_template, tools, opts.generation_prompt, opts.enable_reasoning,
-                                     opts.input_message);
-            }
-
-            // Output analysis if requested
-            if (opts.mode == output_mode::ANALYSIS || opts.mode == output_mode::BOTH) {
-                LOG_ERR("\n");
-                LOG_ERR("================================================================================\n");
-                LOG_ERR("                           TEMPLATE ANALYSIS\n");
-                LOG_ERR("================================================================================\n");
-
-                autoparser::autoparser analysis;
-                analysis.analyze_template(chat_template);
-
-                // Generate Parser
-                parser_data = autoparser::peg_generator::generate_parser(chat_template, params, analysis);
-            }
-        }
-
-        if (!std::empty(parser_data.parser)) {
-            LOG_ERR("\n=== Generated Parser ===\n");
-            common_peg_arena arena;
-            arena.load(parser_data.parser);
-            LOG_ERR("%s\n", arena.dump(arena.root()).c_str());
-
-            LOG_ERR("\n=== Generated Grammar ===\n");
-            LOG_ERR("%s\n", parser_data.grammar.c_str());
-
-            LOG_ERR("\n=== Generated Lazy Grammar ===\n");
-            LOG_ERR("%d\n", parser_data.grammar_lazy);
-
-            LOG_ERR("\n=== Generated Grammar Triggers ===\n");
-            for (const common_grammar_trigger & cgt : parser_data.grammar_triggers) {
-                LOG_ERR("Token: %d | Type: %d | Value: %s\n", cgt.token, cgt.type, cgt.value.c_str());
-            }
-
-            LOG_ERR("\n=== Preserved Tokens ===\n");
-            for (const std::string & token : parser_data.preserved_tokens) {
-                LOG_ERR("  '%s'\n", token.c_str());
-            }
-
-            if (!parser_data.grammar.empty()) {
-                LOG_ERR("\n=== Verifying created grammar ===\n");
-                auto * grammar = llama_grammar_init_impl(nullptr, parser_data.grammar.c_str(), "root",
-                                                         parser_data.grammar_lazy, nullptr, 0, nullptr, 0);
-                if (grammar != nullptr) {
-                    LOG_ERR("\n=== Grammar successfully created ===\n");
-                }
-            }
-        }
-    } catch (const std::exception & e) {
-        LOG_ERR("Analysis failed: %s\n", e.what());
-        return 1;
-    }
-
-    return 0;
-}

tools/server/server-context.cpp

@@ -858,8 +858,10 @@ private:
     // slots / clients
     std::vector<server_slot> slots;
 
-    int trace = 0;
-    int slots_debug = 0;
+    int trace = 0;        // env: LLAMA_TRACE
+    int slots_debug = 0;  // env: LLAMA_SERVER_SLOTS_DEBUG
+    int slots_n_diff = 0; // env: LLAMA_SERVER_SLOTS_N_DIFF
+
     int n_empty_consecutive = 0;
 
     std::unique_ptr<server_prompt_cache> prompt_cache;
@@ -1247,6 +1249,15 @@ private:
             }
         }
 
+        {
+            const char * LLAMA_SERVER_SLOTS_N_DIFF = getenv("LLAMA_SERVER_SLOTS_N_DIFF");
+            slots_n_diff = LLAMA_SERVER_SLOTS_N_DIFF ? atoi(LLAMA_SERVER_SLOTS_N_DIFF) : 0;
+
+            if (slots_n_diff) {
+                SRV_WRN("LLAMA_SERVER_SLOTS_N_DIFF = %d\n", slots_n_diff);
+            }
+        }
+
         // the update_slots() logic will always submit a maximum of n_batch or n_parallel tokens
         // note that n_batch can be > n_ctx (e.g. for non-causal attention models such as BERT where the KV cache is not used)
         {
@@ -3179,8 +3190,8 @@ private:
                                 // when the prompt prefix does not match, print the tokens around the mismatch
                                 // this is useful for debugging prompt caching
                                 if (slots_debug) {
-                                    const int np0 = std::max<int>(n_past - 4, 0);
-                                    const int np1 = std::min<int>(n_past + 6, std::min(slot.prompt.tokens.size(), slot.task->tokens.size()));
+                                    const int np0 = std::max<int>(n_past - slots_n_diff, 0);
+                                    const int np1 = std::min<int>(n_past + slots_n_diff + 2, std::min(slot.prompt.tokens.size(), slot.task->tokens.size()));
 
                                     std::stringstream ss0;
                                     std::stringstream ss1;

tools/server/tests/unit/test_slot_save.py

@@ -319,7 +319,6 @@ def test_slot_save_restore_with_two_images(mmproj_server):
         "prompt": prompt,
     })
     assert res.status_code == 200
-    content = res.body["content"]
     prompt_n_full = res.body["timings"]["prompt_n"]
     assert prompt_n_full > 64
 
@@ -345,6 +344,26 @@ def test_slot_save_restore_with_two_images(mmproj_server):
     assert res.status_code == 200
     assert res.body["timings"]["cache_n"] == prompt_n_full - 1
     assert res.body["timings"]["prompt_n"] == 1
+    content = res.body["content"]
+
+    res = server.make_request("POST", "/slots/1?action=restore", data={
+        "filename": "mm_slot_two_images.bin",
+    })
+    assert res.status_code == 200
+    assert res.body["n_restored"] == n_saved
+
+    res = server.make_request("POST", "/completions", data={
+        "temperature": 0.0,
+        "top_k": 1,
+        "id_slot": 0,
+        "cache_prompt": True,
+        "prompt": prompt,
+    })
+    assert res.status_code == 200
+    assert res.body["timings"]["cache_n"] == prompt_n_full - 1
+    assert res.body["timings"]["prompt_n"] == 1
+    content = res.body["content"]
+
     assert res.body["content"] == content
 
 

tools/server/tests/unit/test_vision_api.py

@@ -121,7 +121,7 @@ def test_vision_chat_completion_token_count():
     "prompt, image_data, success, re_content",
     [
         # test model is trained on CIFAR-10, but it's quite dumb due to small size
-        ("What is this: <__media__>\n", "IMG_BASE64_0",         True, "(cat)+"),
+        ("What is this: <__media__>\n", "IMG_BASE64_0",         True, "(cat)+|(automobile)+"),
         ("What is this: <__media__>\n", "IMG_BASE64_1",         True, "(frog)+"),
         ("What is this: <__media__>\n", "malformed",            False, None), # non-image data
         ("What is this:\n",             "",                     False, None), # empty string

tools/server/tests/utils.py

@@ -623,7 +623,7 @@ class ServerPreset:
         server.model_hf_repo = "ggml-org/tinygemma3-GGUF:Q8_0"
         server.model_alias = "tinygemma3"
         server.n_ctx = 1024
-        server.n_batch = 32
+        server.n_batch = 512
         server.n_slots = 2
         server.n_predict = 4
         server.seed = 42

tools/ui/src/lib/components/app/chat/ChatAttachments/ChatAttachmentsList/ChatAttachmentsList.svelte

@@ -3,8 +3,9 @@
 		ChatAttachmentsListItem,
 		DialogChatAttachmentsPreview,
 		DialogMcpResourcePreview,
-		HorizontalScrollCarousel
+		ScrollCarousel
 	} from '$lib/components/app';
+	import { ScrollCarouselVariant } from '$lib/enums';
 	import type { DatabaseMessageExtraMcpResource } from '$lib/types';
 	import { getAttachmentDisplayItems, isMcpPrompt, isMcpResource } from '$lib/utils';
 
@@ -42,7 +43,7 @@
 		uploadedFiles = $bindable([])
 	}: Props = $props();
 
-	let carouselRef: HorizontalScrollCarousel | undefined = $state();
+	let carouselRef: ScrollCarousel | undefined = $state();
 	let mcpResourcePreviewOpen = $state(false);
 	let mcpResourcePreviewExtra = $state<DatabaseMessageExtraMcpResource | null>(null);
 	let previewFocusIndex = $state(0);
@@ -91,11 +92,11 @@
 {#if displayItems.length > 0}
 	<div class={className} {style}>
 		{#if limitToSingleRow}
-			<HorizontalScrollCarousel bind:this={carouselRef}>
+			<ScrollCarousel bind:this={carouselRef} variant={ScrollCarouselVariant.CENTER}>
 				{#each displayItems as item (item.id)}
 					{@render attachmentitem(item)}
 				{/each}
-			</HorizontalScrollCarousel>
+			</ScrollCarousel>
 		{:else}
 			<div class="flex flex-wrap items-start justify-end gap-3">
 				{#each displayItems as item (item.id)}

tools/ui/src/lib/components/app/chat/ChatAttachments/ChatAttachmentsPreview/ChatAttachmentsPreviewThumbnailStrip.svelte

@@ -1,7 +1,8 @@
 <script lang="ts">
 	import { FileText, Music, Video } from '@lucide/svelte';
-	import { HorizontalScrollCarousel } from '$lib/components/app/misc';
+	import { ScrollCarousel } from '$lib/components/app';
 	import { ICON_CLASS_DEFAULT, UI_DATA_ATTRS } from '$lib/constants';
+	import { ScrollCarouselVariant } from '$lib/enums';
 
 	interface PreviewItem {
 		id: string;
@@ -33,7 +34,7 @@
 
 {#if items.length > 1}
 	<div class="sticky bottom-0 z-10 mt-4 flex-shrink-0">
-		<HorizontalScrollCarousel class="max-w-full">
+		<ScrollCarousel class="max-w-full" variant={ScrollCarouselVariant.CENTER}>
 			{#each items as item, index (item.id)}
 				<button
 					{...{ [UI_DATA_ATTRS.THUMBNAIL_INDEX]: index }}
@@ -64,6 +65,6 @@
 					{/if}
 				</button>
 			{/each}
-		</HorizontalScrollCarousel>
+		</ScrollCarousel>
 	</div>
 {/if}

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormInput/ChatFormInputRich.svelte

@@ -625,7 +625,7 @@
 		}
 
 		if (rootElement && (event.key === 'ArrowLeft' || event.key === 'ArrowRight')) {
-			const isWordJump = (event.altKey || event.ctrlKey) && !event.metaKey;
+			const isWordJump = (event.altKey || event.ctrlKey) && !event.metaKey && !event.shiftKey;
 			const isPlainLeft =
 				event.key === 'ArrowLeft' && !event.altKey && !event.ctrlKey && !event.metaKey;
 

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormMcpResourcesList.svelte

@@ -1,8 +1,6 @@
 <script lang="ts">
-	import {
-		ChatAttachmentsListItemMcpResource,
-		HorizontalScrollCarousel
-	} from '$lib/components/app';
+	import { ChatAttachmentsListItemMcpResource, ScrollCarousel } from '$lib/components/app';
+	import { ScrollCarouselVariant } from '$lib/enums';
 	import { mcpStore } from '$lib/stores';
 
 	interface Props {
@@ -26,7 +24,7 @@
 
 {#if hasAttachments}
 	<div class={className}>
-		<HorizontalScrollCarousel gapSize="2">
+		<ScrollCarousel gapSize="2" variant={ScrollCarouselVariant.CENTER}>
 			{#each attachments as attachment, i (attachment.id)}
 				<ChatAttachmentsListItemMcpResource
 					class={i === 0 ? 'ml-3' : ''}
@@ -35,6 +33,6 @@
 					onclick={() => handleResourceClick(attachment.resource.uri)}
 				/>
 			{/each}
-		</HorizontalScrollCarousel>
+		</ScrollCarousel>
 	</div>
 {/if}

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte

@@ -196,7 +196,7 @@
 		--assistant-min-height-offset: calc(
 			var(--last-user-message-height, 19rem) + var(--chat-form-height, 6rem) +
 				var(--chat-form-bottom-position, 0.5rem) + var(--chat-form-padding-top, 6rem) +
-				var(--assistant-margin-top, 3rem)
+				var(--assistant-margin-top, 3rem) + var(--chat-tabs-offset, 0px)
 		);
 		min-height: calc(100dvh - var(--assistant-min-height-offset));
 
@@ -204,7 +204,7 @@
 			--assistant-min-height-offset: calc(
 				var(--last-user-message-height, 18rem) + var(--chat-form-height, 6rem) +
 					var(--chat-form-bottom-position, 1rem) + var(--chat-form-padding-top, 6rem) +
-					var(--assistant-margin-top, 3rem)
+					var(--assistant-margin-top, 3rem) + var(--chat-tabs-offset, 0px)
 			);
 		}
 	}

tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreen.svelte

@@ -41,10 +41,7 @@
 	let showDeleteDialog = $state(false);
 	let showEmptyFileDialog = $state(false);
 	let isEmpty = $derived(
-		showCenteredEmpty &&
-			!conversationsStore.activeConversation &&
-			conversationsStore.activeMessages.length === 0 &&
-			!chatStore.isLoading
+		showCenteredEmpty && conversationsStore.activeMessages.length === 0 && !chatStore.isLoading
 	);
 	let activeErrorDialog = $derived(chatStore.errorDialogState);
 	let isServerLoading = $derived(serverStore.loading);
@@ -297,7 +294,7 @@
 	<ServerLoadingSplash />
 {:else}
 	<div
-		class="chat-screen flex grow flex-col min-h-[calc(100dvh-1rem)] md:min-h-full px-4 md:py-0 pt-12 pb-48 md:pb-4"
+		class="chat-screen flex grow flex-col min-h-[calc(100dvh-1rem)] md:min-h-[calc(100dvh-1rem-var(--chat-tabs-offset,0px))] px-4 md:py-0 pt-12 pb-48 md:pb-4"
 		style:--chat-form-bottom-position={chatFormBottomPosition}
 		ondragenter={dragAndDrop.dragHandlers.dragenter}
 		ondragleave={dragAndDrop.dragHandlers.dragleave}

tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte

@@ -14,6 +14,6 @@
 		tooltip="Scroll to bottom"
 		size="lg"
 		iconSize={ICON_CLASS_DEFAULT}
-		class="h-9 w-9 rounded-full bg-accent text-accent-foreground absolute bottom-4 shadow-md"
+		class="h-9 w-9 rounded-full bg-muted/60 border border-border/20 shadow-sm text-accent-foreground absolute bottom-4"
 	/>
 </div>

tools/ui/src/lib/components/app/chat/ChatTabs/ChatTabs.svelte

@@ -0,0 +1,136 @@
+<script lang="ts">
+	import ChatTabsItem from './ChatTabsItem.svelte';
+	import ChatTabsNewChatButton from './ChatTabsNewChatButton.svelte';
+	import { page } from '$app/state';
+	import { ScrollCarousel } from '$lib/components/app';
+	import {
+		CHAT_TABS_MAX_WIDTH,
+		NEW_CHAT_LABEL,
+		NEW_CHAT_TAB_ID,
+		UI_DATA_ATTRS,
+		UNNAMED_CHAT_LABEL
+	} from '$lib/constants';
+	import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
+	import { chatStore, conversationsStore, tabsStore, uiStore } from '$lib/stores';
+	import { tick } from 'svelte';
+
+	const carousel = useScrollCarousel();
+
+	let activeId = $derived(page.params.id ?? NEW_CHAT_TAB_ID);
+
+	let tabs = $derived(
+		tabsStore.openTabs.map((id) => ({
+			id,
+			isNewChat: id === NEW_CHAT_TAB_ID,
+			name:
+				id === NEW_CHAT_TAB_ID
+					? NEW_CHAT_LABEL
+					: (conversationsStore.conversations.find((c) => c.id === id)?.name ?? UNNAMED_CHAT_LABEL)
+		}))
+	);
+
+	// hide the New chat button while a new-chat tab is already open
+	let showNewChatButton = $derived(!tabsStore.openTabs.includes(NEW_CHAT_TAB_ID));
+
+	let loadingIds = $derived(new Set(chatStore.getAllLoadingChats()));
+
+	function handleClose(id: string) {
+		void tabsStore.close(id, activeId);
+	}
+
+	function handleStop(id: string, event: MouseEvent) {
+		event.stopPropagation();
+		void chatStore.stopGenerationForChat(id);
+	}
+
+	function handleAuxClick(id: string, event: MouseEvent) {
+		// middle-click closes, like browser tabs
+		if (event.button === 1) {
+			event.preventDefault();
+			handleClose(id);
+		}
+	}
+
+	let previousTabIds = new Set<string>();
+	let previousActiveId: string | null = null;
+
+	$effect(() => {
+		const currentIds = new Set(tabs.map((t) => t.id));
+		const hasAddedTab = tabs.some((t) => !previousTabIds.has(t.id));
+
+		previousTabIds = currentIds;
+
+		const activeChanged = activeId !== previousActiveId;
+
+		previousActiveId = activeId;
+
+		// scroll when the active tab changes (a click) or when a new tab is added
+		if (!hasAddedTab && !activeChanged) return;
+
+		// wait for the new tab to be laid out before scrolling to it
+		void tick().then(() => {
+			const el = carousel.scrollContainer?.querySelector<HTMLElement>(
+				`[${UI_DATA_ATTRS.ACTIVE_TAB}]`
+			);
+
+			if (el) {
+				carousel.scrollToCenter(el);
+			}
+		});
+	});
+</script>
+
+<nav
+	class="group sticky pl-1 top-0 z-10 hidden md:block chat-tabs-fade transition-[padding] duration-200 ease-in-out pt-3.25 {uiStore.isSidebarExpanded
+		? CHAT_TABS_MAX_WIDTH.EXPANDED_SIDEBAR
+		: CHAT_TABS_MAX_WIDTH.COLLAPSED_SIDEBAR}"
+	aria-label="Open conversations"
+>
+	<div class="relative">
+		<ScrollCarousel
+			class="h-10"
+			containerClass="flex h-10 min-w-0 items-center"
+			innerClass="items-center gap-1.25"
+			{carousel}
+		>
+			{#each tabs as tab (tab.id)}
+				<ChatTabsItem
+					{tab}
+					isActive={tab.id === activeId}
+					isLoading={loadingIds.has(tab.id)}
+					onActivate={(id) => tabsStore.activate(id)}
+					onClose={handleClose}
+					onStop={handleStop}
+					onAuxClick={handleAuxClick}
+				/>
+			{/each}
+
+			{#if showNewChatButton}
+				<ChatTabsNewChatButton onclick={() => void conversationsStore.openNewChat()} />
+			{/if}
+		</ScrollCarousel>
+
+		<div
+			class="pointer-events-none absolute inset-y-0 left-0 z-[5] w-8 bg-gradient-to-r from-background to-transparent transition-opacity {carousel.canScrollLeft
+				? 'opacity-100'
+				: 'opacity-0'}"
+		></div>
+		<div
+			class="pointer-events-none absolute inset-y-0 right-0 z-[5] w-8 bg-gradient-to-l from-background to-transparent transition-opacity {carousel.canScrollRight
+				? 'opacity-100'
+				: 'opacity-0'}"
+		></div>
+	</div>
+</nav>
+
+<style>
+	.chat-tabs-fade {
+		background: linear-gradient(
+			to bottom,
+			color-mix(in srgb, var(--background) 100%, transparent) 25%,
+			color-mix(in srgb, var(--background) 80%, transparent) 50%,
+			color-mix(in srgb, var(--background) 40%, transparent) 75%,
+			transparent 100%
+		);
+	}
+</style>

tools/ui/src/lib/components/app/chat/ChatTabs/ChatTabsItem.svelte

@@ -0,0 +1,156 @@
+<script lang="ts">
+	import { Loader2, Square, SquarePen, X } from '@lucide/svelte';
+	import * as Tooltip from '$lib/components/ui/tooltip';
+	import { cn } from '$lib/components/ui/utils';
+	import { ICON_CLASS_SM, ICON_CLASS_XS, ROUTES, UI_DATA_ATTRS } from '$lib/constants';
+	import { RouterService } from '$lib/services/router.service';
+
+	interface Tab {
+		id: string;
+		isNewChat: boolean;
+		name: string;
+	}
+
+	interface Props {
+		tab: Tab;
+		isActive?: boolean;
+		isLoading?: boolean;
+		onActivate?: (id: string) => void;
+		onClose?: (id: string) => void;
+		onStop?: (id: string, event: MouseEvent) => void;
+		onAuxClick?: (id: string, event: MouseEvent) => void;
+	}
+
+	let {
+		isActive = false,
+		isLoading = false,
+		onActivate,
+		onAuxClick,
+		onClose,
+		onStop,
+		tab
+	}: Props = $props();
+
+	let contentOpacity = $derived(isActive ? '' : 'opacity-45 group-hover:opacity-75');
+
+	let href = $derived(tab.isNewChat ? ROUTES.START : RouterService.chat(tab.id));
+
+	function handleActivate(event: MouseEvent) {
+		// let cmd/ctrl/middle-click fall through so the browser keeps its own
+		// behavior (open in a new window); route the plain click ourselves so the
+		// new-chat sentinel and history behave exactly like programmatic nav
+		if (event.metaKey || event.ctrlKey || event.button === 1) return;
+
+		event.preventDefault();
+		onActivate?.(tab.id);
+	}
+
+	// stop/close sit on top of the tab link; swallow their clicks so they do
+	// not also navigate
+	function handleActionClick(event: MouseEvent, action: () => void) {
+		event.preventDefault();
+		event.stopPropagation();
+		action();
+	}
+</script>
+
+<!-- the tab link covers the whole item; stop/close sit on top as siblings so
+     interactive elements are never nested inside the anchor -->
+<div
+	{...{ [UI_DATA_ATTRS.ACTIVE_TAB]: isActive ? 'true' : undefined }}
+	class={cn(
+		'relative flex h-8 max-w-52 min-w-0 shrink-0 items-center gap-1 rounded-lg pr-1 text-sm whitespace-nowrap border backdrop-blur-xl first:ml-2',
+		isLoading ? 'pl-1' : 'pl-3',
+		isActive
+			? 'bg-muted/60 border-border/10 shadow-sm text-accent-foreground hover:bg-primary/15'
+			: 'border-transparent hover:bg-primary/10 hover:border-border/10 hover:shadow-sm'
+	)}
+>
+	<a
+		{href}
+		class="absolute inset-0 z-0 rounded-lg"
+		onclick={handleActivate}
+		onauxclick={(e) => onAuxClick?.(tab.id, e)}
+		aria-current={isActive ? 'page' : undefined}
+		aria-label={tab.name}
+	></a>
+
+	{#if isLoading}
+		<Tooltip.Root>
+			<Tooltip.Trigger>
+				{#snippet child({ props })}
+					<button
+						{...props}
+						class="stop-button relative z-10 flex h-5 w-5 shrink-0 cursor-pointer items-center justify-center rounded-sm text-muted-foreground transition-colors hover:text-foreground"
+						onclick={(e) => handleActionClick(e, () => onStop?.(tab.id, e))}
+						aria-label="Stop generation"
+					>
+						<Loader2
+							class="loading-icon {ICON_CLASS_SM} animate-spin transition-opacity duration-300 {contentOpacity}"
+						/>
+
+						<Square
+							class="stop-icon hidden {ICON_CLASS_XS} fill-current text-destructive transition-opacity {contentOpacity}"
+						/>
+					</button>
+				{/snippet}
+			</Tooltip.Trigger>
+
+			<Tooltip.Content>
+				<p>Stop generation</p>
+			</Tooltip.Content>
+		</Tooltip.Root>
+	{/if}
+
+	{#if tab.isNewChat}
+		<SquarePen
+			class="pointer-events-none {ICON_CLASS_SM} shrink-0 transition-opacity {contentOpacity}"
+		/>
+	{/if}
+
+	<span class="pointer-events-none truncate transition-opacity {contentOpacity}">{tab.name}</span>
+
+	<Tooltip.Root>
+		<Tooltip.Trigger>
+			{#snippet child({ props })}
+				<button
+					{...props}
+					class={cn(
+						'relative z-10 flex h-5 w-5 shrink-0 cursor-pointer items-center justify-center rounded-sm text-muted-foreground transition-opacity hover:bg-foreground/10 hover:text-foreground',
+						contentOpacity
+					)}
+					onclick={(e) => handleActionClick(e, () => onClose?.(tab.id))}
+					aria-label="Close tab"
+				>
+					<X class={ICON_CLASS_SM} />
+				</button>
+			{/snippet}
+		</Tooltip.Trigger>
+
+		<Tooltip.Content>
+			<p>Close tab</p>
+		</Tooltip.Content>
+	</Tooltip.Root>
+</div>
+
+<style>
+	.stop-button {
+		:global(.stop-icon) {
+			display: none;
+		}
+
+		:global(.loading-icon) {
+			display: block;
+		}
+
+		&:is(:hover) {
+			:global(.stop-icon) {
+				display: block;
+			}
+
+			:global(.loading-icon) {
+				display: none;
+			}
+		}
+	}
+</style>

tools/ui/src/lib/components/app/chat/ChatTabs/ChatTabsNewChatButton.svelte

@@ -0,0 +1,30 @@
+<script lang="ts">
+	import { Plus } from '@lucide/svelte';
+	import * as Tooltip from '$lib/components/ui/tooltip';
+	import { ICON_CLASS_DEFAULT } from '$lib/constants';
+
+	interface Props {
+		onclick?: () => void;
+	}
+
+	let { onclick }: Props = $props();
+</script>
+
+<Tooltip.Root>
+	<Tooltip.Trigger>
+		{#snippet child({ props })}
+			<button
+				{...props}
+				class="backdrop-blur-lg flex h-8 w-8 mr-4 shrink-0 cursor-pointer items-center justify-center rounded-md transition-colors hover:bg-foreground/5"
+				{onclick}
+				aria-label="New chat"
+			>
+				<Plus class="{ICON_CLASS_DEFAULT} opacity-40 transition-opacity group-hover:opacity-100" />
+			</button>
+		{/snippet}
+	</Tooltip.Trigger>
+
+	<Tooltip.Content>
+		<p>New chat</p>
+	</Tooltip.Content>
+</Tooltip.Root>

tools/ui/src/lib/components/app/chat/index.ts

@@ -686,6 +686,18 @@ export { default as ChatMessageSystem } from './ChatMessages/ChatMessage/ChatMes
  */
 export { default as ChatScreen } from './ChatScreen/ChatScreen.svelte';
 
+/**
+ * **ChatTabs** - Browser-style tab bar for open conversations
+ *
+ * Horizontal strip of tabs rendered above ChatScreen in the chat layout,
+ * one per conversation tracked by tabsStore. The active tab follows the
+ * route's conversation id; clicking a tab navigates to it, middle-click or
+ * the close button closes it (switching to the left neighbor when closing
+ * the active tab), and a trailing "+" button starts a new chat. Shows a
+ * spinner on tabs with a running generation. Desktop-only.
+ */
+export { default as ChatTabs } from './ChatTabs/ChatTabs.svelte';
+
 /**
  * Visual overlay displayed when user drags files over the chat screen.
  * Shows drop zone indicator to guide users where to release files.

tools/ui/src/lib/components/app/misc/HorizontalScrollCarousel.svelte

@@ -1,96 +0,0 @@
-<script lang="ts">
-	import { ChevronLeft, ChevronRight } from '@lucide/svelte';
-	import { ICON_CLASS_DEFAULT } from '$lib/constants';
-	import type { Snippet } from 'svelte';
-
-	interface Props {
-		class?: string;
-		children?: Snippet;
-		gapSize?: string;
-		onScrollableChange?: (isScrollable: boolean) => void;
-	}
-
-	let { children, class: className = '', gapSize = '3', onScrollableChange }: Props = $props();
-
-	let canScrollLeft = $state(false);
-	let canScrollRight = $state(false);
-	let scrollContainer: HTMLDivElement | undefined = $state();
-
-	function scrollLeft(event?: MouseEvent) {
-		event?.stopPropagation();
-		event?.preventDefault();
-
-		if (!scrollContainer) return;
-
-		scrollContainer.scrollBy({ behavior: 'smooth', left: scrollContainer.clientWidth * -0.67 });
-	}
-
-	function scrollRight(event?: MouseEvent) {
-		event?.stopPropagation();
-		event?.preventDefault();
-
-		if (!scrollContainer) return;
-
-		scrollContainer.scrollBy({ behavior: 'smooth', left: scrollContainer.clientWidth * 0.67 });
-	}
-
-	function updateScrollButtons() {
-		if (!scrollContainer) return;
-
-		const { clientWidth, scrollLeft, scrollWidth } = scrollContainer;
-
-		canScrollLeft = scrollLeft > 0;
-		canScrollRight = scrollLeft < scrollWidth - clientWidth - 1;
-
-		const isScrollable = scrollWidth > clientWidth;
-
-		onScrollableChange?.(isScrollable);
-	}
-
-	export function resetScroll() {
-		if (scrollContainer) {
-			scrollContainer.scrollLeft = 0;
-			setTimeout(() => {
-				updateScrollButtons();
-			}, 0);
-		}
-	}
-
-	$effect(() => {
-		if (!scrollContainer) return;
-
-		const observer = new ResizeObserver(() => updateScrollButtons());
-
-		observer.observe(scrollContainer);
-
-		return () => observer.disconnect();
-	});
-</script>
-
-<div class="relative {className}">
-	<button
-		class="absolute top-1/2 left-4 z-10 flex h-6 w-6 -translate-y-1/2 items-center justify-center rounded-full bg-background/25 shadow-md backdrop-blur-xs transition-opacity hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0"
-		onclick={scrollLeft}
-		disabled={!canScrollLeft}
-		aria-label="Scroll left"
-	>
-		<ChevronLeft class={ICON_CLASS_DEFAULT} />
-	</button>
-
-	<div
-		class="scrollbar-hide flex items-start gap-{gapSize} overflow-x-auto"
-		bind:this={scrollContainer}
-		onscroll={updateScrollButtons}
-	>
-		{@render children?.()}
-	</div>
-
-	<button
-		class="absolute top-1/2 right-4 z-10 flex h-6 w-6 -translate-y-1/2 items-center justify-center rounded-full bg-background/25 shadow-md backdrop-blur-xs transition-opacity hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0"
-		onclick={scrollRight}
-		disabled={!canScrollRight}
-		aria-label="Scroll right"
-	>
-		<ChevronRight class={ICON_CLASS_DEFAULT} />
-	</button>
-</div>

tools/ui/src/lib/components/app/misc/ScrollCarousel.svelte

@@ -0,0 +1,131 @@
+<script lang="ts">
+	import { ChevronLeft, ChevronRight } from '@lucide/svelte';
+	import { cn } from '$lib/components/ui/utils';
+	import { ICON_CLASS_DEFAULT } from '$lib/constants';
+	import { ScrollCarouselVariant } from '$lib/enums';
+	import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
+	import type { Snippet } from 'svelte';
+
+	interface Props {
+		children: Snippet;
+		/** External carousel hook for callers that need to drive it (e.g. scrollToCenter). */
+		carousel?: ReturnType<typeof useScrollCarousel>;
+		/** Classes for the outer relative wrapper. */
+		class?: string;
+		/** Classes for the scrollable overflow container. */
+		containerClass?: string;
+		/** Classes for the min-w-max content wrapper. */
+		innerClass?: string;
+		/** Tailwind gap class applied to the content wrapper. */
+		gapSize?: string;
+		/** Show the arrows whenever the content overflows, even without hover. */
+		alwaysShowArrows?: boolean;
+		/** Arrow placement and styling. */
+		variant?: ScrollCarouselVariant;
+	}
+
+	let {
+		alwaysShowArrows = false,
+		carousel: externalCarousel,
+		children,
+		class: className = '',
+		containerClass = '',
+		gapSize = '3',
+		innerClass = '',
+		variant = ScrollCarouselVariant.TOP
+	}: Props = $props();
+
+	const internalCarousel = useScrollCarousel();
+	const carousel = $derived(externalCarousel ?? internalCarousel);
+
+	const isCenter = $derived(variant === ScrollCarouselVariant.CENTER);
+
+	function scrollLeft(event?: MouseEvent) {
+		event?.stopPropagation();
+		event?.preventDefault();
+
+		const container = carousel.scrollContainer;
+
+		if (!container) return;
+
+		container.scrollBy({ behavior: 'smooth', left: -(container.clientWidth * 0.67) });
+	}
+
+	function scrollRight(event?: MouseEvent) {
+		event?.stopPropagation();
+		event?.preventDefault();
+
+		const container = carousel.scrollContainer;
+
+		if (!container) return;
+
+		container.scrollBy({ behavior: 'smooth', left: container.clientWidth * 0.67 });
+	}
+
+	export function resetScroll() {
+		const container = carousel.scrollContainer;
+
+		if (!container) return;
+
+		container.scrollLeft = 0;
+		setTimeout(() => carousel.updateScrollButtons(), 0);
+	}
+</script>
+
+<div
+	class={cn('group relative', !isCenter && 'flex items-center', className)}
+	style={!isCenter ? 'scroll-padding: 1rem;' : undefined}
+>
+	<button
+		class={cn(
+			'absolute z-10 flex h-6 w-6 items-center justify-center rounded-full shadow-md transition-opacity',
+			isCenter
+				? 'top-1/2 left-4 -translate-y-1/2 bg-background/25 backdrop-blur-xs hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0'
+				: 'left-2 bg-muted backdrop-blur-sm hover:bg-accent',
+			!isCenter &&
+				(carousel.canScrollLeft
+					? alwaysShowArrows
+						? 'opacity-100'
+						: 'opacity-0 group-hover:opacity-100'
+					: 'pointer-events-none opacity-0')
+		)}
+		{...isCenter ? { disabled: !carousel.canScrollLeft } : {}}
+		onclick={scrollLeft}
+		aria-label="Scroll left"
+	>
+		<ChevronLeft class={ICON_CLASS_DEFAULT} />
+	</button>
+
+	<div
+		class={cn('scrollbar-hide overflow-x-auto', containerClass)}
+		bind:this={carousel.scrollContainer}
+		onscroll={carousel.updateScrollButtons}
+	>
+		<div
+			class={cn('flex min-w-max', isCenter && 'items-start', `gap-${gapSize}`, innerClass)}
+			bind:this={carousel.contentContainer}
+		>
+			{@render children?.()}
+		</div>
+	</div>
+
+	<button
+		class={cn(
+			'absolute z-10 flex h-6 w-6 items-center justify-center rounded-full shadow-md transition-opacity',
+			isCenter
+				? 'top-1/2 right-4 -translate-y-1/2 bg-background/25 backdrop-blur-xs hover:bg-background/45 disabled:pointer-events-none disabled:opacity-0'
+				: 'right-2 bg-muted backdrop-blur-sm hover:bg-accent',
+			!isCenter &&
+				(carousel.canScrollRight
+					? alwaysShowArrows
+						? 'opacity-100'
+						: 'opacity-0 group-hover:opacity-100'
+					: 'pointer-events-none opacity-0')
+		)}
+		{...isCenter ? { disabled: !carousel.canScrollRight } : {}}
+		onclick={scrollRight}
+		aria-label="Scroll right"
+	>
+		<ChevronRight class={ICON_CLASS_DEFAULT} />
+	</button>
+</div>

tools/ui/src/lib/components/app/misc/index.ts

@@ -21,13 +21,6 @@
  */
 export { default as ConversationSelection } from './ConversationSelection.svelte';
 
-/**
- * Horizontal scrollable carousel with navigation arrows.
- * Used for displaying items in a horizontally scrollable container
- * with left/right navigation buttons that appear on hover.
- */
-export { default as HorizontalScrollCarousel } from './HorizontalScrollCarousel.svelte';
-
 /**
  * **TruncatedText** - Text with ellipsis and tooltip
  *
@@ -44,6 +37,13 @@ export { default as TruncatedText } from './TruncatedText.svelte';
  */
 export { default as KeyboardShortcutInfo } from './KeyboardShortcutInfo.svelte';
 
+/**
+ * **ScrollCarousel** - Feature/carousel with center-aligned overflow controls
+ *
+ * Horizontal scrollable container with arrows that center the focused item.
+ */
+export { default as ScrollCarousel } from './ScrollCarousel.svelte';
+
 /**
  * **CodeBlockActions** - Actions bar for code blocks (copy, preview)
  *

tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte

@@ -14,7 +14,7 @@
 	import { useKeyboardShortcuts } from '$lib/hooks/use-keyboard-shortcuts.svelte';
 	import { useMarqueeSelection } from '$lib/hooks/use-marquee-selection.svelte';
 	import { RouterService } from '$lib/services/router.service';
-	import { chatStore, conversationsStore, deviceStore, settingsStore } from '$lib/stores';
+	import { chatStore, conversationsStore, deviceStore, settingsStore, uiStore } from '$lib/stores';
 	import { buildConversationTree } from '$lib/utils';
 	import { circIn } from 'svelte/easing';
 	import { SvelteSet } from 'svelte/reactivity';
@@ -31,30 +31,29 @@
 		toggleSidebar: () => toggleExpandedMode()
 	});
 
-	let isExpandedMode = $state(false);
 	let hoveredTooltip = $state<string | null>(null);
 	let logoHovered = $state(false);
 
-	const isStripExpanded = $derived(isExpandedMode || hoveredTooltip !== null);
+	const isStripExpanded = $derived(uiStore.isSidebarExpanded || hoveredTooltip !== null);
 	const isOnMobile = $derived(deviceStore.isMobile);
 	const alwaysShowOnDesktop = $derived(settingsStore.config.alwaysShowSidebarOnDesktop as boolean);
 
 	$effect(() => {
 		if (alwaysShowOnDesktop && !isOnMobile) {
-			isExpandedMode = true;
+			uiStore.isSidebarExpanded = true;
 		}
 	});
 
 	function toggleExpandedMode() {
-		isExpandedMode = !isExpandedMode;
+		uiStore.isSidebarExpanded = !uiStore.isSidebarExpanded;
 
-		if (!isExpandedMode) {
+		if (!uiStore.isSidebarExpanded) {
 			hoveredTooltip = null;
 		}
 	}
 
 	$effect(() => {
-		if (!isExpandedMode) {
+		if (!uiStore.isSidebarExpanded) {
 			isSearchModeActive = false;
 			searchQuery = '';
 
@@ -66,7 +65,7 @@
 
 	$effect(() => {
 		if (deviceStore.isMobile && page.url.hash.includes(ROUTES.SEARCH)) {
-			isExpandedMode = false;
+			uiStore.isSidebarExpanded = false;
 		}
 	});
 
@@ -294,7 +293,7 @@
 		}
 
 		pendingCollapse = setTimeout(() => {
-			isExpandedMode = false;
+			uiStore.isSidebarExpanded = false;
 			pendingCollapse = null;
 		}, 100);
 	}
@@ -314,7 +313,7 @@
 		class={[
 			'fixed md:sticky top-2 left-2 md:left-0 md:ml-2 md:mt-2 pt-2 z-10 w-[calc(100dvw-1rem)]',
 			'md:h-[calc(100dvh-1.125rem)]',
-			isExpandedMode &&
+			uiStore.isSidebarExpanded &&
 				(deviceStore.isStandalone
 					? 'h-[calc(100dvh-2rem)]'
 					: deviceStore.isIOSDevice
@@ -323,9 +322,9 @@
 			'rounded-3xl md:rounded-2xl',
 			'flex flex-col justify-between',
 			'md:transition-[width,padding] duration-200 ease-out',
-			isStripExpanded && 'md:w-72 md:bg-muted/60 md:backdrop-blur-xl border-border shadow-md',
+			isStripExpanded && 'md:w-72 md:bg-muted/60 md:backdrop-blur-xl shadow-md',
 			!isStripExpanded && 'md:w-12',
-			isExpandedMode && 'is-expanded'
+			uiStore.isSidebarExpanded && 'is-expanded'
 		]}
 	>
 		<div class="px-2 flex items-center justify-between">
@@ -337,24 +336,26 @@
 				onmouseleave={() => (logoHovered = false)}
 			>
 				<ActionIcon
-					icon={!isExpandedMode && logoHovered && innerWidth > 768 ? PanelLeftOpen : Logo}
+					icon={!uiStore.isSidebarExpanded && logoHovered && innerWidth > 768
+						? PanelLeftOpen
+						: Logo}
 					size="lg"
 					iconSize="h-4.5 w-4.5 md:h-4 md:w-4"
-					class="{isExpandedMode
+					class="{uiStore.isSidebarExpanded
 						? 'bg-muted! md:bg-foreground/5!'
 						: 'bg-transparent!'} md:h-9 md:w-9 h-10 w-10 rounded-full md:hover:bg-foreground/10! pointer-events-auto"
-					href={isExpandedMode ? ROUTES.START : undefined}
-					onclick={isExpandedMode ? undefined : toggleExpandedMode}
-					tooltip={isExpandedMode ? undefined : 'Open Sidebar'}
+					href={uiStore.isSidebarExpanded ? ROUTES.START : undefined}
+					onclick={uiStore.isSidebarExpanded ? undefined : toggleExpandedMode}
+					tooltip={uiStore.isSidebarExpanded ? undefined : 'Open Sidebar'}
 					tooltipSide={TooltipSide.RIGHT}
-					ariaLabel={isExpandedMode ? 'Go to start' : 'Expand navigation'}
+					ariaLabel={uiStore.isSidebarExpanded ? 'Go to start' : 'Expand navigation'}
 				/>
 			</div>
 
-			{#if isOnMobile || (isExpandedMode && !alwaysShowOnDesktop)}
+			{#if isOnMobile || (uiStore.isSidebarExpanded && !alwaysShowOnDesktop)}
 				<div
 					class="flex items-center transition-all duration-150 ease-out {deviceStore.isMobile &&
-					!isExpandedMode
+					!uiStore.isSidebarExpanded
 						? 'opacity-0 h-0!'
 						: ''}"
 					in:fade={{ delay: 50, duration: 150, easing: circIn }}
@@ -377,12 +378,12 @@
 		<div
 			class="mt-2 flex min-h-0 flex-1 flex-col gap-4 md:gap-1 {deviceStore.isMobile
 				? 'transition-[opacity,height] duration-200 ease-out'
-				: ''} {deviceStore.isMobile && !isExpandedMode ? 'opacity-0 !h-0' : ''}"
+				: ''} {deviceStore.isMobile && !uiStore.isSidebarExpanded ? 'opacity-0 !h-0' : ''}"
 			in:fade={{ duration: 200 }}
 			out:fade={{ duration: 200 }}
 		>
 			<SidebarNavigationActions
-				isExpandedMode={innerWidth > 768 ? isExpandedMode : true}
+				isExpandedMode={innerWidth > 768 ? uiStore.isSidebarExpanded : true}
 				class="px-2"
 				bind:isSearchModeActive
 				bind:searchQuery
@@ -391,7 +392,7 @@
 					searchQuery = '';
 				}}
 				onSearchClick={() => {
-					isExpandedMode = true;
+					uiStore.isSidebarExpanded = true;
 					isSearchModeActive = true;
 				}}
 				onNewChat={() => {
@@ -401,7 +402,7 @@
 				}}
 			/>
 
-			{#if isExpandedMode || isOnMobile}
+			{#if uiStore.isSidebarExpanded || isOnMobile}
 				<div class="flex min-h-0 flex-1 flex-col overflow-y-auto">
 					<SidebarNavigationConversationList
 						class="px-2"

tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigationActions.svelte

@@ -11,8 +11,8 @@
 		ROUTES,
 		SIDEBAR_ACTIONS_ITEMS
 	} from '$lib/constants';
-	import { TooltipSide } from '$lib/enums';
-	import { deviceStore } from '$lib/stores';
+	import { SidebarAction, TooltipSide } from '$lib/enums';
+	import { conversationsStore, deviceStore } from '$lib/stores';
 	import type { Component } from 'svelte';
 	import { onMount } from 'svelte';
 	import { circIn } from 'svelte/easing';
@@ -109,14 +109,20 @@
 			{@const isActive = isItemActive(item)}
 			{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
 			{@const itemHref = isSearchOnMobile ? ROUTES.SEARCH : item.route}
-			{@const itemOnClick = item.route
-				? () => {
-						onNewChat?.();
-						goto(item.route!);
-					}
-				: isSearchOnMobile
-					? undefined
-					: onSearchClick}
+			{@const itemOnClick =
+				item.action === SidebarAction.NEW_CHAT
+					? () => {
+							onNewChat?.();
+							void conversationsStore.openNewChat();
+						}
+					: item.route
+						? () => {
+								onNewChat?.();
+								goto(item.route!);
+							}
+						: isSearchOnMobile
+							? undefined
+							: onSearchClick}
 			{@const itemTransition = {
 				delay: !initialized ? i * ICON_STRIP_TRANSITION_DELAY_MULTIPLIER : 0,
 				duration: ICON_STRIP_TRANSITION_DURATION,
@@ -157,14 +163,20 @@
 		{#each SIDEBAR_ACTIONS_ITEMS as item, i (item.tooltip)}
 			{@const isActive = isItemActive(item)}
 			{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
-			{@const itemOnClick = item.route
-				? () => {
-						onNewChat?.();
-						goto(item.route!);
-					}
-				: isSearchOnMobile
-					? undefined
-					: onSearchClick}
+			{@const itemOnClick =
+				item.action === SidebarAction.NEW_CHAT
+					? () => {
+							onNewChat?.();
+							void conversationsStore.openNewChat();
+						}
+					: item.route
+						? () => {
+								onNewChat?.();
+								goto(item.route!);
+							}
+						: isSearchOnMobile
+							? undefined
+							: onSearchClick}
 			{@const itemTransition = {
 				delay: !initialized ? i * ICON_STRIP_TRANSITION_DELAY_MULTIPLIER : 0,
 				duration: ICON_STRIP_TRANSITION_DURATION,

tools/ui/src/lib/components/app/settings/SettingsChatMobileHeader.svelte

@@ -1,5 +1,6 @@
 <script lang="ts">
-	import { ChevronLeft, ChevronRight, Settings } from '@lucide/svelte';
+	import { Settings } from '@lucide/svelte';
+	import { ScrollCarousel } from '$lib/components/app';
 	import { ICON_CLASS_DEFAULT, UI_DATA_ATTRS } from '$lib/constants';
 	import { BooleanString } from '$lib/enums';
 	import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
@@ -44,70 +45,42 @@
 	</div>
 
 	<div class="border-b border-border/30 py-2">
-		<div class="relative flex items-center" style="scroll-padding: 1rem;">
-			<button
-				class="absolute left-2 z-10 flex h-6 w-6 items-center justify-center rounded-full bg-muted shadow-md backdrop-blur-sm transition-opacity hover:bg-accent {carousel.canScrollLeft
-					? 'opacity-100'
-					: 'pointer-events-none opacity-0'}"
-				onclick={carousel.scrollLeft}
-				aria-label="Scroll left"
-			>
-				<ChevronLeft class={ICON_CLASS_DEFAULT} />
-			</button>
-
-			<div
-				class="scrollbar-hide overflow-x-auto py-2"
-				bind:this={carousel.scrollContainer}
-				onscroll={carousel.updateScrollButtons}
-			>
-				<div class="flex min-w-max gap-2">
-					{#each sections as section (section.title)}
-						{#if getHref}
-							<a
-								class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap no-underline transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
-									section
-								)
-									? 'bg-accent text-accent-foreground'
-									: 'text-muted-foreground'}"
-								{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
-								href={getHref(section)}
-								onclick={(e: MouseEvent) => {
-									carousel.scrollToCenter(e.currentTarget as HTMLElement);
-								}}
-							>
-								<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
-								<span>{section.title}</span>
-							</a>
-						{:else}
-							<button
-								class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
-									section
-								)
-									? 'bg-accent text-accent-foreground'
-									: 'text-muted-foreground'}"
-								{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
-								onclick={(e: MouseEvent) => {
-									onSectionChange?.(section.title);
-									carousel.scrollToCenter(e.currentTarget as HTMLElement);
-								}}
-							>
-								<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
-								<span>{section.title}</span>
-							</button>
-						{/if}
-					{/each}
-				</div>
-			</div>
-
-			<button
-				class="absolute right-2 z-10 flex h-6 w-6 items-center justify-center rounded-full bg-muted shadow-md backdrop-blur-sm transition-opacity hover:bg-accent {carousel.canScrollRight
-					? 'opacity-100'
-					: 'pointer-events-none opacity-0'}"
-				onclick={carousel.scrollRight}
-				aria-label="Scroll right"
-			>
-				<ChevronRight class={ICON_CLASS_DEFAULT} />
-			</button>
-		</div>
+		<ScrollCarousel {carousel} alwaysShowArrows containerClass="py-2" innerClass="gap-2">
+			{#each sections as section (section.title)}
+				{#if getHref}
+					<a
+						class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap no-underline transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
+							section
+						)
+							? 'bg-accent text-accent-foreground'
+							: 'text-muted-foreground'}"
+						{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
+						href={getHref(section)}
+						onclick={(e: MouseEvent) => {
+							carousel.scrollToCenter(e.currentTarget as HTMLElement);
+						}}
+					>
+						<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
+						<span>{section.title}</span>
+					</a>
+				{:else}
+					<button
+						class="flex cursor-pointer items-center gap-2 rounded-lg px-3 py-2 text-sm whitespace-nowrap transition-colors first:ml-4 last:mr-4 hover:bg-accent {isActive(
+							section
+						)
+							? 'bg-accent text-accent-foreground'
+							: 'text-muted-foreground'}"
+						{...{ [UI_DATA_ATTRS.ACTIVE]: isActive(section) }}
+						onclick={(e: MouseEvent) => {
+							onSectionChange?.(section.title);
+							carousel.scrollToCenter(e.currentTarget as HTMLElement);
+						}}
+					>
+						<section.icon class="{ICON_CLASS_DEFAULT} flex-shrink-0" />
+						<span>{section.title}</span>
+					</button>
+				{/if}
+			{/each}
+		</ScrollCarousel>
 	</div>
 </div>

tools/ui/src/lib/constants/chat-tabs.constants.ts

@@ -0,0 +1,18 @@
+/** Sentinel tab id for the bare `#/` new-chat screen */
+export const NEW_CHAT_TAB_ID = 'new-chat';
+
+/** Label shown for the new-chat sentinel tab. */
+export const NEW_CHAT_LABEL = 'New chat';
+
+/** Fallback label for conversations without an auto-generated title. */
+export const UNNAMED_CHAT_LABEL = 'Chat';
+
+/**
+ * Tab bar max width so it stays clear of the sidebar strip. The expanded strip
+ * is `md:w-72` and the collapsed one `md:w-12`; these hold the fully tuned
+ * `max-w-[calc(100vw-?rem)]` classes so the offset has a single source.
+ */
+export const CHAT_TABS_MAX_WIDTH = {
+	COLLAPSED_SIDEBAR: 'max-w-[calc(100vw-5rem)]',
+	EXPANDED_SIDEBAR: 'max-w-[calc(100vw-20rem)]'
+} as const;

tools/ui/src/lib/constants/css-classes.constants.ts

@@ -26,5 +26,11 @@ export const CHAT_INPUT_FOCUS_SELECTOR =
 /** Default Tailwind size class for inline icon components (lucide, etc.). */
 export const ICON_CLASS_DEFAULT = 'h-4 w-4';
 
+/** Small Tailwind size class for inline icons. */
+export const ICON_CLASS_SM = 'h-3.5 w-3.5';
+
+/** Extra-small Tailwind size class for inline icons. */
+export const ICON_CLASS_XS = 'h-3 w-3';
+
 /** Icon size + spinning animation; used for live-streaming tool indicators. */
 export const ICON_CLASS_SPIN = 'h-4 w-4 animate-spin';

tools/ui/src/lib/constants/index.ts

@@ -4,6 +4,7 @@
 export * from './agentic.constants';
 export * from './api-endpoints.constants';
 export * from './app.constants';
+export * from './chat-tabs.constants';
 export * from './database.constants';
 export * from './reasoning-effort.constants';
 export * from './recommended-mcp-servers.constants';

tools/ui/src/lib/constants/routes.constants.ts

@@ -4,8 +4,6 @@ export const URL_PARAMS = {
 	LOAD: 'load',
 	/** Model to select. */
 	MODEL: 'model',
-	/** Start a new chat. */
-	NEW_CHAT: 'new_chat',
 	/** Prompt to send on arrival. */
 	QUERY: 'q'
 } as const;
@@ -15,8 +13,6 @@ export const ROUTES = {
 	CHAT: '#/chat',
 	/** MCP servers. */
 	MCP_SERVERS: '#/mcp-servers',
-	/** New chat — root with new chat query param. */
-	NEW_CHAT: `?${URL_PARAMS.NEW_CHAT}=true#/`,
 	/** Search — mobile-only full-page conversation search. */
 	SEARCH: '#/search',
 	/** Settings base — for dynamic settings URLs use RouterService. */

tools/ui/src/lib/constants/settings-keys.constants.ts

@@ -11,6 +11,7 @@ export const SETTINGS_KEYS = {
 	API_KEY: 'apiKey',
 	AUTO_MIC_ON_EMPTY: 'autoMicOnEmpty',
 	BACKEND_SAMPLING: 'backend_sampling',
+	CONVERSATION_TABS: 'conversationTabs',
 	COPY_TEXT_ATTACHMENTS_AS_PLAIN_TEXT: 'copyTextAttachmentsAsPlainText',
 	CUSTOM_CSS: 'customCss',
 	// PY_INTERPRETER_ENABLED: 'pyInterpreterEnabled',

tools/ui/src/lib/constants/settings.constants.ts

@@ -255,6 +255,13 @@ export const SETTINGS_REGISTRY: SettingsSectionEntry[] = [
 				label: 'Always show sidebar on desktop',
 				type: SettingsFieldType.CHECKBOX
 			},
+			{
+				defaultValue: true,
+				help: 'Show open chats as browser-style tabs above the conversation, one per open chat. When disabled, only one chat is shown at a time.',
+				key: SETTINGS_KEYS.CONVERSATION_TABS,
+				label: 'Conversation tabs',
+				type: SettingsFieldType.CHECKBOX
+			},
 			{
 				defaultValue: false,
 				help: 'Display full raw model identifiers (e.g. "ggml-org/GLM-4.7-Flash-GGUF:Q8_0") instead of parsed names with badges.',

tools/ui/src/lib/constants/storage.constants.ts

@@ -22,6 +22,7 @@ export const DISABLED_TOOLS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.disabledTool
 export const DISABLED_TOOL_KEYS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.disabledToolKeys`;
 export const FAVORITE_MODELS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.favoriteModels`;
 export const REASONING_EFFORT_DEFAULT_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.reasoningEffortDefault`;
+export const CONVERSATION_TABS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.conversationTabs`;
 export const USER_OVERRIDES_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.userOverrides`;
 export const DISMISSED_RECOMMENDED_MCP_SERVERS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.dismissedRecommendedMcpServers`;
 

tools/ui/src/lib/constants/ui.constants.ts

@@ -1,7 +1,7 @@
 import { ROUTES } from './routes.constants';
 import { Package, Search, Settings, SquarePen } from '@lucide/svelte';
 import McpLogo from '$lib/components/app/mcp/McpLogo.svelte';
-import { ToolSource } from '$lib/enums/tools.enums';
+import { SidebarAction, ToolSource } from '$lib/enums';
 import type { DesktopIconStripItem } from '$lib/types';
 
 export const FORK_TREE_DEPTH_PADDING = 8;
@@ -10,6 +10,7 @@ export const SYSTEM_MESSAGE_PLACEHOLDER = 'System message';
 /** Data attributes for app-level DOM contracts. */
 export const UI_DATA_ATTRS = {
 	ACTIVE: 'data-active',
+	ACTIVE_TAB: 'data-active-tab',
 	CONVERSATION_ROW: 'data-conversation-row',
 	HIGHLIGHT_THEME_PREVIEW: 'data-highlight-theme-preview',
 	PICKER_INDEX: 'data-picker-index',
@@ -55,7 +56,12 @@ export const ICON_STRIP_TRANSITION_DELAY_MULTIPLIER = 50;
 export const MAX_HEIGHT_CODE_BLOCK = '22rem';
 
 export const SIDEBAR_ACTIONS_ITEMS: DesktopIconStripItem[] = [
-	{ icon: SquarePen, keys: ['shift', 'cmd', 'o'], route: ROUTES.NEW_CHAT, tooltip: 'New chat' },
+	{
+		action: SidebarAction.NEW_CHAT,
+		icon: SquarePen,
+		keys: ['shift', 'cmd', 'o'],
+		tooltip: 'New chat'
+	},
 	{ icon: Search, keys: ['cmd', 'k'], tooltip: 'Search' },
 	{
 		activeRouteId: '/mcp-servers',

tools/ui/src/lib/enums/index.ts

@@ -79,6 +79,8 @@ export {
 	ColorMode,
 	HtmlInputType,
 	McpPromptVariant,
+	ScrollCarouselVariant,
+	SidebarAction,
 	TooltipSide,
 	UrlProtocol
 } from './ui.enums';

tools/ui/src/lib/enums/keyboard.enums.ts

@@ -7,6 +7,8 @@ export enum KeyboardKey {
 	ARROW_RIGHT = 'ArrowRight',
 	ARROW_UP = 'ArrowUp',
 	B_LOWER = 'b',
+	BRACKET_LEFT = 'BracketLeft',
+	BRACKET_RIGHT = 'BracketRight',
 	D_LOWER = 'd',
 	D_UPPER = 'D',
 	E_UPPER = 'E',
@@ -16,5 +18,7 @@ export enum KeyboardKey {
 	O_LOWER = 'o',
 	O_UPPER = 'O',
 	SPACE = ' ',
-	TAB = 'Tab'
+	TAB = 'Tab',
+	X_LOWER = 'x',
+	X_UPPER = 'X'
 }

tools/ui/src/lib/enums/ui.enums.ts

@@ -11,6 +11,21 @@ export enum TooltipSide {
 	TOP = 'top'
 }
 
+/**
+ * ScrollCarousel arrow placement.
+ */
+export enum ScrollCarouselVariant {
+	CENTER = 'center',
+	TOP = 'top'
+}
+
+/**
+ * Sidebar icon strip actions handled directly by the sidebar.
+ */
+export enum SidebarAction {
+	NEW_CHAT = 'new-chat'
+}
+
 /**
  * MCP prompt display variant
  */

tools/ui/src/lib/hooks/use-keyboard-shortcuts.svelte.ts

@@ -1,6 +1,7 @@
-import { goto } from '$app/navigation';
-import { ROUTES } from '$lib/constants';
+import { page } from '$app/state';
+import { NEW_CHAT_TAB_ID } from '$lib/constants';
 import { KeyboardKey } from '$lib/enums';
+import { conversationsStore, settingsStore, tabsStore } from '$lib/stores';
 
 interface KeyboardShortcutsCallbacks {
 	activateSearchMode?: () => void;
@@ -9,6 +10,8 @@ interface KeyboardShortcutsCallbacks {
 	deleteActiveConversation?: () => void;
 	navigateToPrevConversation?: () => void;
 	navigateToNextConversation?: () => void;
+	navigateToPrevTab?: () => void;
+	navigateToNextTab?: () => void;
 	toggleSidebar?: () => void;
 }
 
@@ -34,7 +37,7 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
 		) {
 			event.preventDefault();
 
-			goto(ROUTES.NEW_CHAT);
+			void conversationsStore.openNewChat();
 		}
 
 		if (event.shiftKey && isCmdOrCtrl && event.key === KeyboardKey.E_UPPER) {
@@ -42,6 +45,28 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
 			callbacks.editActiveConversation?.();
 		}
 
+		if (
+			event.shiftKey &&
+			isCmdOrCtrl &&
+			(event.key === KeyboardKey.X_LOWER || event.key === KeyboardKey.X_UPPER)
+		) {
+			// several components register this shortcut; only let the first handler
+			// act so the synchronous navigation does not cascade-close every tab
+			if (event.defaultPrevented) return;
+
+			// close-tab only makes sense with conversation tabs enabled
+			if (!settingsStore.config.conversationTabs) return;
+
+			event.preventDefault();
+
+			const activeId =
+				page.params.id ?? (page.route.id === '/(chat)' ? NEW_CHAT_TAB_ID : undefined);
+
+			if (activeId) {
+				void tabsStore.close(activeId, activeId);
+			}
+		}
+
 		if (
 			isCmdOrCtrl &&
 			event.shiftKey &&
@@ -60,6 +85,16 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
 			event.preventDefault();
 			callbacks.navigateToNextConversation?.();
 		}
+
+		if (isCmdOrCtrl && event.altKey && event.shiftKey && event.code === KeyboardKey.BRACKET_LEFT) {
+			event.preventDefault();
+			callbacks.navigateToPrevTab?.();
+		}
+
+		if (isCmdOrCtrl && event.altKey && event.shiftKey && event.code === KeyboardKey.BRACKET_RIGHT) {
+			event.preventDefault();
+			callbacks.navigateToNextTab?.();
+		}
 	}
 
 	return { handleKeydown };

tools/ui/src/lib/hooks/use-scroll-carousel.svelte.ts

@@ -2,6 +2,7 @@ export function useScrollCarousel() {
 	let canScrollLeft = $state(false);
 	let canScrollRight = $state(false);
 	let scrollContainer = $state<HTMLDivElement | undefined>();
+	let contentContainer = $state<HTMLDivElement | undefined>();
 
 	function scrollToCenter(element: HTMLElement) {
 		if (!scrollContainer) return;
@@ -15,18 +16,6 @@ export function useScrollCarousel() {
 		scrollContainer.scrollBy({ behavior: 'smooth', left: scrollOffset });
 	}
 
-	function scrollLeft() {
-		if (!scrollContainer) return;
-
-		scrollContainer.scrollBy({ behavior: 'smooth', left: -250 });
-	}
-
-	function scrollRight() {
-		if (!scrollContainer) return;
-
-		scrollContainer.scrollBy({ behavior: 'smooth', left: 250 });
-	}
-
 	function updateScrollButtons() {
 		if (!scrollContainer) return;
 
@@ -36,10 +25,21 @@ export function useScrollCarousel() {
 		canScrollRight = sl < scrollWidth - clientWidth - 1;
 	}
 
+	// Re-evaluate arrow visibility whenever the container or its content resizes,
+	// otherwise the arrows may not appear when overflowing items are added (e.g. new
+	// tabs/attachments) and the user has not scrolled yet.
 	$effect(() => {
-		if (scrollContainer) {
-			updateScrollButtons();
-		}
+		if (!scrollContainer) return;
+
+		updateScrollButtons();
+
+		const observer = new ResizeObserver(() => updateScrollButtons());
+
+		observer.observe(scrollContainer);
+
+		if (contentContainer) observer.observe(contentContainer);
+
+		return () => observer.disconnect();
 	});
 
 	return {
@@ -49,14 +49,18 @@ export function useScrollCarousel() {
 		get canScrollRight() {
 			return canScrollRight;
 		},
+		get contentContainer() {
+			return contentContainer;
+		},
+		set contentContainer(el: HTMLDivElement | undefined) {
+			contentContainer = el;
+		},
 		get scrollContainer() {
 			return scrollContainer;
 		},
 		set scrollContainer(el: HTMLDivElement | undefined) {
 			scrollContainer = el;
 		},
-		scrollLeft,
-		scrollRight,
 		scrollToCenter,
 		updateScrollButtons
 	};

tools/ui/src/lib/services/index.ts

@@ -294,7 +294,7 @@ export { SandboxService } from './sandbox.service';
  * **RouterService** — Dynamic route URL construction utility
  *
  * Stateless utility for building dynamic route URLs from ROUTES base paths.
- * Static routes (START, NEW_CHAT, MCP_SERVERS) live in ROUTES constants;
+ * Static routes (START, MCP_SERVERS) live in ROUTES constants;
  * dynamic routes (CHAT, SETTINGS) are constructed here by appending parameters.
  *
  * **Architecture & Relationships:**

tools/ui/src/lib/stores/conversations/index.svelte.ts

@@ -21,6 +21,7 @@ import {
 	type ConversationsPreferencesHost
 } from '$lib/stores/conversations/preferences.svelte';
 import { settingsStore } from '$lib/stores/settings/index.svelte';
+import { tabsStore } from '$lib/stores/tabs.svelte';
 import { filterByLeafNodeId, findLeafNode, generateConversationTitle } from '$lib/utils';
 import { SvelteSet } from 'svelte/reactivity';
 import { toast } from 'svelte-sonner';
@@ -105,8 +106,8 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 	/**
 	 * Deletes multiple conversations in sequence.
-	 * Mirrors deleteConversation() per-id; navigates to NEW_CHAT only if the
-	 * currently-open chat was among the deleted ones.
+	 * Mirrors deleteConversation() per-id; navigates to the new-chat screen only
+	 * if the currently-open chat was among the deleted ones.
 	 * @param convIds - Conversation IDs to delete
 	 */
 	async bulkDeleteConversations(convIds: string[]): Promise<void> {
@@ -138,8 +139,13 @@ class ConversationsStore implements ConversationsPreferencesHost {
 			this.notifyConversationsDeleted([...idsToRemove]);
 
 			if (activeWasDeleted) {
+				const activeId = this.activeConversation!.id;
+
+				tabsStore.removeTabs([...idsToRemove].filter((id) => id !== activeId));
 				this.clearActiveConversation();
-				await goto(ROUTES.NEW_CHAT);
+				await tabsStore.close(activeId, activeId);
+			} else {
+				tabsStore.removeTabs([...idsToRemove]);
 			}
 
 			toast.success(
@@ -276,11 +282,12 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 			this.clearActiveConversation();
 			this.conversations = [];
+			tabsStore.clear();
 			this.notifyConversationsDeleted(allIds);
 
 			toast.success('All conversations deleted');
 
-			await goto(ROUTES.NEW_CHAT);
+			await goto(ROUTES.START);
 		} catch (error) {
 			console.error('Failed to delete all conversations:', error);
 			toast.error('Failed to delete conversations');
@@ -313,8 +320,13 @@ class ConversationsStore implements ConversationsPreferencesHost {
 				this.conversations = this.conversations.filter((c) => !idsToRemove.has(c.id));
 
 				if (this.activeConversation && idsToRemove.has(this.activeConversation.id)) {
+					const activeId = this.activeConversation.id;
+
+					tabsStore.removeTabs([...idsToRemove].filter((id) => id !== activeId));
 					this.clearActiveConversation();
-					await goto(ROUTES.NEW_CHAT);
+					await tabsStore.close(activeId, activeId);
+				} else {
+					tabsStore.removeTabs([...idsToRemove]);
 				}
 
 				this.notifyConversationsDeleted([...idsToRemove]);
@@ -333,7 +345,9 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 				if (this.activeConversation?.id === convId) {
 					this.clearActiveConversation();
-					await goto(ROUTES.NEW_CHAT);
+					await tabsStore.close(convId, convId);
+				} else {
+					tabsStore.removeTabs([convId]);
 				}
 
 				this.notifyConversationsDeleted([convId]);
@@ -571,6 +585,15 @@ class ConversationsStore implements ConversationsPreferencesHost {
 		return () => this.conversationDeletionListeners.delete(listener);
 	}
 
+	/**
+	 * Start a fresh chat by navigating to the bare `#/` new-chat screen. The
+	 * chat layout opens a new-chat tab for it when Conversation tabs are on.
+	 */
+	async openNewChat(): Promise<void> {
+		this.clearActiveConversation();
+		await goto(ROUTES.START);
+	}
+
 	/**
 	 * Refreshes active messages based on currNode after branch navigation.
 	 */
@@ -676,14 +699,6 @@ class ConversationsStore implements ConversationsPreferencesHost {
 		);
 	}
 
-	/**
-	 *
-	 *
-	 * Import & Export
-	 *
-	 *
-	 */
-
 	/**
 	 * Updates the current node of the active conversation
 	 * @param nodeId - The new current node ID
@@ -717,6 +732,14 @@ class ConversationsStore implements ConversationsPreferencesHost {
 		}
 	}
 
+	/**
+	 *
+	 *
+	 * Import & Export
+	 *
+	 *
+	 */
+
 	private notifyConversationsDeleted(convIds: string[]): void {
 		if (convIds.length === 0) return;
 

tools/ui/src/lib/stores/conversations/preferences.svelte.ts

@@ -152,11 +152,13 @@ export class ConversationPreferences {
 			return;
 		}
 
-		this.host.applyConversationUpdate(this.host.activeConversation.id, {
+		const id = this.host.activeConversation.id;
+
+		this.host.applyConversationUpdate(id, {
 			cwd: trimmed
 		});
 
-		await DatabaseService.updateConversation(this.host.activeConversation.id, {
+		await DatabaseService.updateConversation(id, {
 			cwd: trimmed
 		});
 
@@ -202,12 +204,15 @@ export class ConversationPreferences {
 			}
 		}
 
-		await DatabaseService.updateConversation(this.host.activeConversation.id, {
-			mcpServerOverrides: newOverrides.length > 0 ? newOverrides : undefined
+		const overrides = newOverrides.length > 0 ? newOverrides : undefined;
+		const id = this.host.activeConversation.id;
+
+		this.host.applyConversationUpdate(id, {
+			mcpServerOverrides: overrides
 		});
 
-		this.host.applyConversationUpdate(this.host.activeConversation.id, {
-			mcpServerOverrides: newOverrides.length > 0 ? newOverrides : undefined
+		await DatabaseService.updateConversation(id, {
+			mcpServerOverrides: overrides
 		});
 	}
 
@@ -224,11 +229,13 @@ export class ConversationPreferences {
 			return;
 		}
 
-		this.host.applyConversationUpdate(this.host.activeConversation.id, {
+		const id = this.host.activeConversation.id;
+
+		this.host.applyConversationUpdate(id, {
 			reasoningEffort: effort
 		});
 
-		await DatabaseService.updateConversation(this.host.activeConversation.id, {
+		await DatabaseService.updateConversation(id, {
 			reasoningEffort: effort
 		});
 	}

tools/ui/src/lib/stores/index.ts

@@ -22,6 +22,9 @@ export { chatStore } from './chat/index.svelte';
 
 export { draftMessagesStore } from './chat/drafts.svelte';
 
+// CONVERSATION TABS
+export { tabsStore } from './tabs.svelte';
+
 // CONTEXT STATS (active conversation context window usage)
 export { contextStatsStore } from './chat/context-stats.svelte';
 
@@ -40,6 +43,9 @@ export { modelsStore } from './models/index.svelte';
 // SERVER
 export { serverStore } from './server.svelte';
 
+// UI / LAYOUT
+export { uiStore } from './ui.svelte';
+
 // SETTINGS / UI PREFERENCES
 export { settingsStore } from './settings/index.svelte';
 

tools/ui/src/lib/stores/init.ts

@@ -1,21 +1,8 @@
-/**
- * Explicit store initialization, run once and shared by every caller.
- *
- * Order matters: migrations run first because they rename and rewrite
- * localStorage keys, so every store that reads localStorage initializes
- * only after they complete. Constructors and module-level side effects
- * stay empty so import order can no longer change startup behavior.
- *
- * The returned promise resolves once the persisted state is in memory, which
- * route loads await before reading settings: they run ahead of the root layout
- * script. The conversation list loads in the background, awaited by the chat
- * page that renders it.
- */
-
 // direct imports, not via the barrel, to avoid circular deps
 import { conversationsStore } from './conversations/index.svelte';
 import { permissionsStore } from './permissions.svelte';
 import { settingsStore } from './settings/index.svelte';
+import { tabsStore } from './tabs.svelte';
 import { toolsStore } from './tools.svelte';
 import { versionStore } from './version.svelte';
 import { browser } from '$app/environment';
@@ -33,7 +20,12 @@ export function initStores(): Promise<void> {
 		permissionsStore.initialize();
 		toolsStore.initialize();
 		void versionStore.initialize();
-		void conversationsStore.initialize();
+
+		// the full conversation list loads in the background; once it is back,
+		// prune persisted tabs against the conversations that still exist
+		void conversationsStore.initialize().then(() => {
+			tabsStore.init(conversationsStore.conversations.map((c) => c.id));
+		});
 	})();
 
 	return startup;

tools/ui/src/lib/stores/tabs.svelte.ts

@@ -0,0 +1,154 @@
+/**
+ * tabsStore - Reactive State Store for Browser-Style Conversation Tabs
+ *
+ * Tracks which conversations and the new-chat screen are open as tabs in
+ * the chat layout, in order. Real conversation tabs are `#/chat/<id>`
+ * routes; the new-chat tab is the bare `#/` route, represented here by the
+ * `NEW_CHAT_TAB_ID` sentinel (see {@link NEW_CHAT_TAB_ID}).
+ *
+ * **Architecture & Relationships:**
+ * - **conversationsStore**: owns conversation data; calls `removeTabs()` /
+ *   `close()` when conversations are deleted. This store never imports it,
+ *   so there is no circular dependency - tab names are resolved by the
+ *   ChatTabs component from conversationsStore.
+ * - Tab order persists to localStorage and is pruned against the loaded
+ *   conversation list on init. The new-chat tab is kept across reloads.
+ */
+
+import { browser } from '$app/environment';
+import { goto } from '$app/navigation';
+import { CONVERSATION_TABS_LOCALSTORAGE_KEY, NEW_CHAT_TAB_ID, ROUTES } from '$lib/constants';
+import { RouterService } from '$lib/services/router.service';
+import { untrack } from 'svelte';
+
+class TabsStore {
+	/** Ordered tab ids: conversation ids and the `NEW_CHAT_TAB_ID` sentinel */
+	openTabs = $state<string[]>([]);
+
+	/** False until init() has read the persisted tabs; save() is a no-op before that */
+	private initialized = false;
+
+	/** Navigate to a tab (the new-chat sentinel maps to the bare `#/` route) */
+	async activate(id: string): Promise<void> {
+		await goto(id === NEW_CHAT_TAB_ID ? ROUTES.START : RouterService.chat(id));
+	}
+
+	/** Remove all tabs (e.g. after deleting all conversations) */
+	clear(): void {
+		this.openTabs = [];
+		this.save();
+	}
+
+	/**
+	 * Close a tab. When it belongs to the active route, navigate to the left
+	 * neighbor (or the right one when the closed tab was leftmost), falling
+	 * back to the new-chat screen when no tabs remain.
+	 * @param id - Tab id to close
+	 * @param activeTabId - Tab id of the current route, if any
+	 */
+	async close(id: string, activeTabId: string | null): Promise<void> {
+		const idx = this.openTabs.indexOf(id);
+
+		if (idx === -1) {
+			// tab not tracked (e.g. Conversation tabs are off); still fall back to
+			// the new-chat screen when closing the active conversation
+			if (id === activeTabId) {
+				await goto(ROUTES.START);
+			}
+
+			return;
+		}
+
+		this.openTabs = this.openTabs.filter((tabId) => tabId !== id);
+		this.save();
+
+		if (id !== activeTabId) return;
+
+		const target = (idx > 0 ? this.openTabs[idx - 1] : this.openTabs[0]) ?? null;
+
+		if (target) {
+			await goto(target === NEW_CHAT_TAB_ID ? ROUTES.START : RouterService.chat(target));
+		} else {
+			await goto(ROUTES.START);
+		}
+	}
+
+	/**
+	 * Load persisted tabs, dropping conversation ids that no longer exist.
+	 * Called once from initStores() after conversations are loaded.
+	 * Merges with (rather than replaces) current openTabs: the chat layout
+	 * syncs the route's tab before this async init completes, and replacing
+	 * here would drop it.
+	 * @param validIds - Ids of conversations present in the database
+	 */
+	init(validIds: string[]): void {
+		if (!browser) return;
+
+		// the new-chat sentinel is a pseudo-tab, not a conversation, but it is
+		// still kept so a reload on `#/` does not drop the tab the user is on
+		const isLive = (id: string) => validIds.includes(id) || id === NEW_CHAT_TAB_ID;
+		const persisted = this.load().filter(isLive);
+		// tabs already in openTabs come from the live route, so they stay as they
+		// are: `validIds` is a snapshot and a conversation created while the list
+		// was loading is not in it
+		const extras = this.openTabs.filter((id) => !persisted.includes(id));
+
+		this.openTabs = [...persisted, ...extras];
+		this.initialized = true;
+		this.save();
+	}
+
+	/**
+	 * Remove tabs without navigating. Used when conversations are deleted
+	 * while some other conversation stays open.
+	 * @param ids - Tab ids to drop
+	 */
+	removeTabs(ids: string[]): void {
+		const removed = new Set(ids);
+		const next = this.openTabs.filter((id) => !removed.has(id));
+
+		if (next.length !== this.openTabs.length) {
+			this.openTabs = next;
+			this.save();
+		}
+	}
+
+	/**
+	 * Sync the tab strip with the route. Called from the chat layout on every
+	 * navigation, so any way of reaching a conversation or new-chat tab opens
+	 * a tab for it.
+	 * @param id - The conversation (or temporary new-chat) id of the route
+	 */
+	syncWithRoute(id: string): void {
+		// untrack: callers invoke this from an effect keyed on the route, and
+		// reading openTabs here would subscribe that effect to openTabs too -
+		// closing the active tab would then re-run the effect and re-add the tab
+		untrack(() => {
+			if (!this.openTabs.includes(id)) {
+				this.openTabs = [...this.openTabs, id];
+				this.save();
+			}
+		});
+	}
+
+	private load(): string[] {
+		try {
+			const raw = localStorage.getItem(CONVERSATION_TABS_LOCALSTORAGE_KEY);
+			const parsed: unknown = raw ? JSON.parse(raw) : [];
+
+			return Array.isArray(parsed) ? parsed.filter((id) => typeof id === 'string') : [];
+		} catch {
+			return [];
+		}
+	}
+
+	private save(): void {
+		// never write before init has read the persisted tabs, or an early
+		// route sync (layout effect runs before async init) would clobber them
+		if (!browser || !this.initialized) return;
+
+		localStorage.setItem(CONVERSATION_TABS_LOCALSTORAGE_KEY, JSON.stringify(this.openTabs));
+	}
+}
+
+export const tabsStore = new TabsStore();

tools/ui/src/lib/stores/ui.svelte.ts

@@ -0,0 +1,14 @@
+/**
+ * uiStore - Shared UI/layout state
+ *
+ * Holds cross-component UI state that does not belong to a single component
+ * (e.g. the desktop sidebar's expanded/collapsed state, which the sidebar
+ * controls and the chat tab bar reacts to).
+ */
+
+class UiStore {
+	/** Whether the desktop sidebar is expanded (open). */
+	isSidebarExpanded = $state(false);
+}
+
+export const uiStore = new UiStore();

tools/ui/src/lib/types/navigation.d.ts

@@ -1,3 +1,4 @@
+import type { SidebarAction } from '$lib/enums';
 import type { Component } from 'svelte';
 
 /**
@@ -7,6 +8,8 @@ export interface DesktopIconStripItem {
 	icon: Component;
 	tooltip: string;
 	route?: string;
+	/** Custom action handled by the sidebar, e.g. opening a new-chat tab */
+	action?: SidebarAction;
 	activeRouteId?: string;
 	activeRoutePrefix?: string;
 	activeUrlIncludes?: string;

tools/ui/src/routes/(chat)/+layout.svelte

@@ -1,12 +1,34 @@
 <script lang="ts">
 	import { page } from '$app/state';
-	import { ChatScreen } from '$lib/components/app';
+	import { ChatScreen, ChatTabs } from '$lib/components/app';
+	import { NEW_CHAT_TAB_ID } from '$lib/constants';
+	import { settingsStore, tabsStore } from '$lib/stores';
 
 	let { children } = $props();
 
+	// the new-chat screen is the bare `#/` route (no conversation id)
 	let showCenteredEmpty = $derived(!page.params.id);
+
+	let showTabs = $derived(
+		Boolean(settingsStore.config.conversationTabs) &&
+			(page.params.id || tabsStore.openTabs.some((id) => id !== NEW_CHAT_TAB_ID))
+	);
+
+	$effect(() => {
+		const id = page.params.id ?? (page.route.id === '/(chat)' ? NEW_CHAT_TAB_ID : undefined);
+
+		if (id && settingsStore.config.conversationTabs) {
+			tabsStore.syncWithRoute(id);
+		}
+	});
 </script>
 
-<ChatScreen {showCenteredEmpty} />
+<div class={showTabs ? 'md:[--chat-tabs-offset:1.25rem]' : ''}>
+	{#if showTabs}
+		<ChatTabs />
+	{/if}
+
+	<ChatScreen {showCenteredEmpty} />
+</div>
 
 {@render children?.()}

tools/ui/src/routes/(chat)/+page.svelte

@@ -8,24 +8,19 @@
 
 	let qParam = $derived(page.url.searchParams.get(URL_PARAMS.QUERY));
 	let modelParam = $derived(page.url.searchParams.get(URL_PARAMS.MODEL));
-	let newChatParam = $derived(page.url.searchParams.get(URL_PARAMS.NEW_CHAT));
 	let loadParam = $derived(page.url.searchParams.get(URL_PARAMS.LOAD));
 
-	// Dialog state for model not available error
 	let showModelNotAvailable = $state(false);
 	let requestedModelName = $state('');
 	let availableModelNames = $derived(modelsStore.models.map((m) => m.model));
 
-	/**
-	 * Clear URL params after message is sent to prevent re-sending on refresh
-	 */
+	// Clear params after handling the deep link so a refresh does not replay them
 	function clearUrlParams() {
 		const url = new URL(page.url);
 
 		url.searchParams.delete(URL_PARAMS.QUERY);
 		url.searchParams.delete(URL_PARAMS.MODEL);
 		url.searchParams.delete(URL_PARAMS.LOAD);
-		url.searchParams.delete(URL_PARAMS.NEW_CHAT);
 
 		replaceState(url.toString(), {});
 	}
@@ -40,8 +35,8 @@
 				try {
 					await modelsStore.selectModelById(model.id);
 
-					// with ?load=true, start loading right away so the model is ready sooner;
-					// not awaited, so the UI stays usable during the load
+					// with ?load=true in router mode, start loading right away so the
+					// model is ready sooner; not awaited so the UI stays usable
 					if (
 						loadParam === 'true' &&
 						serverStore.isRouterMode &&
@@ -66,11 +61,12 @@
 			}
 		}
 
-		// Handle ?q= parameter - create new conversation and send message
+		// ?q= creates the conversation, the chat route sends the prompt once the
+		// conversation id is in the URL
 		if (qParam !== null) {
 			await conversationsStore.createConversation();
 			clearUrlParams();
-		} else if (modelParam || newChatParam === 'true') {
+		} else if (modelParam) {
 			clearUrlParams();
 		}
 	}
@@ -85,7 +81,7 @@
 
 		await modelsStore.fetch();
 
-		if (qParam !== null || modelParam !== null || newChatParam === 'true') {
+		if (qParam !== null || modelParam !== null) {
 			await handleUrlParams();
 		}
 

tools/ui/src/routes/+layout.svelte

@@ -11,6 +11,7 @@
 		FAVICON_PATHS,
 		FAVICON_SELECTORS,
 		HEADERS,
+		NEW_CHAT_TAB_ID,
 		ROUTES,
 		SETTINGS_KEYS,
 		TOOLTIP_DELAY_DURATION
@@ -26,6 +27,7 @@
 		modelsStore,
 		serverStore,
 		settingsStore,
+		tabsStore,
 		versionStore
 	} from '$lib/stores';
 	import { initStores } from '$lib/stores/init';
@@ -74,6 +76,27 @@
 		}
 	}
 
+	function navigateToTab(direction: -1 | 1) {
+		// only makes sense with conversation tabs enabled
+		if (!settingsStore.config.conversationTabs) return;
+
+		const openTabs = tabsStore.openTabs;
+
+		if (openTabs.length === 0) return;
+
+		const activeId = page.params.id ?? NEW_CHAT_TAB_ID;
+		const idx = openTabs.indexOf(activeId);
+		// active tab not in list (e.g. a non-chat route): start from an edge
+		const targetIdx =
+			idx === -1
+				? direction === 1
+					? 0
+					: openTabs.length - 1
+				: (idx + direction + openTabs.length) % openTabs.length;
+
+		void tabsStore.activate(openTabs[targetIdx]);
+	}
+
 	function navigateToConversation(direction: -1 | 1) {
 		const allConvs = conversationsStore.conversations;
 
@@ -96,15 +119,31 @@
 		if (targetIdx >= 0 && targetIdx < allConvs.length) {
 			goto(RouterService.chat(allConvs[targetIdx].id));
 		} else {
-			goto(ROUTES.NEW_CHAT);
+			conversationsStore.openNewChat();
 		}
 	}
 
+	// navigating away from the new-chat screen drops its tab, so it does not
+	// linger once the user moves to a real conversation or another route
+	let previousChatId = $state<string | undefined>(undefined);
+
+	$effect(() => {
+		const id = page.params.id ?? (page.route.id === '/(chat)' ? NEW_CHAT_TAB_ID : undefined);
+		const prev = untrack(() => previousChatId);
+
+		previousChatId = id;
+
+		if (id !== prev && prev && settingsStore.config.conversationTabs && prev === NEW_CHAT_TAB_ID) {
+			untrack(() => tabsStore.removeTabs([NEW_CHAT_TAB_ID]));
+		}
+	});
 	// Global keyboard shortcuts
 	const { handleKeydown } = useKeyboardShortcuts({
 		editActiveConversation: () => chatSidebar?.editActiveConversation?.(),
 		navigateToNextConversation: () => navigateToConversation(1),
-		navigateToPrevConversation: () => navigateToConversation(-1)
+		navigateToNextTab: () => navigateToTab(1),
+		navigateToPrevConversation: () => navigateToConversation(-1),
+		navigateToPrevTab: () => navigateToTab(-1)
 	});
 
 	function checkApiKey() {

tools/ui/src/routes/search/+page.svelte

@@ -21,10 +21,10 @@
 	});
 
 	// Search page is intended for mobile; on desktop the sidebar already exposes
-	// in-place search, so bounce back to a chat.
+	// in-place search, so bounce back to a new-chat tab without a history entry.
 	$effect(() => {
 		if (browser && !deviceStore.isMobile) {
-			goto(ROUTES.NEW_CHAT, { replaceState: true });
+			goto(ROUTES.START, { replaceState: true });
 		}
 	});
 
@@ -66,7 +66,7 @@
 		if (history.length > 1) {
 			history.back();
 		} else {
-			goto(ROUTES.NEW_CHAT);
+			conversationsStore.openNewChat();
 		}
 	}
 </script>

tools/ui/tests/stories/a11y/ScrollCarousel.a11y.stories.svelte

@@ -1,15 +1,16 @@
 <script module lang="ts">
 	import { defineMeta } from '@storybook/addon-svelte-csf';
-	import HorizontalScrollCarousel from '$lib/components/app/misc/HorizontalScrollCarousel.svelte';
+	import { ScrollCarousel } from '$lib/components/app';
+	import { ScrollCarouselVariant } from '$lib/enums';
 	import { expect, waitFor } from 'storybook/test';
 
 	const { Story } = defineMeta({
-		component: HorizontalScrollCarousel,
+		component: ScrollCarousel,
 		parameters: {
 			layout: 'centered'
 		},
 		tags: ['!dev'],
-		title: 'Components/HorizontalScrollCarousel/Accessibility'
+		title: 'Components/ScrollCarousel/Accessibility'
 	});
 </script>
 
@@ -33,10 +34,10 @@
 >
 	<div>
 		<button type="button">before</button>
-		<HorizontalScrollCarousel class="w-96">
+		<ScrollCarousel class="w-96" variant={ScrollCarouselVariant.CENTER}>
 			<div class="h-12 w-12 shrink-0 bg-muted"></div>
 			<div class="h-12 w-12 shrink-0 bg-muted"></div>
-		</HorizontalScrollCarousel>
+		</ScrollCarousel>
 		<button type="button">after</button>
 	</div>
 </Story>
@@ -60,10 +61,10 @@
 >
 	<div>
 		<button type="button">before</button>
-		<HorizontalScrollCarousel class="w-48">
+		<ScrollCarousel class="w-48" variant={ScrollCarouselVariant.CENTER}>
 			{#each [...Array(20).keys()] as i (i)}
 				<div class="h-12 w-24 shrink-0 bg-muted">{i}</div>
 			{/each}
-		</HorizontalScrollCarousel>
+		</ScrollCarousel>
 	</div>
 </Story>

vendor/sheredom/subprocess.h

@@ -275,13 +275,46 @@ subprocess_weak int subprocess_alive(struct subprocess_s *const process);
 #include <unistd.h>
 #endif
 
+#if defined(__NetBSD__)
+#include <sys/param.h>
+#endif
+
+/* Which spelling of the chdir file action the platform provides, if any.
+   POSIX 2024 standardised posix_spawn_file_actions_addchdir; implementations
+   that shipped it earlier called it ..._np. macOS 26 and NetBSD 10 use the
+   standard name, glibc 2.29+, macOS 10.15+ and FreeBSD 13.1+ use the _np name,
+   and AIX, NetBSD 9 and older, and OpenBSD provide neither. */
+#if !defined(SUBPROCESS_ADDCHDIR_IS_POSIX)
+#if (defined(__APPLE__) && MAC_OS_X_VERSION_MIN_REQUIRED >= 260000) ||        \
+    (defined(__NetBSD__) && __NetBSD_Version__ >= 1000000000)
+#define SUBPROCESS_ADDCHDIR_IS_POSIX 1
+#else
+#define SUBPROCESS_ADDCHDIR_IS_POSIX 0
+#endif
+#endif
+
+/* Whether to launch the child with fork()+exec() instead of posix_spawn(),
+   for platforms with no posix_spawn_file_actions_addchdir under either
+   spelling: the child chdir()s before exec, and a close-on-exec pipe carries
+   exec's errno back. Define this yourself to force either implementation. */
+#if !defined(SUBPROCESS_SPAWN_VIA_FORK)
+#if defined(_AIX) || defined(__OpenBSD__) ||                                  \
+    (defined(__NetBSD__) && (__NetBSD_Version__ < 1000000000))
+#define SUBPROCESS_SPAWN_VIA_FORK 1
+#else
+#define SUBPROCESS_SPAWN_VIA_FORK 0
+#endif
+#endif
+
 /* Whether subprocess_create_ex can honour process_cwd. glibc only gained
    posix_spawn_file_actions_addchdir_np in 2.29, and macOS in 10.15; the SDKs
    mark it unavailable on iOS, tvOS and watchOS, where the undefined version
    macro folds to 0 and so answers correctly. Define this yourself to override
    the detection, for instance on musl older than 1.1.24. */
 #if !defined(SUBPROCESS_HAVE_CWD)
-#if defined(__GLIBC__)
+#if SUBPROCESS_SPAWN_VIA_FORK
+#define SUBPROCESS_HAVE_CWD 1
+#elif defined(__GLIBC__)
 #if __GLIBC_PREREQ(2, 29)
 #define SUBPROCESS_HAVE_CWD 1
 #else
@@ -294,10 +327,13 @@ subprocess_weak int subprocess_alive(struct subprocess_s *const process);
 #endif
 #endif
 
-/* Whether posix_spawn reports a failed exec back to the caller. glibc only
-   started doing so in 2.24; before that the child silently exits with 127. */
+/* Whether a failed exec is reported back to the caller. The fork() path always
+   reports it through its error pipe. glibc's posix_spawn only started doing so
+   in 2.24; before that the child silently exits with 127. */
 #if !defined(SUBPROCESS_SPAWN_REPORTS_EXEC_ERRORS)
-#if defined(__GLIBC__)
+#if SUBPROCESS_SPAWN_VIA_FORK
+#define SUBPROCESS_SPAWN_REPORTS_EXEC_ERRORS 1
+#elif defined(__GLIBC__)
 #if __GLIBC_PREREQ(2, 24)
 #define SUBPROCESS_SPAWN_REPORTS_EXEC_ERRORS 1
 #else
@@ -342,6 +378,14 @@ typedef intptr_t subprocess_intptr_t;
 typedef size_t subprocess_size_t;
 #endif
 
+/* SIZE_T is ULONG_PTR, which is not size_t: on Win32 both are 32 bits wide but
+   unsigned long and unsigned int are still distinct types. */
+#ifdef _WIN64
+typedef subprocess_size_t subprocess_ulongptr_t;
+#else
+typedef unsigned long subprocess_ulongptr_t;
+#endif
+
 #ifdef __clang__
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Wreserved-identifier"
@@ -351,6 +395,7 @@ typedef struct _PROCESS_INFORMATION *LPPROCESS_INFORMATION;
 typedef struct _SECURITY_ATTRIBUTES *LPSECURITY_ATTRIBUTES;
 typedef struct _STARTUPINFOW *LPSTARTUPINFOW;
 typedef struct _OVERLAPPED *LPOVERLAPPED;
+typedef struct _PROC_THREAD_ATTRIBUTE_LIST *LPPROC_THREAD_ATTRIBUTE_LIST;
 
 #ifdef __clang__
 #pragma clang diagnostic pop
@@ -402,6 +447,11 @@ struct subprocess_startup_info_s {
   void *hStdError;
 };
 
+struct subprocess_startup_info_ex_s {
+  struct subprocess_startup_info_s startupInfo;
+  void *attributeList;
+};
+
 struct subprocess_overlapped_s {
   uintptr_t Internal;
   uintptr_t InternalHigh;
@@ -451,6 +501,14 @@ __declspec(dllimport) int __stdcall CreateProcessW(
     const subprocess_wchar_t *, subprocess_wchar_t *, LPSECURITY_ATTRIBUTES,
     LPSECURITY_ATTRIBUTES, int, unsigned long, void *,
     const subprocess_wchar_t *, LPSTARTUPINFOW, LPPROCESS_INFORMATION);
+__declspec(dllimport) int __stdcall
+InitializeProcThreadAttributeList(LPPROC_THREAD_ATTRIBUTE_LIST, unsigned long,
+                                  unsigned long, subprocess_ulongptr_t *);
+__declspec(dllimport) int __stdcall UpdateProcThreadAttribute(
+    LPPROC_THREAD_ATTRIBUTE_LIST, unsigned long, subprocess_ulongptr_t, void *,
+    subprocess_ulongptr_t, void *, subprocess_ulongptr_t *);
+__declspec(dllimport) void __stdcall
+DeleteProcThreadAttributeList(LPPROC_THREAD_ATTRIBUTE_LIST);
 __declspec(dllimport) int __stdcall MultiByteToWideChar(
     unsigned int, unsigned long, const char *, int, subprocess_wchar_t *, int);
 __declspec(dllimport) int __stdcall CloseHandle(void *);
@@ -667,12 +725,104 @@ int subprocess_create_named_pipe_helper(void **rd, void **wr) {
 }
 #endif
 
+#if !defined(_WIN32)
+/* Move a pipe end off 0, 1 or 2. Duplicating a descriptor onto itself is a
+   no-op, so a pipe end already sitting on a standard descriptor would keep its
+   FD_CLOEXEC and be closed by exec, leaving the child without that stream. */
+static int subprocess_fds_above_std(int fds[2]) {
+  int fd_flags;
+  int index;
+  int moved;
+  int saved_errno;
+
+  for (index = 0; index < 2; index++) {
+    if (fds[index] > STDERR_FILENO) {
+      continue;
+    }
+
+    moved = fcntl(fds[index], F_DUPFD, STDERR_FILENO + 1);
+    if (-1 != moved) {
+      fd_flags = fcntl(moved, F_GETFD, 0);
+      if ((-1 == fd_flags) ||
+          (-1 == fcntl(moved, F_SETFD, fd_flags | FD_CLOEXEC))) {
+        saved_errno = errno;
+        close(moved);
+        errno = saved_errno;
+        moved = -1;
+      }
+    }
+
+    if (-1 == moved) {
+      saved_errno = errno;
+      close(fds[0]);
+      close(fds[1]);
+      fds[0] = -1;
+      fds[1] = -1;
+      errno = saved_errno;
+      return -1;
+    }
+
+    close(fds[index]);
+    fds[index] = moved;
+  }
+
+  return 0;
+}
+
+/* Create pipes with close-on-exec set so later subprocesses do not inherit
+   descriptors belonging to subprocesses which are already running. */
+static int subprocess_pipe_cloexec(int fds[2]) {
+  int fd_flags;
+  int index;
+  int saved_errno;
+
+#if defined(__linux__) || defined(__FreeBSD__) || defined(__NetBSD__) ||       \
+    defined(__OpenBSD__) || defined(__DragonFly__) ||                         \
+    (defined(__sun) && defined(__SVR4))
+  if (0 == pipe2(fds, O_CLOEXEC)) {
+    return subprocess_fds_above_std(fds);
+  }
+
+  /* Older kernels can lack pipe2 even when the C library declares it. */
+  if (ENOSYS != errno) {
+    return -1;
+  }
+#endif
+
+  if (0 != pipe(fds)) {
+    return -1;
+  }
+
+  for (index = 0; index < 2; index++) {
+    fd_flags = fcntl(fds[index], F_GETFD, 0);
+    if ((-1 == fd_flags) ||
+        (-1 == fcntl(fds[index], F_SETFD, fd_flags | FD_CLOEXEC))) {
+      saved_errno = errno;
+      close(fds[0]);
+      close(fds[1]);
+      fds[0] = -1;
+      fds[1] = -1;
+      errno = saved_errno;
+      return -1;
+    }
+  }
+
+  return subprocess_fds_above_std(fds);
+}
+#endif
+
 int subprocess_create(const char *const commandLine[], int options,
                       struct subprocess_s *const out_process) {
   return subprocess_create_ex(commandLine, options, SUBPROCESS_NULL,
                               SUBPROCESS_NULL, out_process);
 }
 
+#if SUBPROCESS_SPAWN_VIA_FORK
+/* Not every platform declares execvpe: AIX exports it from libc without ever
+   naming it in a header, and glibc hides it behind _GNU_SOURCE. */
+extern int execvpe(const char *, char *const *, char *const *);
+#endif
+
 int subprocess_create_ex(const char *const commandLine[], int options,
                          const char *const environment[],
                          const char *const process_cwd,
@@ -692,6 +842,7 @@ int subprocess_create_ex(const char *const commandLine[], int options,
   subprocess_size_t bs_run;
   unsigned long flags = 0;
   unsigned long last_error = 0;
+  int attribute_list_initialized = 0;
   int result = subprocess_error_unknown;
   const unsigned int codePageUtf8 = 65001;
   const unsigned long mbErrInvalidChars = 0x00000008;
@@ -699,6 +850,8 @@ int subprocess_create_ex(const char *const commandLine[], int options,
   const unsigned long handleFlagInherit = 0x00000001;
   const unsigned long createNoWindow = 0x08000000;
   const unsigned long createUnicodeEnvironment = 0x00000400;
+  const unsigned long extendedStartupInfoPresent = 0x00080000;
+  const subprocess_size_t procThreadAttributeHandleList = 0x00020002;
   struct subprocess_subprocess_information_s processInfo = {SUBPROCESS_NULL,
                                                             SUBPROCESS_NULL, 0,
                                                             0};
@@ -706,6 +859,11 @@ int subprocess_create_ex(const char *const commandLine[], int options,
                                                     SUBPROCESS_NULL, 1};
   subprocess_wchar_t empty_environment[2] = {0, 0};
   subprocess_wchar_t *used_environment = SUBPROCESS_NULL;
+  subprocess_ulongptr_t attribute_list_size = 0;
+  subprocess_size_t inherited_handle_count = 0;
+  LPPROC_THREAD_ATTRIBUTE_LIST attribute_list = SUBPROCESS_NULL;
+  void *inherited_handles[3];
+  struct subprocess_startup_info_ex_s startInfoEx;
   struct subprocess_startup_info_s startInfo = {0,
                                                 SUBPROCESS_NULL,
                                                 SUBPROCESS_NULL,
@@ -1080,6 +1238,44 @@ int subprocess_create_ex(const char *const commandLine[], int options,
     }
   }
 
+  /* Restrict inheritance to this subprocess's standard streams. Without a
+     handle list, concurrent subprocess_create calls can inherit each other's
+     temporarily-inheritable child pipe handles. */
+  inherited_handles[inherited_handle_count++] = startInfo.hStdInput;
+  inherited_handles[inherited_handle_count++] = startInfo.hStdOutput;
+  if (startInfo.hStdError != startInfo.hStdOutput) {
+    inherited_handles[inherited_handle_count++] = startInfo.hStdError;
+  }
+
+  InitializeProcThreadAttributeList(SUBPROCESS_NULL, 1, 0,
+                                    &attribute_list_size);
+  if (0 == attribute_list_size) {
+    result = subprocess_error_spawn;
+    goto cleanup;
+  }
+
+  attribute_list = SUBPROCESS_PTR_CAST(LPPROC_THREAD_ATTRIBUTE_LIST,
+                                       _alloca(attribute_list_size));
+  if (!attribute_list || !InitializeProcThreadAttributeList(
+                             attribute_list, 1, 0, &attribute_list_size)) {
+    result = subprocess_error_spawn;
+    goto cleanup;
+  }
+  attribute_list_initialized = 1;
+
+  if (!UpdateProcThreadAttribute(
+          attribute_list, 0, procThreadAttributeHandleList, inherited_handles,
+          inherited_handle_count * sizeof(inherited_handles[0]),
+          SUBPROCESS_NULL, SUBPROCESS_NULL)) {
+    result = subprocess_error_spawn;
+    goto cleanup;
+  }
+
+  startInfoEx.startupInfo = startInfo;
+  startInfoEx.startupInfo.cb = sizeof(startInfoEx);
+  startInfoEx.attributeList = attribute_list;
+  flags |= extendedStartupInfoPresent;
+
   if (!CreateProcessW(
           SUBPROCESS_NULL,
           commandLineCombinedWide, // command line
@@ -1090,7 +1286,7 @@ int subprocess_create_ex(const char *const commandLine[], int options,
           used_environment,        // used environment
           process_cwd_wide,        // use specified current directory
           SUBPROCESS_PTR_CAST(LPSTARTUPINFOW,
-                              &startInfo), // STARTUPINFO pointer
+                              &startInfoEx), // STARTUPINFOEX pointer
           SUBPROCESS_PTR_CAST(LPPROCESS_INFORMATION, &processInfo))) {
     result = subprocess_error_from_windows_error(GetLastError());
     if (subprocess_error_unknown == result) {
@@ -1099,6 +1295,9 @@ int subprocess_create_ex(const char *const commandLine[], int options,
     goto cleanup;
   }
 
+  DeleteProcThreadAttributeList(attribute_list);
+  attribute_list_initialized = 0;
+
   out_process->hProcess = processInfo.hProcess;
   processInfo.hProcess = SUBPROCESS_NULL;
 
@@ -1128,6 +1327,10 @@ int subprocess_create_ex(const char *const commandLine[], int options,
 cleanup:
   last_error = GetLastError();
 
+  if (attribute_list_initialized) {
+    DeleteProcThreadAttributeList(attribute_list);
+  }
+
   if (subprocess_error_unknown == result) {
     result = subprocess_error_from_windows_error(last_error);
   }
@@ -1173,15 +1376,20 @@ cleanup:
   int stderrfd[2] = {-1, -1};
   int fd, fd_flags;
   int async_no_wait;
-  int actions_created = 0;
   int result = subprocess_error_unknown;
   int saved_errno = 0;
-  int posix_error;
   pid_t child = 0;
   extern char **environ;
   char *const empty_environment[1] = {SUBPROCESS_NULL};
-  posix_spawn_file_actions_t actions;
   char *const *used_environment;
+#if SUBPROCESS_SPAWN_VIA_FORK
+  /* Pipe used to relay the child's exec() errno back to the parent. */
+  int exec_errfd[2] = {-1, -1};
+#else
+  int actions_created = 0;
+  int posix_error;
+  posix_spawn_file_actions_t actions;
+#endif
 
   async_no_wait = subprocess_option_enable_async_no_wait ==
                   (options & subprocess_option_enable_async_no_wait);
@@ -1202,13 +1410,13 @@ cleanup:
 
   memset(out_process, 0, sizeof(*out_process));
 
-  if (0 != pipe(stdinfd)) {
+  if (0 != subprocess_pipe_cloexec(stdinfd)) {
     saved_errno = errno;
     result = subprocess_error_pipe;
     goto cleanup;
   }
 
-  if (0 != pipe(stdoutfd)) {
+  if (0 != subprocess_pipe_cloexec(stdoutfd)) {
     saved_errno = errno;
     result = subprocess_error_pipe;
     goto cleanup;
@@ -1216,7 +1424,7 @@ cleanup:
 
   if (subprocess_option_combined_stdout_stderr !=
       (options & subprocess_option_combined_stdout_stderr)) {
-    if (0 != pipe(stderrfd)) {
+    if (0 != subprocess_pipe_cloexec(stderrfd)) {
       saved_errno = errno;
       result = subprocess_error_pipe;
       goto cleanup;
@@ -1240,6 +1448,136 @@ cleanup:
     used_environment = empty_environment;
   }
 
+#if SUBPROCESS_SPAWN_VIA_FORK
+  /* fork()+exec() instead of posix_spawn, so the child can chdir() first.
+     exec_errfd[1] is close-on-exec: a successful exec closes it and the parent
+     reads EOF; a failed exec writes errno through it before _exit. */
+  if (0 != pipe(exec_errfd)) {
+    saved_errno = errno;
+    result = subprocess_error_pipe;
+    goto cleanup;
+  }
+
+  if (-1 == fcntl(exec_errfd[1], F_SETFD, FD_CLOEXEC)) {
+    saved_errno = errno;
+    result = subprocess_error_spawn;
+    goto cleanup;
+  }
+
+  child = fork();
+
+  if (child < 0) {
+    saved_errno = errno;
+    result = subprocess_error_spawn;
+    goto cleanup;
+  }
+
+  if (0 == child) {
+    /* Child. Everything below must stay async-signal-safe: after fork() in a
+       threaded process only such functions may be called before exec. */
+    int child_errno;
+
+    close(exec_errfd[0]);
+
+    if ((-1 == dup2(stdinfd[0], STDIN_FILENO)) ||
+        (-1 == dup2(stdoutfd[1], STDOUT_FILENO))) {
+      goto child_failed;
+    }
+
+    if (subprocess_option_combined_stdout_stderr ==
+        (options & subprocess_option_combined_stdout_stderr)) {
+      if (-1 == dup2(STDOUT_FILENO, STDERR_FILENO)) {
+        goto child_failed;
+      }
+    } else {
+      if (-1 == dup2(stderrfd[1], STDERR_FILENO)) {
+        goto child_failed;
+      }
+    }
+
+    /* The originals are only closed once they have been duplicated, so that a
+       pipe end that already sits on 0, 1 or 2 is not closed out from under us. */
+    if (stdinfd[0] > STDERR_FILENO) {
+      close(stdinfd[0]);
+    }
+    if (stdinfd[1] > STDERR_FILENO) {
+      close(stdinfd[1]);
+    }
+    if (stdoutfd[0] > STDERR_FILENO) {
+      close(stdoutfd[0]);
+    }
+    if (stdoutfd[1] > STDERR_FILENO) {
+      close(stdoutfd[1]);
+    }
+    if (stderrfd[0] > STDERR_FILENO) {
+      close(stderrfd[0]);
+    }
+    if (stderrfd[1] > STDERR_FILENO) {
+      close(stderrfd[1]);
+    }
+
+    if (process_cwd && (0 != chdir(process_cwd))) {
+      goto child_failed;
+    }
+
+#ifdef __clang__
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wcast-qual"
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#endif
+    if (subprocess_option_search_user_path ==
+        (options & subprocess_option_search_user_path)) {
+      execvpe(commandLine[0],
+              SUBPROCESS_CONST_CAST(char *const *, commandLine),
+              SUBPROCESS_CONST_CAST(char *const *, used_environment));
+    } else {
+      execve(commandLine[0],
+             SUBPROCESS_CONST_CAST(char *const *, commandLine),
+             SUBPROCESS_CONST_CAST(char *const *, used_environment));
+    }
+#ifdef __clang__
+#pragma clang diagnostic pop
+#endif
+
+  child_failed:
+    child_errno = errno;
+    /* Nothing useful can be done if this write fails; the parent then sees EOF
+       and reports success, exactly as posix_spawn would without exec reporting. */
+    (void)!write(exec_errfd[1], &child_errno, sizeof(child_errno));
+    /* 127 is what POSIX requires posix_spawn's child to exit with when exec
+       fails, so both implementations look the same to a caller. */
+    _exit(127);
+  }
+
+  /* Parent. */
+  close(exec_errfd[1]);
+  exec_errfd[1] = -1;
+
+  {
+    int child_errno = 0;
+    ssize_t bytes_read;
+
+    do {
+      bytes_read = read(exec_errfd[0], &child_errno, sizeof(child_errno));
+    } while ((-1 == bytes_read) && (EINTR == errno));
+
+    close(exec_errfd[0]);
+    exec_errfd[0] = -1;
+
+    if (bytes_read == (ssize_t)sizeof(child_errno)) {
+      /* exec failed in the child. Reap it and surface the reason. */
+      while ((-1 == waitpid(child, SUBPROCESS_NULL, 0)) && (EINTR == errno)) {
+      }
+      child = 0;
+      saved_errno = child_errno;
+      result = subprocess_error_from_errno(child_errno);
+      if (subprocess_error_unknown == result) {
+        result = subprocess_error_spawn;
+      }
+      goto cleanup;
+    }
+  }
+#else
   posix_error = posix_spawn_file_actions_init(&actions);
   if (0 != posix_error) {
     saved_errno = posix_error;
@@ -1253,7 +1591,7 @@ cleanup:
 
   // Set working directory
   if (process_cwd) {
-#if defined(__NetBSD__) || (defined(__APPLE__) && MAC_OS_X_VERSION_MIN_REQUIRED >= 260000)
+#if SUBPROCESS_ADDCHDIR_IS_POSIX
     posix_error = posix_spawn_file_actions_addchdir(&actions, process_cwd);
 #elif !SUBPROCESS_HAVE_CWD
     posix_error = ENOSYS;
@@ -1406,6 +1744,7 @@ cleanup:
 #ifdef __clang__
 #pragma clang diagnostic pop
 #endif
+#endif /* SUBPROCESS_SPAWN_VIA_FORK */
 
   // Close the stdin read end
   close(stdinfd[0]);
@@ -1480,9 +1819,21 @@ cleanup:
     result = subprocess_error_from_errno(saved_errno);
   }
 
+#if SUBPROCESS_SPAWN_VIA_FORK
+  if (-1 != exec_errfd[0]) {
+    close(exec_errfd[0]);
+    exec_errfd[0] = -1;
+  }
+
+  if (-1 != exec_errfd[1]) {
+    close(exec_errfd[1]);
+    exec_errfd[1] = -1;
+  }
+#else
   if (actions_created) {
     posix_spawn_file_actions_destroy(&actions);
   }
+#endif
 
   if (0 != result) {
     if (child) {