test: move tools/parser to tests

CONTRIBUTING.md

@@ -74,7 +74,6 @@ 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)
 
-[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)
+[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)
 
 </div>
 

ci/run.sh

@@ -300,40 +300,6 @@ 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 {
@@ -785,8 +751,6 @@ 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,11 +402,10 @@ 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);
 
-    const int verbosity = common_log_get_verbosity_thold();
-    COM_INF("%s: verbosity = %d (adjust with the `-lv N` CLI arg)\n", __func__, verbosity);
+    COM_INF("%s: verbosity = %d (adjust with the `-lv N` CLI arg)\n", __func__, common_log_get_verbosity_thold());
 
     // device enumeration creates a primary context on CUDA backends, skip it when the caller does not own any device
-    if (print_devices && verbosity >= LOG_LEVEL_TRACE) {
+    if (print_devices) {
         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,21 +78,19 @@ 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>
-static std::shared_ptr<common_json> set_json(const std::set<T> & vals) {
+common_json_value::common_json_value(const std::set<T> & vals) : type(VAL_JSON) {
     common_json out = common_json::array();
 
     for (const auto & val : vals) {
         out.push_back(val);
     }
 
-    return std::make_shared<common_json>(std::move(out));
+    val_json = 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__> & vals) : type(VAL_JSON), val_json(set_json(vals)) {}
+#define COMMON_JSON_SET(...) template common_json_value::common_json_value(const std::set<__VA_ARGS__> &);
 
 COMMON_JSON_SET(int)
 COMMON_JSON_SET(std::string)
@@ -100,45 +98,56 @@ COMMON_JSON_SET(std::string)
 #undef COMMON_JSON_SET
 
 template <typename T>
-static std::shared_ptr<common_json> map_json(const T & vals) {
+common_json_value::common_json_value(const std::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 });
     }
 
-    return std::make_shared<common_json>(std::move(out));
+    val_json = 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__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
+#define COMMON_JSON_MAP(...) template common_json_value::common_json_value(const std::map<std::string, __VA_ARGS__> &);
 
 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__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
+#define COMMON_JSON_UMAP(...) template common_json_value::common_json_value(const std::unordered_map<std::string, __VA_ARGS__> &);
 
 COMMON_JSON_UMAP(size_t)
 
 #undef COMMON_JSON_UMAP
 
 template <typename T>
-static std::shared_ptr<common_json> vec_json(const std::vector<T> & vals) {
+common_json_value::common_json_value(const std::vector<T> & vals) : type(VAL_JSON) {
     common_json out = common_json::array();
 
     for (const auto & val : vals) {
         out.push_back(val);
     }
 
-    return std::make_shared<common_json>(std::move(out));
+    val_json = 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__> & vals) : type(VAL_JSON), val_json(vec_json(vals)) {}
+#define COMMON_JSON_VEC(...) template common_json_value::common_json_value(const std::vector<__VA_ARGS__> &);
 
 COMMON_JSON_VEC(int)
 COMMON_JSON_VEC(unsigned char)
@@ -395,6 +404,10 @@ 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;
@@ -402,7 +415,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 { return guard([&] { return as_json(this).get<__VA_ARGS__>(); }); }
+#define COMMON_JSON_GET(...) template __VA_ARGS__ common_json::get<__VA_ARGS__>() const;
 
 COMMON_JSON_GET(bool)
 COMMON_JSON_GET(int)
@@ -422,12 +435,3 @@ 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,14 +221,16 @@ 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;
+    operator std::string() const { return get<std::string>(); }
 
     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;
+    std::string value(const std::string & key, const char * def) const {
+        return contains(key) ? at(key).get<std::string>() : std::string(def);
+    }
 
     // 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,36 +112,12 @@ 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)
-        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
+        # 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)
 
     def set_vocab(self):
         return self._set_vocab_glm()
@@ -177,22 +153,10 @@ 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)
 
-        if not self.no_mtp and (num_nextn_predict_layers := self.hparams.get("num_nextn_predict_layers")) is not None:
+        # NextN/MTP prediction layers
+        if (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
@@ -384,7 +348,6 @@ 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

ggml/src/ggml-backend-meta.cpp

@@ -602,40 +602,27 @@ 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: {
-                int64_t base_ne_in = 1;
-                for (int dim = 0; dim <= src_ss[0].axis; dim++) {
+                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++) {
                     base_ne_in *= tensor->src[0]->ne[dim];
                 }
-                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.
+                base_ne_in /= src_ss[0].nr[0];
                 int64_t base_ne_out = 1;
                 for (int dim = 0; dim < GGML_MAX_DIMS; dim++) {
-                    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};
+                    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};
                     }
-                    if (base_ne_out > base_ne_in) {
+                    if (base_ne_out_next > 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));
             }
@@ -805,7 +792,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++) {
@@ -815,9 +802,6 @@ 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;
         }
@@ -1368,29 +1352,15 @@ 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);
-            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);
+            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);
             }
             for (size_t j = 0; j < n_bufs; j++) {
                 ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j);
-                const float * partial = has_contributor_mask && split_state.ne[j] == 0 ? zero.data() : tmp.data();
-                ggml_backend_tensor_set(simple_tensor, partial, offset, size);
+                ggml_backend_tensor_set(simple_tensor, tmp.data(), offset, size);
             }
         } break;
         default: {

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

@@ -38,7 +38,6 @@
 #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"
@@ -2327,9 +2326,6 @@ 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;
@@ -5249,7 +5245,6 @@ 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

@@ -1,85 +0,0 @@
-#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

@@ -1,5 +0,0 @@
-#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,7 +955,6 @@ 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;
@@ -5631,7 +5630,6 @@ 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);
 
@@ -11338,11 +11336,6 @@ 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;
@@ -12246,7 +12239,6 @@ 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:
@@ -13119,17 +13111,6 @@ 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);
@@ -15539,10 +15520,6 @@ 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);
@@ -18469,7 +18446,6 @@ 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:
@@ -19252,8 +19228,6 @@ 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

@@ -1,43 +0,0 @@
-#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,7 +896,6 @@ 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
+        # NextN/MTP tensors - preserved but unused
         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 [TAG_LLAMA_SEQ_ID_NEG]
+    // seq_id < 0 : match any sequence
     // 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/0dccaa9aa176dd6d7ef8afeca3c18d6e80a32795/subprocess.h": "vendor/sheredom/subprocess.h",
+    "https://raw.githubusercontent.com/sheredom/subprocess.h/9ce0d701b6fb10f8f8c4445edd31e7c60a1237e3/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",

src/llama-context.cpp

@@ -3218,6 +3218,8 @@ 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);
     }
@@ -3226,6 +3228,8 @@ 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,33 +599,6 @@ 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());
@@ -635,10 +608,7 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
                 overlap_cur_reads.begin(), overlap_cur_reads.end());
     }
 
