DeepseekV4: fix rollback with multi-seq (#26756)

* DeepseekV4: fix rollback with multi-seq

* fix model loading

* make pending rollback single use

* only clear cache for seq_id for full load

* add assert for compress ratio

* make graph topology static

* pass true instead of flags in clear_compressed

* cont : clean-up + TODOs

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>

ci/run.sh

@@ -300,6 +300,31 @@ 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
+
+    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
+
+    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 CUDA devices\n'
+    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
+    gg_printf '```\n'
+    gg_printf '%s\n' "$(cat $OUT/${ci}.log)"
+    gg_printf '```\n'
+}
+
 # test_scripts
 
 function gg_run_test_scripts {
@@ -751,6 +776,10 @@ ret=0
 test $ret -eq 0 && gg_run ctest_debug
 test $ret -eq 0 && gg_run ctest_release
 
+if [ ! -z ${GG_BUILD_CUDA} ]; then
+    test $ret -eq 0 && gg_run test_llama_archs_tensor_split
+fi
+
 if [ ! -z ${GG_BUILD_HIGH_PERF} ]; then
     test $ret -eq 0 && gg_run test_backend_ops_cpu
 fi

common/json.cpp

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

common/json.h

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

ggml/src/ggml-backend-meta.cpp

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

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

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

ggml/src/ggml-cuda/pool1d.cu

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

ggml/src/ggml-cuda/pool1d.cuh

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

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

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

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

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

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

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

include/llama.h

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

scripts/sync_vendor.py

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

src/llama-context.cpp

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

src/llama-kv-cache-dsv4.cpp

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

src/llama-kv-cache.cpp

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

src/llama-memory-recurrent.cpp

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

src/llama-model-saver.cpp

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

src/llama-model.cpp

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

src/models/deepseek4.cpp

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

tests/CMakeLists.txt

@@ -228,6 +228,15 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
     set_tests_properties(test-recurrent-state-rollback-nemotron-h PROPERTIES
         FIXTURES_REQUIRED generate-models
     )
+    llama_test(
+        test-recurrent-state-rollback
+        NAME test-recurrent-state-rollback-dsv4
+        LABEL main
+        ARGS -m "${MODEL_DIR}/deepseek4-moe.gguf"
+    )
+    set_tests_properties(test-recurrent-state-rollback-dsv4 PROPERTIES
+        FIXTURES_REQUIRED generate-models
+    )
 endif()
 
 llama_build_and_test(test-chat-peg-parser.cpp peg-parser/simple-tokenize.cpp)

tests/test-llama-archs.cpp

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

tests/test-recurrent-state-rollback.cpp

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

@@ -16,5 +16,7 @@ export enum KeyboardKey {
 	O_LOWER = 'o',
 	O_UPPER = 'O',
 	SPACE = ' ',
-	TAB = 'Tab'
+	TAB = 'Tab',
+	X_LOWER = 'x',
+	X_UPPER = 'X'
 }

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

vendor/sheredom/subprocess.h

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