-    // 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);
+    plan.n_kv = GGML_PAD(plan.n_kv, 256u);
 
     std::sort(persist_rows.begin(), persist_rows.end(),
             [](const persist_row & a, const persist_row & b) {
@@ -650,26 +620,16 @@ 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) {
-        // 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;
-                    }
-                }
+        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;
             }
 
-            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;
+            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;
             // 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) {
@@ -679,33 +639,35 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
 
             std::vector<uint32_t> token_idxs;
             token_idxs.reserve(ubatch.n_tokens);
-            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);
-                    }
+            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 = (int32_t) (stream_off + r);
+                    int32_t src;
+                    if (d <= n_seq_tokens) {
+                        const uint32_t prefix = n_seq_tokens - d;
+                        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);
+                            }
                         }
-                    }
-
-                    if (n_seq_tokens == 0) {
-                        // no-op: copy the snapshot plane onto itself
-                        src = (int32_t) (dst_plane + stream_off + r);
+                    } else {
+                        const int64_t src_plane = (int64_t) (d - n_seq_tokens)*state_rows;
+                        src = (int32_t) (src_plane + stream_off + r);
                     }
 
                     plan.state_snapshot_src_idxs.push_back(src);
@@ -721,16 +683,10 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_comp_plan(
     }();
 
     if (debug) {
-        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,
+        LLAMA_LOG_INFO("%s: ratio=%u, n_tokens=%u, state_persist_dst=%s, state_write_pos=%s\n",
+                __func__, ratio, ubatch.n_tokens,
                 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;
@@ -748,17 +704,8 @@ 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));
-
-        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;
-            }
-        }
+        plans.push_back(dsv4_build_comp_plan(ubatch, ratio, overlap, state_size, kv_size, n_stream, n_rs_seq, rs_idx));
     }
 
     return plans;
@@ -856,15 +803,16 @@ static llama_kv_cache_dsv4_context::comp_plan dsv4_build_reserve_comp_plan(
         return plan;
     }
 
-    // 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 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));
 
     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*n_stream : 0;
-    const size_t n_snapshot = (size_t) n_rs_seq*state_size*n_stream;
+    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);
 
     plan.state_pos .resize(ubatch.n_tokens);
     plan.state_persist_src_idxs.resize(n_persist);
@@ -1408,9 +1356,7 @@ llama_memory_context_ptr llama_kv_cache_dsv4::init_batch(
             if (has_coupled) {
                 ubatch = balloc.split_seq(n_ubatch);
             } else {
-                // [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);
+                ubatch = balloc.split_equal(n_ubatch, raw_per_seq || comp_per_seq, 0);
             }
 
             if (ubatch.n_tokens == 0) {
@@ -1487,11 +1433,6 @@ 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;
@@ -1653,7 +1594,9 @@ 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) {
-        clear_compressed(seq_id, true);
+        kv_csa->clear(true);
+        kv_hca->clear(true);
+        kv_lid->clear(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,7 +383,6 @@ 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) {
@@ -2044,7 +2043,6 @@ 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,14 +158,13 @@ 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) {
-        set_rs_idx(seq_id, 0);
+        if (seq_id >= 0) {
+            set_rs_idx(seq_id, 0);
+        } else {
+            std::fill(rs_idx.begin(), rs_idx.end(), 0);
+        }
     }
 
     // models like Mamba or RWKV can't have a state partially erased at the end
@@ -182,9 +181,7 @@ 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);
-                // pending rollback is single-use
-                const bool pending = rs_idx[seq_id] != 0;
-                if (!pending && rollback >= 1 && rollback <= (llama_pos) n_rs_seq) {
+                if (rollback >= 1 && rollback <= (llama_pos) n_rs_seq) {
                     set_rs_idx(seq_id, (uint32_t) rollback);
                     cell.pos = p0 - 1;
                     return true;
@@ -393,17 +390,10 @@ 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) {
-        std::fill(rs_idx.begin(), rs_idx.end(), 0);
+    if (seq_id < 0 || (size_t) seq_id >= rs_idx.size()) {
         return;
     }
-
-    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;
+    rs_idx[seq_id] = (idx > n_rs_seq) ? n_rs_seq : idx;
 }
 
 std::map<ggml_backend_buffer_type_t, size_t> llama_memory_recurrent::memory_breakdown() const {
@@ -752,7 +742,6 @@ 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;
@@ -838,7 +827,6 @@ 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 {
@@ -848,7 +836,11 @@ void llama_memory_recurrent::state_read(llama_io_read_i & io, llama_seq_id seq_i
     }
 
     if (n_rs_seq != 0) {
-        set_rs_idx(seq_id, 0);
+        if (seq_id == -1) {
+            std::fill(rs_idx.begin(), rs_idx.end(), 0);
+        } else {
+            set_rs_idx(seq_id, 0);
+        }
     }
 }
 

src/llama-model-saver.cpp

@@ -293,14 +293,6 @@ 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;
 
@@ -430,9 +422,6 @@ 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, "ffn_down_exps.weight");
+            return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_PARTIAL);
         }
 
         // output
@@ -554,9 +554,6 @@ 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)) {
@@ -645,12 +642,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 {granularity_head};
+                return {std::lcm(n_embd_q, blck_size_perf)/n_embd_q * n_gqa};
             }
+
+            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:
@@ -663,13 +660,6 @@ 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) ||
@@ -738,16 +728,6 @@ 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,4 +1,3 @@
-#include "llama-hparams.h"
 #include "models.h"
 
 #include "llama-kv-cache-dsv4.h"
@@ -59,7 +58,6 @@ 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,19 +29,10 @@ void llama_model_glm4_moe::load_arch_hparams(llama_model_loader & ml) {
     }
 }
 
-void llama_model_glm4_moe::load_arch_tensors(llama_model_loader & ml) {
+void llama_model_glm4_moe::load_arch_tensors(llama_model_loader &) {
     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");
@@ -56,9 +47,16 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader & ml) {
         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);
 
@@ -112,186 +110,24 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader & ml) {
             layer.ffn_up   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), { n_embd, n_ff }, flags);
         }
 
-        // NextN/MTP tensors
+        // NextN/MTP tensors (preserved but unused) - conditionally load for last nextn_predict_layers
         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 }, 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);
+            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);
         }
     }
 }
 
 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();
 
@@ -318,7 +154,8 @@ llama_model_glm4_moe::graph::graph(const llama_model & model, const llm_graph_pa
 
     ggml_tensor * inp_out_ids = build_inp_out_ids();
 
-    // NextN layers are processed by graph_mtp.
+    // Only process up to last layer (skip final NextN layer)
+    // Final layer tensors are loaded but not processed in forward pass
     for (int il = 0; il < n_layer; ++il) {
         ggml_tensor * inpSA = inpL;
 
@@ -368,7 +205,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 && (!cparams.embeddings_nextn || cparams.embeddings_nextn_masked)) {
+        if (il == n_layer - 1 && inp_out_ids) {
             cur   = ggml_get_rows(ctx0, cur, inp_out_ids);
             inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
         }
@@ -428,13 +265,6 @@ 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,10 +1412,6 @@ 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,15 +228,6 @@ 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)

tests/test-llama-archs.cpp

@@ -101,15 +101,6 @@ 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
@@ -129,12 +120,6 @@ 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));
@@ -175,7 +160,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_kv);
+        ms.add_kv(LLM_KV_ATTENTION_HEAD_COUNT_KV, n_head);
     }
 
     ms.add_kv(LLM_KV_ATTENTION_MAX_ALIBI_BIAS, 8.0f);
@@ -208,10 +193,6 @@ 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));
@@ -248,20 +229,6 @@ 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));
@@ -280,7 +247,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,         arch == LLM_ARCH_DEEPSEEK4 ? uint32_t(4) : uint32_t(2)); // sqrtsoftplus : sigmoid
+        ms.add_kv(LLM_KV_EXPERT_GATING_FUNC,         uint32_t(2)); // sigmoid
         ms.add_kv(LLM_KV_EXPERT_GROUP_SCALE,         1.0f);
         ms.add_kv(LLM_KV_EXPERTS_PER_GROUP,          uint32_t(1));
     }
@@ -418,7 +385,6 @@ 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:
@@ -504,6 +470,9 @@ 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,178 +35,6 @@ 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");
 
@@ -392,10 +220,5 @@ 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/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,
-    RESIZE_ALGO_BICUBIC,
+    RESIZE_ALGO_BILINEAR, // stretch to target resolution
+    RESIZE_ALGO_BICUBIC, // center-crop when aspect ratio doesn't match
+    RESIZE_ALGO_BICUBIC_PILLOW,
     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_BICUBIC;
+    resize_algo image_resize_algo_ov = RESIZE_ALGO_BILINEAR;
     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,18 +1420,20 @@ 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_BICUBIC;
+                            hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         } 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_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                     } break;
                 case PROJECTOR_TYPE_MINICPMV:
                     {
@@ -1488,7 +1490,6 @@ 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();
@@ -1515,7 +1516,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_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         hparams.rope_theta = 10000.0f;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge, false);
                         hparams.set_limit_image_tokens(8, 1024);
@@ -1543,7 +1544,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_DOTS3NOTE_V:
                     {
                         hparams.rope_theta = 10000.0f;
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
                         get_u32(KEY_SPATIAL_MERGE_SIZE, hparams.n_merge);
                         get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);
                         get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels);
@@ -1561,7 +1562,7 @@ struct clip_model_loader {
                     } break;
                 case PROJECTOR_TYPE_KIMIVL:
                     {
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         hparams.rope_theta = 10000.0f;
                         get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
                         // TODO: check kimivl preprocessor for exact values
@@ -1600,7 +1601,7 @@ struct clip_model_loader {
                     {
                         hparams.rope_theta = 100.0f;
                         hparams.n_merge = 3; // pooling_kernel_size
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         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
@@ -1617,7 +1618,6 @@ 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_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         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;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
                         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;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
                         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 the image is stretched to fit
-                        // fixed square size (no padding)
+                        // reka model performs better when using resize_bicubic, which stretches
+                        // the image to fit fixed square size
                         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_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         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,7 +1729,6 @@ 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;
@@ -1841,7 +1840,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_PADDLEOCR:
                     {
                         hparams.n_merge = 2;
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                         get_u32(KEY_IMAGE_MIN_PIXELS, hparams.image_min_pixels);
                         get_u32(KEY_IMAGE_MAX_PIXELS, hparams.image_max_pixels);
 
@@ -1853,7 +1852,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;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
                         hparams.image_pad_color = {127, 127, 127};
 
                         get_u32(KEY_SAM_N_BLOCK, hparams.sam_n_layer, true);
@@ -1883,7 +1882,7 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_HUNYUANVL:
                     {
                         hparams.n_merge = 2;
-                        hparams.image_resize_algo = RESIZE_ALGO_LANCZOS;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
                         hparams.image_resize_pad = PAD_NONE;
                         hparams.ffn_op = FFN_GELU;
                         hparams.set_limit_image_tokens(256, 16384);
@@ -1956,12 +1955,12 @@ struct clip_model_loader {
                 case PROJECTOR_TYPE_JANUS_PRO:
                     {
                         hparams.image_pad_color   = {127, 127, 127};
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
                     } break;
                 case PROJECTOR_TYPE_GRANITE4_VISION:
                     {
                         // SigLIP tower.
-                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC;
+                        hparams.image_resize_algo = RESIZE_ALGO_BICUBIC_PILLOW;
                         hparams.image_resize_pad = PAD_CEIL;
 
                         // NOTE: feature_layers loaded in common path as optional

tools/mtmd/mtmd-image.cpp

@@ -58,7 +58,22 @@ struct img_tool {
 
         if (padding == PAD_NONE) {
             // direct resize
-            resize_pillow(src, dst, target_resolution.width, target_resolution.height, algo);
+            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");
+            }
         } else {
             // resize with padding
             clip_image_u8 resized_image;
@@ -75,7 +90,22 @@ struct img_tool {
                 new_height = std::min(static_cast<int>(std::ceil(src.get_size().height * scale)), target_resolution.height);
             }
 
-            resize_pillow(src, resized_image, new_width, new_height, algo);
+            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");
+            }
 
             // fill dst with pad_color
             fill(dst, pad_color);
@@ -194,37 +224,152 @@ struct img_tool {
     }
 
 private:
-    // Pillow-compatible separable resampling (Bilinear, Bicubic and Lanczos)
+    // 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)
     // 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,
-            resize_algo algo) {
+            bool use_lanczos) {
         // 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;
 
-        // 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");
-        }
-
+        // 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
         // Returns filter weight for distance x from pixel center
-        // 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) {
+        auto resample_filter = [use_lanczos](double x) -> double {
+            if (use_lanczos) {
                 if (-3.0 <= x && x < 3.0) {
                     auto sinc = [](double v) {
                         if (v == 0.0) {
@@ -238,15 +383,10 @@ 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;
             }
@@ -256,6 +396,9 @@ 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;
@@ -350,92 +493,100 @@ private:
             const double fxp_scale = std::ldexp(1.0, PRECISION_BITS); // 1.0 * 2^PRECISION_BITS
 
             for (int i = 0; i < outSize * ksize; i++) {
-                // 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);
+                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);
             }
 
             return ksize;
         };
 
         // Horizontal resampling pass
-        // Resizes width from src to out_nx, preserving height
-        auto resample_horizontal = [&](const uint8_t * src, int in_nx, int in_ny,
+        // Resizes width from imIn to out_nx, preserving height
+        auto resample_horizontal = [&](const clip_image_u8 & imIn, clip_image_u8 & imOut,
                                        int out_nx,
                                        int ksize, const std::vector<int> & bounds, const std::vector<int32_t> & weights) {
-            std::vector<uint8_t> out((size_t) out_nx * in_ny * 3);
+            const int in_ny = imIn.get_size().height;
+            imOut.set_size({out_nx, in_ny}, false);
 
             // 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++) {
-                    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;
+                    // 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
 
-                    // Accumulators for RGB channels, with rounding bias (0.5 in fixed-point)
+                    // Initialize 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++) {
-                        ss0 += p[0] * k[x];
-                        ss1 += p[1] * k[x];
-                        ss2 += p[2] * k[x];
-                        p += 3;
+                        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
                     }
 
                     // Convert back from fixed-point (divide by 2^PRECISION_BITS) and clamp to [0,255]
-                    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);
+                    imOut.set_pixel(xx, yy, {clip8(ss0 >> PRECISION_BITS),
+                                             clip8(ss1 >> PRECISION_BITS),
+                                             clip8(ss2 >> PRECISION_BITS)});
                 }
             }
-
-            return out;
         };
 
         // Vertical resampling pass
-        // 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,
+        // Resizes height from imIn to out_ny, preserving width
+        auto resample_vertical = [&](const clip_image_u8 & imIn, clip_image_u8 & imOut,
                                      int out_ny,
                                      int ksize, const std::vector<int> & bounds, const std::vector<int32_t> & weight) {
-            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);
+            const int in_nx = imIn.get_size().width;
+            imOut.set_size({in_nx, out_ny}, false);
 
             // For each output row
             for (int yy = 0; yy < out_ny; yy++) {
-                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];
+                // 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
 
-                // Rounding bias (0.5 in fixed-point)
-                std::fill(acc.begin(), acc.end(), 1 << (PRECISION_BITS - 1));
+                // 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);
 
-                // 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;
+                    // 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
                     }
-                }
 
-                // 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);
+                    // 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)});
                 }
             }
-
-            return out;
         };
 
         // Main resampling logic using separable two-pass approach
@@ -459,25 +610,36 @@ private:
         }
 
         // Perform two-pass resampling
-        const uint8_t * src = img.get_ro_buf().data();
         if (need_horizontal && need_vertical) {
-            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));
+            // 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);
         } else if (need_horizontal) {
-            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));
+            // Only horizontal
+            resample_horizontal(img, dst, target_width, ksize_horiz, bounds_horiz, weights_horiz);
         } else if (need_vertical) {
-            dst.set_size({src_width, target_height}, false);
-            dst.cpy_buf(resample_vertical(src, src_width, target_height, ksize_vert, bounds_vert, weights_vert));
+            // Only vertical
+            resample_vertical(img, dst, target_height, ksize_vert, bounds_vert, weights_vert);
         } else {
             // No resizing needed - direct copy
-            dst.set_size(img.get_size(), false);
-            dst.cpy_buf(img.get_ro_buf());
+            dst.set_size(img.get_size(), img.is_placeholder());
+            if (!img.is_placeholder()) {
+                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;
+    }
 };
 
 
@@ -1102,7 +1264,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,
+                    RESIZE_ALGO_BICUBIC_PILLOW,
                     PAD_NEAREST,
                     hparams.image_pad_color);
     output.append_overview(hparams, padded, true);
@@ -1118,7 +1280,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,
+        img_tool::resize(img, refined, { tile_size * grid_w, tile_size * grid_h }, RESIZE_ALGO_BICUBIC_PILLOW,
                          PAD_NONE);
 
         for (int row = 0; row < grid_h; row++) {

tools/server/server-context.cpp

@@ -858,10 +858,8 @@ private:
     // slots / clients
     std::vector<server_slot> slots;
 
-    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 trace = 0;
+    int slots_debug = 0;
     int n_empty_consecutive = 0;
 
     std::unique_ptr<server_prompt_cache> prompt_cache;
@@ -1249,15 +1247,6 @@ 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)
         {
@@ -3190,8 +3179,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 - 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()));
+                                    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()));
 
                                     std::stringstream ss0;
                                     std::stringstream ss1;

tools/server/tests/unit/test_slot_save.py

@@ -319,6 +319,7 @@ 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
 
@@ -344,26 +345,6 @@ 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)+|(automobile)+"),
+        ("What is this: <__media__>\n", "IMG_BASE64_0",         True, "(cat)+"),
         ("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 = 512
+        server.n_batch = 32
         server.n_slots = 2
         server.n_predict = 4
         server.seed = 42

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

@@ -3,9 +3,8 @@
 		ChatAttachmentsListItem,
 		DialogChatAttachmentsPreview,
 		DialogMcpResourcePreview,
-		ScrollCarousel
+		HorizontalScrollCarousel
 	} from '$lib/components/app';
-	import { ScrollCarouselVariant } from '$lib/enums';
 	import type { DatabaseMessageExtraMcpResource } from '$lib/types';
 	import { getAttachmentDisplayItems, isMcpPrompt, isMcpResource } from '$lib/utils';
 
@@ -43,7 +42,7 @@
 		uploadedFiles = $bindable([])
 	}: Props = $props();
 
-	let carouselRef: ScrollCarousel | undefined = $state();
+	let carouselRef: HorizontalScrollCarousel | undefined = $state();
 	let mcpResourcePreviewOpen = $state(false);
 	let mcpResourcePreviewExtra = $state<DatabaseMessageExtraMcpResource | null>(null);
 	let previewFocusIndex = $state(0);
@@ -92,11 +91,11 @@
 {#if displayItems.length > 0}
 	<div class={className} {style}>
 		{#if limitToSingleRow}
-			<ScrollCarousel bind:this={carouselRef} variant={ScrollCarouselVariant.CENTER}>
+			<HorizontalScrollCarousel bind:this={carouselRef}>
 				{#each displayItems as item (item.id)}
 					{@render attachmentitem(item)}
 				{/each}
-			</ScrollCarousel>
+			</HorizontalScrollCarousel>
 		{: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,8 +1,7 @@
 <script lang="ts">
 	import { FileText, Music, Video } from '@lucide/svelte';
-	import { ScrollCarousel } from '$lib/components/app';
+	import { HorizontalScrollCarousel } from '$lib/components/app/misc';
 	import { ICON_CLASS_DEFAULT, UI_DATA_ATTRS } from '$lib/constants';
-	import { ScrollCarouselVariant } from '$lib/enums';
 
 	interface PreviewItem {
 		id: string;
@@ -34,7 +33,7 @@
 
 {#if items.length > 1}
 	<div class="sticky bottom-0 z-10 mt-4 flex-shrink-0">
-		<ScrollCarousel class="max-w-full" variant={ScrollCarouselVariant.CENTER}>
+		<HorizontalScrollCarousel class="max-w-full">
 			{#each items as item, index (item.id)}
 				<button
 					{...{ [UI_DATA_ATTRS.THUMBNAIL_INDEX]: index }}
@@ -65,6 +64,6 @@
 					{/if}
 				</button>
 			{/each}
-		</ScrollCarousel>
+		</HorizontalScrollCarousel>
 	</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 && !event.shiftKey;
+			const isWordJump = (event.altKey || event.ctrlKey) && !event.metaKey;
 			const isPlainLeft =
 				event.key === 'ArrowLeft' && !event.altKey && !event.ctrlKey && !event.metaKey;
 

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

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

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

@@ -41,7 +41,10 @@
 	let showDeleteDialog = $state(false);
 	let showEmptyFileDialog = $state(false);
 	let isEmpty = $derived(
-		showCenteredEmpty && conversationsStore.activeMessages.length === 0 && !chatStore.isLoading
+		showCenteredEmpty &&
+			!conversationsStore.activeConversation &&
+			conversationsStore.activeMessages.length === 0 &&
+			!chatStore.isLoading
 	);
 	let activeErrorDialog = $derived(chatStore.errorDialogState);
 	let isServerLoading = $derived(serverStore.loading);
@@ -294,7 +297,7 @@
 	<ServerLoadingSplash />
 {:else}
 	<div
-		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"
+		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"
 		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-muted/60 border border-border/20 shadow-sm text-accent-foreground absolute bottom-4"
+		class="h-9 w-9 rounded-full bg-accent text-accent-foreground absolute bottom-4 shadow-md"
 	/>
 </div>

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

@@ -1,136 +0,0 @@
-<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

@@ -1,156 +0,0 @@
-<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

@@ -1,30 +0,0 @@
-<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,18 +686,6 @@ 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

@@ -0,0 +1,96 @@
+<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

@@ -1,131 +0,0 @@
-<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,6 +21,13 @@
  */
 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
  *
@@ -37,13 +44,6 @@ 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, uiStore } from '$lib/stores';
+	import { chatStore, conversationsStore, deviceStore, settingsStore } from '$lib/stores';
 	import { buildConversationTree } from '$lib/utils';
 	import { circIn } from 'svelte/easing';
 	import { SvelteSet } from 'svelte/reactivity';
@@ -31,29 +31,30 @@
 		toggleSidebar: () => toggleExpandedMode()
 	});
 
+	let isExpandedMode = $state(false);
 	let hoveredTooltip = $state<string | null>(null);
 	let logoHovered = $state(false);
 
-	const isStripExpanded = $derived(uiStore.isSidebarExpanded || hoveredTooltip !== null);
+	const isStripExpanded = $derived(isExpandedMode || hoveredTooltip !== null);
 	const isOnMobile = $derived(deviceStore.isMobile);
 	const alwaysShowOnDesktop = $derived(settingsStore.config.alwaysShowSidebarOnDesktop as boolean);
 
 	$effect(() => {
 		if (alwaysShowOnDesktop && !isOnMobile) {
-			uiStore.isSidebarExpanded = true;
+			isExpandedMode = true;
 		}
 	});
 
 	function toggleExpandedMode() {
-		uiStore.isSidebarExpanded = !uiStore.isSidebarExpanded;
+		isExpandedMode = !isExpandedMode;
 
-		if (!uiStore.isSidebarExpanded) {
+		if (!isExpandedMode) {
 			hoveredTooltip = null;
 		}
 	}
 
 	$effect(() => {
-		if (!uiStore.isSidebarExpanded) {
+		if (!isExpandedMode) {
 			isSearchModeActive = false;
 			searchQuery = '';
 
@@ -65,7 +66,7 @@
 
 	$effect(() => {
 		if (deviceStore.isMobile && page.url.hash.includes(ROUTES.SEARCH)) {
-			uiStore.isSidebarExpanded = false;
+			isExpandedMode = false;
 		}
 	});
 
@@ -293,7 +294,7 @@
 		}
 
 		pendingCollapse = setTimeout(() => {
-			uiStore.isSidebarExpanded = false;
+			isExpandedMode = false;
 			pendingCollapse = null;
 		}, 100);
 	}
@@ -313,7 +314,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)]',
-			uiStore.isSidebarExpanded &&
+			isExpandedMode &&
 				(deviceStore.isStandalone
 					? 'h-[calc(100dvh-2rem)]'
 					: deviceStore.isIOSDevice
@@ -322,9 +323,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 shadow-md',
+			isStripExpanded && 'md:w-72 md:bg-muted/60 md:backdrop-blur-xl border-border shadow-md',
 			!isStripExpanded && 'md:w-12',
-			uiStore.isSidebarExpanded && 'is-expanded'
+			isExpandedMode && 'is-expanded'
 		]}
 	>
 		<div class="px-2 flex items-center justify-between">
@@ -336,26 +337,24 @@
 				onmouseleave={() => (logoHovered = false)}
 			>
 				<ActionIcon
-					icon={!uiStore.isSidebarExpanded && logoHovered && innerWidth > 768
-						? PanelLeftOpen
-						: Logo}
+					icon={!isExpandedMode && logoHovered && innerWidth > 768 ? PanelLeftOpen : Logo}
 					size="lg"
 					iconSize="h-4.5 w-4.5 md:h-4 md:w-4"
-					class="{uiStore.isSidebarExpanded
+					class="{isExpandedMode
 						? '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={uiStore.isSidebarExpanded ? ROUTES.START : undefined}
-					onclick={uiStore.isSidebarExpanded ? undefined : toggleExpandedMode}
-					tooltip={uiStore.isSidebarExpanded ? undefined : 'Open Sidebar'}
+					href={isExpandedMode ? ROUTES.START : undefined}
+					onclick={isExpandedMode ? undefined : toggleExpandedMode}
+					tooltip={isExpandedMode ? undefined : 'Open Sidebar'}
 					tooltipSide={TooltipSide.RIGHT}
-					ariaLabel={uiStore.isSidebarExpanded ? 'Go to start' : 'Expand navigation'}
+					ariaLabel={isExpandedMode ? 'Go to start' : 'Expand navigation'}
 				/>
 			</div>
 
-			{#if isOnMobile || (uiStore.isSidebarExpanded && !alwaysShowOnDesktop)}
+			{#if isOnMobile || (isExpandedMode && !alwaysShowOnDesktop)}
 				<div
 					class="flex items-center transition-all duration-150 ease-out {deviceStore.isMobile &&
-					!uiStore.isSidebarExpanded
+					!isExpandedMode
 						? 'opacity-0 h-0!'
 						: ''}"
 					in:fade={{ delay: 50, duration: 150, easing: circIn }}
@@ -378,12 +377,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 && !uiStore.isSidebarExpanded ? 'opacity-0 !h-0' : ''}"
+				: ''} {deviceStore.isMobile && !isExpandedMode ? 'opacity-0 !h-0' : ''}"
 			in:fade={{ duration: 200 }}
 			out:fade={{ duration: 200 }}
 		>
 			<SidebarNavigationActions
-				isExpandedMode={innerWidth > 768 ? uiStore.isSidebarExpanded : true}
+				isExpandedMode={innerWidth > 768 ? isExpandedMode : true}
 				class="px-2"
 				bind:isSearchModeActive
 				bind:searchQuery
@@ -392,7 +391,7 @@
 					searchQuery = '';
 				}}
 				onSearchClick={() => {
-					uiStore.isSidebarExpanded = true;
+					isExpandedMode = true;
 					isSearchModeActive = true;
 				}}
 				onNewChat={() => {
@@ -402,7 +401,7 @@
 				}}
 			/>
 
-			{#if uiStore.isSidebarExpanded || isOnMobile}
+			{#if isExpandedMode || 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 { SidebarAction, TooltipSide } from '$lib/enums';
-	import { conversationsStore, deviceStore } from '$lib/stores';
+	import { TooltipSide } from '$lib/enums';
+	import { deviceStore } from '$lib/stores';
 	import type { Component } from 'svelte';
 	import { onMount } from 'svelte';
 	import { circIn } from 'svelte/easing';
@@ -109,20 +109,14 @@
 			{@const isActive = isItemActive(item)}
 			{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
 			{@const itemHref = isSearchOnMobile ? ROUTES.SEARCH : item.route}
-			{@const itemOnClick =
-				item.action === SidebarAction.NEW_CHAT
-					? () => {
-							onNewChat?.();
-							void conversationsStore.openNewChat();
-						}
-					: item.route
-						? () => {
-								onNewChat?.();
-								goto(item.route!);
-							}
-						: isSearchOnMobile
-							? undefined
-							: onSearchClick}
+			{@const itemOnClick = 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,
@@ -163,20 +157,14 @@
 		{#each SIDEBAR_ACTIONS_ITEMS as item, i (item.tooltip)}
 			{@const isActive = isItemActive(item)}
 			{@const isSearchOnMobile = item.icon === Search && deviceStore.isMobile}
-			{@const itemOnClick =
-				item.action === SidebarAction.NEW_CHAT
-					? () => {
-							onNewChat?.();
-							void conversationsStore.openNewChat();
-						}
-					: item.route
-						? () => {
-								onNewChat?.();
-								goto(item.route!);
-							}
-						: isSearchOnMobile
-							? undefined
-							: onSearchClick}
+			{@const itemOnClick = 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,6 +1,5 @@
 <script lang="ts">
-	import { Settings } from '@lucide/svelte';
-	import { ScrollCarousel } from '$lib/components/app';
+	import { ChevronLeft, ChevronRight, Settings } from '@lucide/svelte';
 	import { ICON_CLASS_DEFAULT, UI_DATA_ATTRS } from '$lib/constants';
 	import { BooleanString } from '$lib/enums';
 	import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
@@ -45,42 +44,70 @@
 	</div>
 
 	<div class="border-b border-border/30 py-2">
-		<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 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>
 	</div>
 </div>

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

@@ -1,18 +0,0 @@
-/** 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,11 +26,5 @@ 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,7 +4,6 @@
 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,6 +4,8 @@ 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;
@@ -13,6 +15,8 @@ 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,7 +11,6 @@ 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,13 +255,6 @@ 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,7 +22,6 @@ 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 { SidebarAction, ToolSource } from '$lib/enums';
+import { ToolSource } from '$lib/enums/tools.enums';
 import type { DesktopIconStripItem } from '$lib/types';
 
 export const FORK_TREE_DEPTH_PADDING = 8;
@@ -10,7 +10,6 @@ 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',
@@ -56,12 +55,7 @@ export const ICON_STRIP_TRANSITION_DELAY_MULTIPLIER = 50;
 export const MAX_HEIGHT_CODE_BLOCK = '22rem';
 
 export const SIDEBAR_ACTIONS_ITEMS: DesktopIconStripItem[] = [
-	{
-		action: SidebarAction.NEW_CHAT,
-		icon: SquarePen,
-		keys: ['shift', 'cmd', 'o'],
-		tooltip: 'New chat'
-	},
+	{ icon: SquarePen, keys: ['shift', 'cmd', 'o'], route: ROUTES.NEW_CHAT, tooltip: 'New chat' },
 	{ icon: Search, keys: ['cmd', 'k'], tooltip: 'Search' },
 	{
 		activeRouteId: '/mcp-servers',

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

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

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

@@ -7,8 +7,6 @@ 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',
@@ -18,7 +16,5 @@ export enum KeyboardKey {
 	O_LOWER = 'o',
 	O_UPPER = 'O',
 	SPACE = ' ',
-	TAB = 'Tab',
-	X_LOWER = 'x',
-	X_UPPER = 'X'
+	TAB = 'Tab'
 }

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

@@ -11,21 +11,6 @@ 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,7 +1,6 @@
-import { page } from '$app/state';
-import { NEW_CHAT_TAB_ID } from '$lib/constants';
+import { goto } from '$app/navigation';
+import { ROUTES } from '$lib/constants';
 import { KeyboardKey } from '$lib/enums';
-import { conversationsStore, settingsStore, tabsStore } from '$lib/stores';
 
 interface KeyboardShortcutsCallbacks {
 	activateSearchMode?: () => void;
@@ -10,8 +9,6 @@ interface KeyboardShortcutsCallbacks {
 	deleteActiveConversation?: () => void;
 	navigateToPrevConversation?: () => void;
 	navigateToNextConversation?: () => void;
-	navigateToPrevTab?: () => void;
-	navigateToNextTab?: () => void;
 	toggleSidebar?: () => void;
 }
 
@@ -37,7 +34,7 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
 		) {
 			event.preventDefault();
 
-			void conversationsStore.openNewChat();
+			goto(ROUTES.NEW_CHAT);
 		}
 
 		if (event.shiftKey && isCmdOrCtrl && event.key === KeyboardKey.E_UPPER) {
@@ -45,28 +42,6 @@ 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 &&
@@ -85,16 +60,6 @@ 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,7 +2,6 @@ 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;
@@ -16,6 +15,18 @@ 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;
 
@@ -25,21 +36,10 @@ 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) return;
-
-		updateScrollButtons();
-
-		const observer = new ResizeObserver(() => updateScrollButtons());
-
-		observer.observe(scrollContainer);
-
-		if (contentContainer) observer.observe(contentContainer);
-
-		return () => observer.disconnect();
+		if (scrollContainer) {
+			updateScrollButtons();
+		}
 	});
 
 	return {
@@ -49,18 +49,14 @@ 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, MCP_SERVERS) live in ROUTES constants;
+ * Static routes (START, NEW_CHAT, 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,7 +21,6 @@ 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';
@@ -106,8 +105,8 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 	/**
 	 * Deletes multiple conversations in sequence.
-	 * Mirrors deleteConversation() per-id; navigates to the new-chat screen only
-	 * if the currently-open chat was among the deleted ones.
+	 * Mirrors deleteConversation() per-id; navigates to NEW_CHAT only if the
+	 * currently-open chat was among the deleted ones.
 	 * @param convIds - Conversation IDs to delete
 	 */
 	async bulkDeleteConversations(convIds: string[]): Promise<void> {
@@ -139,13 +138,8 @@ class ConversationsStore implements ConversationsPreferencesHost {
 			this.notifyConversationsDeleted([...idsToRemove]);
 
 			if (activeWasDeleted) {
-				const activeId = this.activeConversation!.id;
-
-				tabsStore.removeTabs([...idsToRemove].filter((id) => id !== activeId));
 				this.clearActiveConversation();
-				await tabsStore.close(activeId, activeId);
-			} else {
-				tabsStore.removeTabs([...idsToRemove]);
+				await goto(ROUTES.NEW_CHAT);
 			}
 
 			toast.success(
@@ -282,12 +276,11 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 			this.clearActiveConversation();
 			this.conversations = [];
-			tabsStore.clear();
 			this.notifyConversationsDeleted(allIds);
 
 			toast.success('All conversations deleted');
 
-			await goto(ROUTES.START);
+			await goto(ROUTES.NEW_CHAT);
 		} catch (error) {
 			console.error('Failed to delete all conversations:', error);
 			toast.error('Failed to delete conversations');
@@ -320,13 +313,8 @@ 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 tabsStore.close(activeId, activeId);
-				} else {
-					tabsStore.removeTabs([...idsToRemove]);
+					await goto(ROUTES.NEW_CHAT);
 				}
 
 				this.notifyConversationsDeleted([...idsToRemove]);
@@ -345,9 +333,7 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 				if (this.activeConversation?.id === convId) {
 					this.clearActiveConversation();
-					await tabsStore.close(convId, convId);
-				} else {
-					tabsStore.removeTabs([convId]);
+					await goto(ROUTES.NEW_CHAT);
 				}
 
 				this.notifyConversationsDeleted([convId]);
@@ -585,15 +571,6 @@ 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.
 	 */
@@ -699,6 +676,14 @@ class ConversationsStore implements ConversationsPreferencesHost {
 		);
 	}
 
+	/**
+	 *
+	 *
+	 * Import & Export
+	 *
+	 *
+	 */
+
 	/**
 	 * Updates the current node of the active conversation
 	 * @param nodeId - The new current node ID
@@ -732,14 +717,6 @@ 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,13 +152,11 @@ export class ConversationPreferences {
 			return;
 		}
 
-		const id = this.host.activeConversation.id;
-
-		this.host.applyConversationUpdate(id, {
+		this.host.applyConversationUpdate(this.host.activeConversation.id, {
 			cwd: trimmed
 		});
 
-		await DatabaseService.updateConversation(id, {
+		await DatabaseService.updateConversation(this.host.activeConversation.id, {
 			cwd: trimmed
 		});
 
@@ -204,15 +202,12 @@ export class ConversationPreferences {
 			}
 		}
 
-		const overrides = newOverrides.length > 0 ? newOverrides : undefined;
-		const id = this.host.activeConversation.id;
-
-		this.host.applyConversationUpdate(id, {
-			mcpServerOverrides: overrides
+		await DatabaseService.updateConversation(this.host.activeConversation.id, {
+			mcpServerOverrides: newOverrides.length > 0 ? newOverrides : undefined
 		});
 
-		await DatabaseService.updateConversation(id, {
-			mcpServerOverrides: overrides
+		this.host.applyConversationUpdate(this.host.activeConversation.id, {
+			mcpServerOverrides: newOverrides.length > 0 ? newOverrides : undefined
 		});
 	}
 
@@ -229,13 +224,11 @@ export class ConversationPreferences {
 			return;
 		}
 
-		const id = this.host.activeConversation.id;
-
-		this.host.applyConversationUpdate(id, {
+		this.host.applyConversationUpdate(this.host.activeConversation.id, {
 			reasoningEffort: effort
 		});
 
-		await DatabaseService.updateConversation(id, {
+		await DatabaseService.updateConversation(this.host.activeConversation.id, {
 			reasoningEffort: effort
 		});
 	}

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

@@ -22,9 +22,6 @@ 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';
 
@@ -43,9 +40,6 @@ 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,8 +1,21 @@
+/**
+ * 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';
@@ -20,12 +33,7 @@ export function initStores(): Promise<void> {
 		permissionsStore.initialize();
 		toolsStore.initialize();
 		void versionStore.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));
-		});
+		void conversationsStore.initialize();
 	})();
 
 	return startup;

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

@@ -1,154 +0,0 @@
-/**
- * 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

@@ -1,14 +0,0 @@
-/**
- * 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,4 +1,3 @@
-import type { SidebarAction } from '$lib/enums';
 import type { Component } from 'svelte';
 
 /**
@@ -8,8 +7,6 @@ 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,34 +1,12 @@
 <script lang="ts">
 	import { page } from '$app/state';
-	import { ChatScreen, ChatTabs } from '$lib/components/app';
-	import { NEW_CHAT_TAB_ID } from '$lib/constants';
-	import { settingsStore, tabsStore } from '$lib/stores';
+	import { ChatScreen } from '$lib/components/app';
 
 	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>
 
-<div class={showTabs ? 'md:[--chat-tabs-offset:1.25rem]' : ''}>
-	{#if showTabs}
-		<ChatTabs />
-	{/if}
-
-	<ChatScreen {showCenteredEmpty} />
-</div>
+<ChatScreen {showCenteredEmpty} />
 
 {@render children?.()}

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

@@ -8,19 +8,24 @@
 
 	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 params after handling the deep link so a refresh does not replay them
+	/**
+	 * Clear URL params after message is sent to prevent re-sending on refresh
+	 */
 	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(), {});
 	}
@@ -35,8 +40,8 @@
 				try {
 					await modelsStore.selectModelById(model.id);
 
-					// with ?load=true in router mode, start loading right away so the
-					// model is ready sooner; not awaited so the UI stays usable
+					// with ?load=true, start loading right away so the model is ready sooner;
+					// not awaited, so the UI stays usable during the load
 					if (
 						loadParam === 'true' &&
 						serverStore.isRouterMode &&
@@ -61,12 +66,11 @@
 			}
 		}
 
-		// ?q= creates the conversation, the chat route sends the prompt once the
-		// conversation id is in the URL
+		// Handle ?q= parameter - create new conversation and send message
 		if (qParam !== null) {
 			await conversationsStore.createConversation();
 			clearUrlParams();
-		} else if (modelParam) {
+		} else if (modelParam || newChatParam === 'true') {
 			clearUrlParams();
 		}
 	}
@@ -81,7 +85,7 @@
 
 		await modelsStore.fetch();
 
-		if (qParam !== null || modelParam !== null) {
+		if (qParam !== null || modelParam !== null || newChatParam === 'true') {
 			await handleUrlParams();
 		}
 

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

@@ -11,7 +11,6 @@
 		FAVICON_PATHS,
 		FAVICON_SELECTORS,
 		HEADERS,
-		NEW_CHAT_TAB_ID,
 		ROUTES,
 		SETTINGS_KEYS,
 		TOOLTIP_DELAY_DURATION
@@ -27,7 +26,6 @@
 		modelsStore,
 		serverStore,
 		settingsStore,
-		tabsStore,
 		versionStore
 	} from '$lib/stores';
 	import { initStores } from '$lib/stores/init';
@@ -76,27 +74,6 @@
 		}
 	}
 
-	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;
 
@@ -119,31 +96,15 @@
 		if (targetIdx >= 0 && targetIdx < allConvs.length) {
 			goto(RouterService.chat(allConvs[targetIdx].id));
 		} else {
-			conversationsStore.openNewChat();
+			goto(ROUTES.NEW_CHAT);
 		}
 	}
 
-	// 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),
-		navigateToNextTab: () => navigateToTab(1),
-		navigateToPrevConversation: () => navigateToConversation(-1),
-		navigateToPrevTab: () => navigateToTab(-1)
+		navigateToPrevConversation: () => navigateToConversation(-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 new-chat tab without a history entry.
+	// in-place search, so bounce back to a chat.
 	$effect(() => {
 		if (browser && !deviceStore.isMobile) {
-			goto(ROUTES.START, { replaceState: true });
+			goto(ROUTES.NEW_CHAT, { replaceState: true });
 		}
 	});
 
@@ -66,7 +66,7 @@
 		if (history.length > 1) {
 			history.back();
 		} else {
-			conversationsStore.openNewChat();
+			goto(ROUTES.NEW_CHAT);
 		}
 	}
 </script>

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

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

vendor/sheredom/subprocess.h

@@ -275,46 +275,13 @@ 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 SUBPROCESS_SPAWN_VIA_FORK
-#define SUBPROCESS_HAVE_CWD 1
-#elif defined(__GLIBC__)
+#if defined(__GLIBC__)
 #if __GLIBC_PREREQ(2, 29)
 #define SUBPROCESS_HAVE_CWD 1
 #else
@@ -327,13 +294,10 @@ subprocess_weak int subprocess_alive(struct subprocess_s *const process);
 #endif
 #endif
 
-/* 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. */
+/* 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. */
 #if !defined(SUBPROCESS_SPAWN_REPORTS_EXEC_ERRORS)
-#if SUBPROCESS_SPAWN_VIA_FORK
-#define SUBPROCESS_SPAWN_REPORTS_EXEC_ERRORS 1
-#elif defined(__GLIBC__)
+#if defined(__GLIBC__)
 #if __GLIBC_PREREQ(2, 24)
 #define SUBPROCESS_SPAWN_REPORTS_EXEC_ERRORS 1
 #else
@@ -378,14 +342,6 @@ 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"
@@ -395,7 +351,6 @@ 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
@@ -447,11 +402,6 @@ 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;
@@ -501,14 +451,6 @@ __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 *);
@@ -725,104 +667,12 @@ 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,
@@ -842,7 +692,6 @@ 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;
@@ -850,8 +699,6 @@ 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};
@@ -859,11 +706,6 @@ 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,
@@ -1238,44 +1080,6 @@ 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
@@ -1286,7 +1090,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,
-                              &startInfoEx), // STARTUPINFOEX pointer
+                              &startInfo), // STARTUPINFO pointer
           SUBPROCESS_PTR_CAST(LPPROCESS_INFORMATION, &processInfo))) {
     result = subprocess_error_from_windows_error(GetLastError());
     if (subprocess_error_unknown == result) {
@@ -1295,9 +1099,6 @@ 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;
 
@@ -1327,10 +1128,6 @@ 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);
   }
@@ -1376,20 +1173,15 @@ 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};
-  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
+  char *const *used_environment;
 
   async_no_wait = subprocess_option_enable_async_no_wait ==
                   (options & subprocess_option_enable_async_no_wait);
@@ -1410,13 +1202,13 @@ cleanup:
 
   memset(out_process, 0, sizeof(*out_process));
 
-  if (0 != subprocess_pipe_cloexec(stdinfd)) {
+  if (0 != pipe(stdinfd)) {
     saved_errno = errno;
     result = subprocess_error_pipe;
     goto cleanup;
   }
 
-  if (0 != subprocess_pipe_cloexec(stdoutfd)) {
+  if (0 != pipe(stdoutfd)) {
     saved_errno = errno;
     result = subprocess_error_pipe;
     goto cleanup;
@@ -1424,7 +1216,7 @@ cleanup:
 
   if (subprocess_option_combined_stdout_stderr !=
       (options & subprocess_option_combined_stdout_stderr)) {
-    if (0 != subprocess_pipe_cloexec(stderrfd)) {
+    if (0 != pipe(stderrfd)) {
       saved_errno = errno;
       result = subprocess_error_pipe;
       goto cleanup;
@@ -1448,136 +1240,6 @@ 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;
@@ -1591,7 +1253,7 @@ cleanup:
 
   // Set working directory
   if (process_cwd) {
-#if SUBPROCESS_ADDCHDIR_IS_POSIX
+#if defined(__NetBSD__) || (defined(__APPLE__) && MAC_OS_X_VERSION_MIN_REQUIRED >= 260000)
     posix_error = posix_spawn_file_actions_addchdir(&actions, process_cwd);
 #elif !SUBPROCESS_HAVE_CWD
     posix_error = ENOSYS;
@@ -1744,7 +1406,6 @@ cleanup:
 #ifdef __clang__
 #pragma clang diagnostic pop
 #endif
-#endif /* SUBPROCESS_SPAWN_VIA_FORK */
 
   // Close the stdin read end
   close(stdinfd[0]);
@@ -1819,21 +1480,9 @@ 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) {