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>
[Tensor parallel] Fix meta tensor split state propagation (#27574 )
2026-08-23 13:12:34 +02:00 · 2026-08-23 13:57:49 +03:00 · 2026-08-23 18:49:12 +08:00 · 2026-08-23 10:46:49 +02:00 · 2026-08-23 10:38:29 +03:00 · 2026-08-23 10:37:32 +03:00
125 changed files with 3349 additions and 934 deletions
@@ -774,6 +774,7 @@ jobs:
        with:
          key: windows-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
          evict-old-files: 1d
+          max-size: "1G"

      # - name: Cache ROCm Installation
      #   id: cache-rocm
@@ -1286,123 +1287,134 @@ jobs:
        with:
          key: release-ubuntu-24.04-sycl-${{ matrix.build }}

-#   ubuntu-22-rocm:
-#     needs: [check-release, get-version]
-#     if: ${{ needs.check-release.outputs.should_release == 'true' }}
+  ubuntu-22-rocm:
+    needs: [check-release, get-version]
+    if: ${{ needs.check-release.outputs.should_release == 'true' }}

-#     runs-on: ubuntu-22.04
+    runs-on: ubuntu-22.04

-#     permissions:
-#       actions: write
+    permissions:
+      actions: write

-#     strategy:
-#       matrix:
-#         include:
-#           - ROCM_VERSION: "7.14.0"
-#             gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
-#             build: 'x64'
+    strategy:
+      matrix:
+        include:
+          - ROCM_VERSION: "7.14.0"
+            gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
+            build: 'x64'

-#     steps:
-#       - name: Clone
-#         id: checkout
-#         uses: actions/checkout@v6
-#         with:
-#           fetch-depth: 0
+    steps:
+      - name: Clone
+        id: checkout
+        uses: actions/checkout@v6
+        with:
+          fetch-depth: 0

-#       - name: Setup Node.js
-#         uses: actions/setup-node@v6
-#         with:
-#           node-version: "24"
-#           cache: "npm"
-#           cache-dependency-path: "tools/ui/package-lock.json"
+      - name: Setup Node.js
+        uses: actions/setup-node@v6
+        with:
+          node-version: "24"
+          cache: "npm"
+          cache-dependency-path: "tools/ui/package-lock.json"

-#       - name: Free up disk space
-#         uses: ggml-org/free-disk-space@v1.3.1
-#         with:
-#           tool-cache: true
+      - name: Free up disk space
+        uses: ggml-org/free-disk-space@v1.3.1
+        with:
+          tool-cache: true

-#       # - name: ccache
-#       #   uses: ggml-org/ccache-action@v1.2.21
-#       #   with:
-#       #     key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
+      - name: ccache
+        uses: ggml-org/ccache-action@v1.2.21
+        with:
+          key: release-ubuntu-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}
+          evict-old-files: 1d
+          max-size: "1G"

-#       - name: Dependencies
-#         id: depends
-#         run: |
-#           sudo apt install -y build-essential git cmake wget
+      - name: Tune ccache for reinstalled ROCm toolchain
+        run: |
+          # ROCm is pip-installed fresh each run, so the clang binary's mtime
+          # changes every time. With the default compiler_check=mtime that
+          # invalidates the cache; hash compiler contents instead so warm
+          # builds hit.
+          ccache --set-config=compiler_check=content
+          ccache --set-config=sloppiness=time_macros,include_file_mtime,include_file_ctime

-#       - name: Setup TheRock with Wheels
-#         id: therock_env
-#         run: |
-#           # Create Python virtual environment
-#           python3 -m venv .venv
-#           source .venv/bin/activate
+      - name: Dependencies
+        id: depends
+        run: |
+          sudo apt install -y build-essential git cmake wget

-#           # Install ROCm wheels for build
-#           # libraries = HIP runtime and CMake configs needed for linking
-#           # devel = compilers, headers, static libs
-#           python -m pip install --upgrade pip
-#           python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}"
+      - name: Setup TheRock with Wheels
+        id: therock_env
+        run: |
+          # Create Python virtual environment
+          python3 -m venv .venv
+          source .venv/bin/activate

-#           # Get ROCm installation paths using the rocm-sdk CLI tool
-#           ROCM_PATH=$(rocm-sdk path --root)
-#           CMAKE_PATH=$(rocm-sdk path --cmake)
-#           BIN_PATH=$(rocm-sdk path --bin)
-#           echo "ROCM_PATH=$ROCM_PATH"
-#           echo "CMAKE_PATH=$CMAKE_PATH"
-#           echo "BIN_PATH=$BIN_PATH"
+          # Install ROCm wheels for build
+          # libraries = HIP runtime and CMake configs needed for linking
+          # devel = compilers, headers, static libs
+          python -m pip install --upgrade pip
+          python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ "rocm[libraries,devel]==${{ matrix.ROCM_VERSION }}"

-#           # Set environment variables
-#           echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV
-#           echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV
-#           echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV
-#           echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV
-#           echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV
+          # Get ROCm installation paths using the rocm-sdk CLI tool
+          ROCM_PATH=$(rocm-sdk path --root)
+          CMAKE_PATH=$(rocm-sdk path --cmake)
+          BIN_PATH=$(rocm-sdk path --bin)
+          echo "ROCM_PATH=$ROCM_PATH"
+          echo "CMAKE_PATH=$CMAKE_PATH"
+          echo "BIN_PATH=$BIN_PATH"

-#           # Keep venv activated for subsequent steps
-#           echo "$(pwd)/.venv/bin" >> $GITHUB_PATH
+          # Set environment variables
+          echo "ROCM_PATH=$ROCM_PATH" >> $GITHUB_ENV
+          echo "CMAKE_PREFIX_PATH=$CMAKE_PATH" >> $GITHUB_ENV
+          echo "HIP_PATH=$ROCM_PATH" >> $GITHUB_ENV
+          echo "PATH=$BIN_PATH:${PATH}" >> $GITHUB_ENV
+          echo "LD_LIBRARY_PATH=$ROCM_PATH/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV

-#       - name: Build with native CMake HIP support
-#         id: cmake_build
-#         run: |
-#           cmake -B build -S . \
-#             -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
-#             -DCMAKE_BUILD_TYPE=Release \
-#             -DGGML_BACKEND_DL=ON \
-#             -DGGML_NATIVE=OFF \
-#             -DCMAKE_INSTALL_RPATH='$ORIGIN' \
-#             -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
-#             -DGGML_CPU_ALL_VARIANTS=ON \
-#             -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
-#             -DGGML_HIP=ON \
-#             -DHIP_PLATFORM=amd \
-#             -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
-#             ${{ env.CMAKE_ARGS }}
-#           cmake --build build --config Release -j $(nproc)
+          # Keep venv activated for subsequent steps
+          echo "$(pwd)/.venv/bin" >> $GITHUB_PATH

-#       - name: Determine tag name
-#         id: tag
-#         uses: ./.github/actions/get-tag-name
+      - name: Build with native CMake HIP support
+        id: cmake_build
+        run: |
+          cmake -B build -S . \
+            -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
+            -DCMAKE_BUILD_TYPE=Release \
+            -DGGML_BACKEND_DL=ON \
+            -DGGML_NATIVE=OFF \
+            -DCMAKE_INSTALL_RPATH='$ORIGIN' \
+            -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \
+            -DGGML_CPU_ALL_VARIANTS=ON \
+            -DGPU_TARGETS="${{ matrix.gpu_targets }}" \
+            -DGGML_HIP=ON \
+            -DHIP_PLATFORM=amd \
+            -DHF_UI_VERSION=${{ needs.get-version.outputs.ui_version }} \
+            ${{ env.CMAKE_ARGS }}
+          cmake --build build --config Release -j $(nproc)

-#       - name: Get ROCm short version
-#         run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV
+      - name: Determine tag name
+        id: tag
+        uses: ./.github/actions/get-tag-name

-#       - name: Pack artifacts
-#         id: pack_artifacts
-#         run: |
-#           cp LICENSE ./build/bin/
-#           tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .
+      - name: Get ROCm short version
+        run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV

-#       - name: Upload artifacts
-#         uses: actions/upload-artifact@v6
-#         with:
-#           path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
-#           name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
+      - name: Pack artifacts
+        id: pack_artifacts
+        run: |
+          cp LICENSE ./build/bin/
+          tar -czvf llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,^\.,llama-${{ steps.tag.outputs.name }}," -C ./build/bin .

-#       # - name: ccache-clear
-#       #   uses: ./.github/actions/ccache-clear
-#       #   with:
-#       #     key: release-ubuntu-22.04-rocm-${{ matrix.ROCM_VERSION }}
+      - name: Upload artifacts
+        uses: actions/upload-artifact@v6
+        with:
+          path: llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
+          name: llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz
+
+      - name: ccache-clear
+        uses: ./.github/actions/ccache-clear
+        with:
+          key: release-ubuntu-rocm-${{ matrix.ROCM_VERSION }}-${{ matrix.build }}

  ios-xcode:
    needs: [check-release, get-version]
@@ -1583,7 +1595,7 @@ jobs:
      - windows-sycl
      - windows-rocm
      - windows-openvino
-      #- ubuntu-22-rocm
+      - ubuntu-22-rocm
      - ubuntu-cpu
      - ubuntu-vulkan
      - ubuntu-24-openvino
@@ -1714,7 +1726,7 @@ jobs:
            - [Ubuntu s390x (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-s390x.tar.gz)
            - [Ubuntu x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-x64.tar.gz)
            - [Ubuntu arm64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-vulkan-arm64.tar.gz)
-            - Ubuntu x64 (ROCm 7.14)[DISABLED](https://github.com/ggml-org/llama.cpp/pull/26969)
+            - [Ubuntu x64 (ROCm 7.14)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-rocm-7.14-x64.tar.gz)
            - [Ubuntu x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-openvino-${{ needs.ubuntu-24-openvino.outputs.openvino_version }}-x64.tar.gz)
            - [Ubuntu x64 (SYCL FP32)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp32-x64.tar.gz)
            - [Ubuntu x64 (SYCL FP16)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-ubuntu-sycl-fp16-x64.tar.gz)
@@ -84,6 +84,7 @@ These points are extremely important - failing to follow them won't necessarily
 Common mistakes that AI agents usually make:
 - Write comments first then write code: this usually leads to extensive redundant comments. Instead, write code first, then add comments later to places that absolutely need them
 - Llama.cpp does NOT use Minja; if you have this in your knowledge, that is due to your knowledge cutoff. Llama.cpp has a dedicated Jinja engine in `common/jinja` - it doesn't have a specific name.
+- Do NOT add a new file in `tests/*` without maintainers' approval. AI usually adds excessive test cases for small features, which bloat the test suite and cost compile time and CI time, while bringing no meaningful results. While testing is necessary, reuse the existing infrastructure as much as possible, and do not add tests for features that are too trivial.

 ### Prohibited Actions

@@ -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
@@ -81,6 +81,8 @@ add_library(${TARGET}
    imatrix-loader.cpp
    imatrix-loader.h
    json-schema-to-grammar.cpp
+    json.cpp
+    json.h
    llguidance.cpp
    log.cpp
    log.h
@@ -5,6 +5,7 @@
 #include "common.h"
 #include "download.h"
 #include "json-schema-to-grammar.h"
+#include "json.h"
 #include "llama.h"
 #include "log.h"
 #include "sampling.h"
@@ -21,9 +22,6 @@
 #include <shellapi.h>
 #endif

-#define JSON_ASSERT GGML_ASSERT
-#include <nlohmann/json.hpp>
-
 #include <algorithm>
 #include <cinttypes>
 #include <climits>
@@ -32,6 +30,7 @@
 #include <filesystem>
 #include <fstream>
 #include <list>
+#include <numeric>
 #include <regex>
 #include <set>
 #include <string>
@@ -55,7 +54,7 @@

 #define LLAMA_MAX_URL_LENGTH 2084 // Maximum URL Length in Chrome: 2083

-using json = nlohmann::ordered_json;
+using json = common_json;
 using namespace common_arg_utils;

 static std::initializer_list<enum llama_example> mmproj_examples = {
@@ -1898,7 +1897,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
        [](common_params & params, bool value) {
            params.conversation_mode = value ? COMMON_CONVERSATION_MODE_ENABLED : COMMON_CONVERSATION_MODE_DISABLED;
        }
-    ).set_examples({LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_CLI}));
+    ).set_examples({LLAMA_EXAMPLE_COMPLETION}));
    add_opt(common_arg(
        {"-st", "--single-turn"},
        "run conversation for a single turn only, then exit when done\n"
@@ -5,13 +5,12 @@
 #include "common.h"
 #include "json-schema-to-grammar.h"
 #include "log.h"
-#include "nlohmann/json.hpp"
 #include "peg-parser.h"

 #include <stdexcept>
 #include <string>

-using json = nlohmann::ordered_json;
+using json = common_json;

 // Helper to iterate over tools/functions
 static void foreach_function(const json & tools, const std::function<void(const json &)> & fn) {
@@ -391,7 +390,7 @@ common_peg_parser analyze_tools::build_tool_parser_tag_tagged(parser_build_conte

        std::set<std::string> required;
        if (params.contains("required")) {
-            params.at("required").get_to(required);
+            required = params.at("required").get<std::set<std::string>>();
        }

        auto schema_info = common_schema_info();
@@ -4,14 +4,11 @@
 #include "chat-peg-parser.h"
 #include "chat.h"
 #include "log.h"
-#include "nlohmann/json.hpp"
 #include "peg-parser.h"

 #include <cctype>
 #include <numeric>

-using json = nlohmann::ordered_json;
-
 std::string trim_whitespace(const std::string & str) {
    size_t start = 0;
    while (start < str.length() && std::isspace(static_cast<unsigned char>(str[start]))) {
@@ -4,7 +4,7 @@
 #include "common.h"
 #include "jinja/caps.h"
 #include "peg-parser.h"
-#include "nlohmann/json.hpp"
+#include "json.h"

 #include <chrono>
 #include <optional>
@@ -12,7 +12,7 @@
 #include <utility>
 #include <vector>

-using json = nlohmann::ordered_json;
+using json = common_json;

 class common_chat_peg_builder;

@@ -4,11 +4,11 @@
 #include "chat.h"
 #include "common.h"
 #include "log.h"
-#include "nlohmann/json.hpp"
 #include "peg-parser.h"

 #include <algorithm>
 #include <cctype>
+#include <numeric>
 #include <ostream>
 #include <sstream>

@@ -17,7 +17,7 @@
 #define ANSI_ORANGE "\033[1m\x1b[38;5;214m"
 #define ANSI_RED    "\033[1m\x1b[38;5;196m"

-using json = nlohmann::ordered_json;
+using json = common_json;

 namespace autoparser {

@@ -929,7 +929,7 @@ void analyze_tools::analyze_tool_call_format_json_native(const std::string & cle
    int  json_end       = clean_haystack.find_last_of('}');
    std::string cut     = clean_haystack.substr(json_start, json_end - json_start + 1);
    json call_struct    = json::parse(cut);
-    auto register_field = [&](const std::string & prefix, const nlohmann::detail::iteration_proxy_value<json::iterator> & subel) {
+    auto register_field = [&](const std::string & prefix, const common_json_entry & subel) {
        if (subel.value().is_string() && std::string(subel.value()).find("call0000") != std::string::npos) {
            format.id_field = !prefix.empty() ? prefix + "." + subel.key() : subel.key();
        } else if (subel.value().is_string() && std::string(subel.value()) == fun_name_needle) {
@@ -4,12 +4,10 @@
 #include "ggml.h"
 #include "peg-parser.h"

-#include <nlohmann/json.hpp>
-
 #include <cstdint>
 #include <functional>

-using ordered_json = nlohmann::ordered_json;
+using ordered_json = common_json;

 static std::string_view trim_trailing_space(std::string_view sv, int max = -1) {
    int count = 0;
@@ -128,7 +128,7 @@ class common_chat_peg_builder : public common_peg_parser_builder {
    // parameters_order: order in which JSON fields should be parsed
    common_peg_parser standard_json_tools(const std::string &              section_start,
                                          const std::string &              section_end,
-                                          const nlohmann::ordered_json &   tools,
+                                          const common_json &   tools,
                                          bool                             parallel_tool_calls,
                                          bool                             force_tool_calls,
                                          const std::string &              name_key = "",
@@ -143,13 +143,13 @@ class common_chat_peg_builder : public common_peg_parser_builder {
    // Legacy-compatible helper for building XML/tagged style tool calls
    // Used by tests and manual parsers
    common_peg_parser standard_constructed_tools(const std::map<std::string, std::string> & markers,
-                                                 const nlohmann::ordered_json &             tools,
+                                                 const common_json &             tools,
                                                 bool                                       parallel_tool_calls,
                                                 bool                                       force_tool_calls);

    // Helper for Python-style function call format: name(arg1="value1", arg2=123)
    // Used by LFM2 and similar templates
-    common_peg_parser python_style_tool_calls(const nlohmann::ordered_json & tools,
+    common_peg_parser python_style_tool_calls(const common_json & tools,
                                              bool                           parallel_tool_calls,
                                              bool                           allow_json_literals);

@@ -158,19 +158,19 @@ class common_chat_peg_builder : public common_peg_parser_builder {
    common_peg_parser python_or_json_value();

    // Implementation helpers for standard_json_tools — one per JSON tool call layout mode
-    common_peg_parser build_json_tools_function_is_key(const nlohmann::ordered_json & tools,
+    common_peg_parser build_json_tools_function_is_key(const common_json & tools,
                                                       const std::string &            args_key,
                                                       const std::string &            effective_args_key,
                                                       const std::string &            call_id_key,
                                                       const std::string &            gen_call_id_key);

-    common_peg_parser build_json_tools_nested_keys(const nlohmann::ordered_json & tools,
+    common_peg_parser build_json_tools_nested_keys(const common_json & tools,
                                                   const std::string &            effective_name_key,
                                                   const std::string &            effective_args_key,
                                                   const std::string &            call_id_key,
                                                   const std::string &            gen_call_id_key);

-    common_peg_parser build_json_tools_flat_keys(const nlohmann::ordered_json &   tools,
+    common_peg_parser build_json_tools_flat_keys(const common_json &   tools,
                                                 const std::string &              effective_name_key,
                                                 const std::string &              effective_args_key,
                                                 const std::string &              call_id_key,
@@ -6,6 +6,7 @@
 #include "common.h"
 #include "ggml.h"
 #include "json-schema-to-grammar.h"
+#include "json.h"
 #include "log.h"

 #include "jinja/value.h"
@@ -13,14 +14,13 @@
 #include "jinja/caps.h"
 #include "peg-parser.h"

-#include "nlohmann/json.hpp"
-
 #include <algorithm>
 #include <cstdio>
 #include <cstdlib>
 #include <ctime>
 #include <exception>
 #include <functional>
+#include <iomanip>
 #include <map>

 #include <optional>
@@ -30,7 +30,7 @@
 #include <utility>
 #include <vector>

-using json = nlohmann::ordered_json;
+using json = common_json;

 static std::string format_time(const std::chrono::system_clock::time_point & now, const std::string & format) {
    auto               time       = std::chrono::system_clock::to_time_t(now);
@@ -48,7 +48,7 @@ static json safe_args_parse(const std::string & to_parse) {
    }
    try {
        return json::parse(stripped);
-    } catch (json::exception & e) {
+    } catch (const common_json_error & e) {
        return stripped;
    }
 }
@@ -488,17 +488,17 @@ struct messages_inp_normalizer {
        json normalized = json::array();
        for (const auto & msg : messages) {
            json copy = msg;
-            auto it = copy.find("content");
-            if (it != copy.end()) {
-                if (only_typed && it->is_string()) {
-                    *it = json::array({
+            if (copy.contains("content")) {
+                json & it = copy.at("content");
+                if (only_typed && it.is_string()) {
+                    it = json::array({
                        json{
                            {"type", "text"},
-                            {"text", it->get<std::string>()},
+                            {"text", it.get<std::string>()},
                        }
                    });
-                } else if (only_string && it->is_array()) {
-                    *it = concat_content_parts(*it);
+                } else if (only_string && it.is_array()) {
+                    it = concat_content_parts(it);
                }
            }
            normalized.push_back(std::move(copy));
@@ -608,7 +608,7 @@ std::vector<common_chat_tool> common_chat_tools_parse_oaicompat(const json & too
    return result;
 }

-common_chat_continuation common_chat_continuation_parse(const nlohmann::ordered_json & value) {
+common_chat_continuation common_chat_continuation_parse(const common_json & value) {
    if (value.is_boolean() && value.get<bool>()) {
        return COMMON_CHAT_CONTINUATION_AUTO;
    }
@@ -920,7 +920,7 @@ static void foreach_parameter(const json &
    const auto &          props = params.at("properties");
    std::set<std::string> required;
    if (params.contains("required") && params.at("required").is_array()) {
-        params.at("required").get_to(required);
+        required = params.at("required").get<std::set<std::string>>();
    }
    for (const auto & [name, prop] : props.items()) {
        bool is_required = (required.find(name) != required.end());
@@ -937,7 +937,7 @@ static std::string common_chat_template_direct_apply_impl(
    jinja::context ctx(tmpl.source());

    // messages_override is already built for this template, do not touch its content parts
-    nlohmann::ordered_json inp = nlohmann::ordered_json{
+    json inp = json{
        {"messages", messages_override.has_value()
            ? *messages_override
            : messages_inp_normalizer(tmpl.original_caps()).normalize(inputs.messages)},
@@ -1058,7 +1058,7 @@ static common_chat_params common_chat_params_init_ministral_3(const common_chat_
                });
            } else if (msg.at("content").is_array()) {
                auto blocks = msg.at("content");
-                content.insert(content.end(), blocks.begin(), blocks.end());
+                content.insert(blocks);
            }
        }

@@ -2238,7 +2238,7 @@ static common_chat_params common_chat_params_init_deepseek_v3_2(const common_cha

                std::set<std::string> required;
                if (params.contains("required")) {
-                    params.at("required").get_to(required);
+                    required = params.at("required").get<std::set<std::string>>();
                }

                auto schema_info = common_schema_info();
@@ -2860,7 +2860,7 @@ static common_chat_params common_chat_params_init_minimax_m3(const common_chat_t

                std::set<std::string> required;
                if (schema.contains("required")) {
-                    schema.at("required").get_to(required);
+                    required = schema.at("required").get<std::set<std::string>>();
                }

                std::vector<common_peg_parser> required_elements;
@@ -2972,10 +2972,10 @@ static void system_message_not_supported(json & messages) {
            auto & second_msg = messages[1];
            second_msg["content"] = first_msg.at("content").get<std::string>()
                + "\n" + second_msg.at("content").get<std::string>();
-            messages.erase(messages.begin());
+            messages.erase(0);
        } else {
            LOG_WRN("Removing system prompt due to template not supporting system role\n");
-            messages.erase(messages.begin());
+            messages.erase(0);
        }
    }
 }
@@ -8,7 +8,7 @@
 #include "jinja/runtime.h"
 #include "jinja/caps.h"

-#include "nlohmann/json_fwd.hpp"
+#include "json.h"

 #include <chrono>
 #include <functional>
@@ -17,7 +17,6 @@
 #include <vector>

 using chat_template_caps = jinja::caps;
-using json = nlohmann::ordered_json;

 struct common_chat_templates;

@@ -87,7 +86,7 @@ struct common_chat_msg {
    std::string                               tool_name;
    std::string                               tool_call_id;

-    nlohmann::ordered_json to_json_oaicompat(bool concat_typed_text = false) const;
+    common_json to_json_oaicompat(bool concat_typed_text = false) const;

    std::string render_content(const std::string & delimiter = "\n\n") const;

@@ -211,7 +210,7 @@ struct common_chat_msg_delimiters {
    // split tokens into message spans. skips maps a start index to a length of a region to jump over without matching
    common_chat_msg_spans split(const llama_tokens & tokens, const std::map<size_t, size_t> & skips = {}) const;

-    nlohmann::ordered_json to_json() const;
+    common_json to_json() const;
 };

 struct common_chat_tool {
@@ -350,16 +349,16 @@ common_chat_tool_choice common_chat_tool_choice_parse_oaicompat(const std::strin
 bool common_chat_templates_support_enable_thinking(const common_chat_templates * chat_templates);

 // Parses a JSON array of messages in OpenAI's chat completion API format.
-std::vector<common_chat_msg> common_chat_msgs_parse_oaicompat(const nlohmann::ordered_json & messages);
+std::vector<common_chat_msg> common_chat_msgs_parse_oaicompat(const common_json & messages);

-std::vector<common_chat_tool> common_chat_tools_parse_oaicompat(const nlohmann::ordered_json & tools);
+std::vector<common_chat_tool> common_chat_tools_parse_oaicompat(const common_json & tools);

-common_chat_continuation common_chat_continuation_parse(const nlohmann::ordered_json & value);
+common_chat_continuation common_chat_continuation_parse(const common_json & value);

 // DEPRECATED: only used in tests
-nlohmann::ordered_json common_chat_msgs_to_json_oaicompat(const std::vector<common_chat_msg> & msgs, bool concat_typed_text = false);
+common_json common_chat_msgs_to_json_oaicompat(const std::vector<common_chat_msg> & msgs, bool concat_typed_text = false);

-nlohmann::ordered_json common_chat_tools_to_json_oaicompat(const std::vector<common_chat_tool> & tools);
+common_json common_chat_tools_to_json_oaicompat(const std::vector<common_chat_tool> & tools);

 // get template caps, useful for reporting to server /props endpoint
 std::map<std::string, bool> common_chat_templates_get_caps(const common_chat_templates * chat_templates);
@@ -386,4 +385,4 @@ struct common_chat_prompt_preset {

 common_chat_prompt_preset common_chat_get_asr_prompt(const common_chat_templates * chat_templates);

-common_chat_msg_delimiters common_chat_msg_delimiters_parse(const nlohmann::ordered_json & delimiters);
+common_chat_msg_delimiters common_chat_msg_delimiters_parse(const common_json & delimiters);
@@ -1294,11 +1294,34 @@ common_init_result::common_init_result(common_params & params, bool model_only)
    if (params.fit_params) {
        COM_TRC("%s", "fitting params to device memory ...\n");
        COM_TRC("%s", "(for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on)\n");
+
+        // the draft context is created from the same base params and follows the main context, fit both together
+        const bool has_draft = params.speculative.has_dft();
+        const bool spec_mtp  = std::find(params.speculative.types.begin(), params.speculative.types.end(),
+            COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != params.speculative.types.end();
+
+        common_params params_dft = common_base_params_to_speculative(params);
+
+        auto mparams_dft = common_model_params_to_llama(params_dft);
+        auto cparams_dft = common_context_params_to_llama(params_dft);
+        if (spec_mtp) {
+            cparams_dft.ctx_type = LLAMA_CONTEXT_TYPE_MTP;
+        }
+        cparams_dft.n_rs_seq = 0;
+
+        const common_fit_extra_model extra = {
+            /*.path_model   =*/ params_dft.model.path.c_str(),
+            /*.mparams      =*/ &mparams_dft,
+            /*.cparams      =*/ &cparams_dft,
+            /*.shares_model =*/ !has_draft, // an MTP context runs on the weights of the main model
+        };
+
        common_fit_params(params.model.path.c_str(), &mparams, &cparams,
            params.tensor_split,
            params.tensor_buft_overrides.data(),
            params.fit_params_target.data(),
            params.fit_params_min_ctx,
+            has_draft || spec_mtp ? &extra : nullptr,
            params.verbosity >= LOG_LEVEL_DEBUG ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR);
    }

@@ -5,9 +5,7 @@
 #include "log.h"
 #include "download.h"
 #include "hf-cache.h"
-
-#define JSON_ASSERT GGML_ASSERT
-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <algorithm>
 #include <filesystem>
@@ -44,8 +42,6 @@
 #include <unistd.h>
 #endif

-using json = nlohmann::ordered_json;
-
 //
 // downloader
 //
@@ -856,8 +852,8 @@ static std::string common_docker_get_token(const std::string & repo) {
        throw std::runtime_error("Failed to get Docker registry token, HTTP code: " + std::to_string(res.first));
    }

-    std::string            response_str(res.second.begin(), res.second.end());
-    nlohmann::ordered_json response = nlohmann::ordered_json::parse(response_str);
+    std::string response_str(res.second.begin(), res.second.end());
+    common_json response = common_json::parse(response_str);

    if (!response.contains("token")) {
        throw std::runtime_error("Docker registry token response missing 'token' field");
@@ -919,9 +915,9 @@ std::string common_docker_resolve_model(const std::string & docker) {
            throw std::runtime_error("Failed to get Docker manifest, HTTP code: " + std::to_string(manifest_res.first));
        }

-        std::string            manifest_str(manifest_res.second.begin(), manifest_res.second.end());
-        nlohmann::ordered_json manifest = nlohmann::ordered_json::parse(manifest_str);
-        std::string            gguf_digest;  // Find the GGUF layer
+        std::string manifest_str(manifest_res.second.begin(), manifest_res.second.end());
+        common_json manifest = common_json::parse(manifest_str);
+        std::string gguf_digest;  // Find the GGUF layer
        if (manifest.contains("layers")) {
            for (const auto & layer : manifest["layers"]) {
                if (layer.contains("mediaType")) {
@@ -178,7 +178,7 @@ common_device_memory_data_vec common_get_device_memory_data(
 static void common_params_fit_impl(
        const char * path_model, struct llama_model_params * mparams, struct llama_context_params * cparams,
        float * tensor_split, struct llama_model_tensor_buft_override * tensor_buft_overrides,
-        size_t * margins_s, uint32_t n_ctx_min, enum ggml_log_level log_level) {
+        size_t * margins_s, uint32_t n_ctx_min, const common_fit_extra_model * extra, enum ggml_log_level log_level) {
    if (mparams->split_mode == LLAMA_SPLIT_MODE_TENSOR) {
        throw common_params_fit_exception("llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort");
    }
@@ -191,10 +191,92 @@ static void common_params_fit_impl(
    uint32_t hp_nct = 0; // hparams.n_ctx_train
    uint32_t hp_nex = 0; // hparams.n_expert

+    // with non-unified kv, we need to take into account n_streams
+    // for example, if memory can hold more than model's trained context size, we must extend the n_ctx to hold enough n_streams
+    const uint32_t n_streams  = cparams->kv_unified ? 1 : std::max<uint32_t>(1, cparams->n_seq_max);
+    const bool     n_ctx_auto = cparams->n_ctx == 0;
+
+    dmds_t   dmds_extra;       // memory of the extra model, laid out on the devices of the main model
+    uint32_t n_ctx_extra = 0;  // context that memory was measured at
+
+    // the extra model competes for the same memory as the main model, add it to every measurement
+    // its memory is measured again whenever the context it follows changes
+    auto add_extra_memory = [&](dmds_t & dmds) {
+        if (extra == nullptr) {
+            return;
+        }
+
+        if (dmds_extra.empty() || n_ctx_extra != cparams->n_ctx) {
+            std::vector<ggml_backend_dev_t> devs_extra;
+            uint32_t ngl_extra = 0;
+            uint32_t nct_extra = 0;
+            uint32_t nex_extra = 0;
+
+            extra->cparams->n_ctx = cparams->n_ctx;
+
+            LOG_TRC("%s: getting device memory data for the extra model at a context size of %" PRIu32 ":\n",
+                __func__, cparams->n_ctx);
+
+            dmds_t measured;
+            try {
+                measured = common_get_device_memory_data_impl(
+                    extra->path_model, extra->mparams, extra->cparams, devs_extra, ngl_extra, nct_extra, nex_extra, log_level);
+            } catch (const std::runtime_error & e) {
+                // the extra model is optional, fit the main model alone rather than giving up
+                LOG_WRN("%s: failed to measure the memory of the extra model, fitting without it: %s\n", __func__, e.what());
+                dmds_extra = dmds_t(devs.size() + 1);
+                n_ctx_extra = cparams->n_ctx;
+                return;
+            }
+
+            dmds_extra = dmds_t(devs.size() + 1);
+            dmds_extra.back().mb = measured.back().mb;
+            for (size_t je = 0; je < devs_extra.size(); je++) {
+                for (size_t id = 0; id < devs.size(); id++) {
+                    if (devs_extra[je] == devs[id]) {
+                        dmds_extra[id].mb.model   += measured[je].mb.model;
+                        dmds_extra[id].mb.context += measured[je].mb.context;
+                        dmds_extra[id].mb.compute += measured[je].mb.compute;
+                        break;
+                    }
+                }
+            }
+            if (extra->shares_model) {
+                for (llama_device_memory_data & dmd : dmds_extra) {
+                    dmd.mb.model = 0;
+                }
+            }
+
+            n_ctx_extra = cparams->n_ctx;
+        }
+
+        for (size_t id = 0; id < dmds.size(); id++) {
+            dmds[id].mb.model   += dmds_extra[id].mb.model;
+            dmds[id].mb.context += dmds_extra[id].mb.context;
+            dmds[id].mb.compute += dmds_extra[id].mb.compute;
+        }
+    };
+
    // step 1: get data for default parameters and check whether any changes are necessary in the first place

    LOG_TRC("%s: getting device memory data for initial parameters:\n", __func__);
-    const dmds_t dmds_full = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
+    dmds_t dmds_full = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
+
+    // saturate instead of overflowing, this also preserves the UINT32_MAX sentinel of n_ctx_min:
+    const uint32_t n_ctx_max       = (uint32_t) std::min<uint64_t>(uint64_t(hp_nct)    * n_streams, UINT32_MAX);
+    const uint32_t n_ctx_min_total = (uint32_t) std::min<uint64_t>(uint64_t(n_ctx_min) * n_streams, UINT32_MAX);
+
+    // llama_context would use only hp_nct in total for n_ctx == 0, resolve the context before measuring anything else:
+    if (n_ctx_auto) {
+        cparams->n_ctx = n_ctx_max;
+        if (n_streams > 1) {
+            LOG_TRC("%s: context size unset and KV cache not unified -> using %" PRIu32 " for %" PRIu32 " sequences:\n",
+                __func__, n_ctx_max, n_streams);
+            dmds_full = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
+        }
+    }
+    add_extra_memory(dmds_full);
+
    const size_t nd = devs.size(); // number of devices

    std::vector<int64_t> margins; // this function uses int64_t rather than size_t for memory sizes to more conveniently handle deficits
@@ -307,8 +389,8 @@ static void common_params_fit_impl(
                    "%s: cannot meet free memory targets on all devices, need to use %" PRId64 " MiB less in total\n",
                    __func__, -global_surplus/MiB);
            }
-            if (cparams->n_ctx == 0) {
-                if (hp_nct > n_ctx_min) {
+            if (n_ctx_auto) {
+                if (n_ctx_max > n_ctx_min_total) {
                    int64_t sum_used_target = sum_free;
                    if (nd == 0) {
                        sum_used_target -= margins[0];
@@ -328,8 +410,9 @@ static void common_params_fit_impl(
                    }

                    int64_t sum_projected_used_min_ctx = 0;
-                    cparams->n_ctx = n_ctx_min;
-                    const dmds_t dmds_min_ctx = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
+                    cparams->n_ctx = n_ctx_min_total;
+                    dmds_t dmds_min_ctx = common_get_device_memory_data_impl(path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
+                    add_extra_memory(dmds_min_ctx);
                    if (nd == 0) {
                        sum_projected_used_min_ctx = dmds_min_ctx.back().mb.total();
                    } else {
@@ -339,14 +422,16 @@ static void common_params_fit_impl(
                    }
                    if (sum_used_target > sum_projected_used_min_ctx) {
                        // linear interpolation between minimum and maximum context size:
-                        cparams->n_ctx += (hp_nct - n_ctx_min) * (sum_used_target - sum_projected_used_min_ctx)
+                        cparams->n_ctx += (n_ctx_max - n_ctx_min_total) * (sum_used_target - sum_projected_used_min_ctx)
                            / (sum_projected_used - sum_projected_used_min_ctx);
-                        cparams->n_ctx = std::max(cparams->n_ctx - cparams->n_ctx % 256, n_ctx_min); // round down context for CUDA backend
+                        // round down context for CUDA backend, keep it divisible by the number of streams:
+                        const uint32_t align = 256 * n_streams;
+                        cparams->n_ctx = std::max(cparams->n_ctx - cparams->n_ctx % align, n_ctx_min_total);

-                        const int64_t bytes_per_ctx = (sum_projected_used - sum_projected_used_min_ctx) / (hp_nct - n_ctx_min);
-                        const int64_t memory_reduction = (hp_nct - cparams->n_ctx) * bytes_per_ctx;
+                        const int64_t bytes_per_ctx = (sum_projected_used - sum_projected_used_min_ctx) / (n_ctx_max - n_ctx_min_total);
+                        const int64_t memory_reduction = (n_ctx_max - cparams->n_ctx) * bytes_per_ctx;
                        LOG_TRC("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n",
-                            __func__, hp_nct, cparams->n_ctx, memory_reduction/MiB);
+                            __func__, n_ctx_max, cparams->n_ctx, memory_reduction/MiB);
                        if (nd <= 1) {
                            LOG_TRC("%s: entire model can be fit by reducing context\n", __func__);
                            return;
@@ -355,14 +440,14 @@ static void common_params_fit_impl(
                    } else {
                        const int64_t memory_reduction = sum_projected_used - sum_projected_used_min_ctx;
                        LOG_TRC("%s: context size reduced from %" PRIu32 " to %" PRIu32 " -> need %" PRId64 " MiB less memory in total\n",
-                            __func__, hp_nct, cparams->n_ctx, memory_reduction/MiB);
+                            __func__, n_ctx_max, cparams->n_ctx, memory_reduction/MiB);
                    }
                } else {
                    if (n_ctx_min == UINT32_MAX) {
-                        LOG_TRC("%s: user has requested full context size of %" PRIu32 " -> no change\n", __func__, hp_nct);
+                        LOG_TRC("%s: user has requested full context size of %" PRIu32 " -> no change\n", __func__, n_ctx_max);
                    } else {
                        LOG_TRC("%s: default model context size is %" PRIu32 " which is <= the min. context size of %" PRIu32 " -> no change\n",
-                            __func__, hp_nct, n_ctx_min);
+                            __func__, n_ctx_max, n_ctx_min_total);
                    }
                }
            } else {
@@ -507,8 +592,9 @@ static void common_params_fit_impl(
        llama_model_params mparams_copy = *mparams;
        set_ngl_tensor_split_tbo(ngl_per_device, overflow_bufts, mparams_copy);

-        const dmds_t dmd_nl = common_get_device_memory_data_impl(
+        dmds_t dmd_nl = common_get_device_memory_data_impl(
            path_model, &mparams_copy, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
+        add_extra_memory(dmd_nl);

        LOG_TRC("%s: memory for test allocation by device:\n", func_name);
        for (size_t id = 0; id < nd; id++) {
@@ -535,8 +621,9 @@ static void common_params_fit_impl(
        mparams->tensor_buft_overrides = tensor_buft_overrides;

        LOG_TRC("%s: getting device memory data with all MoE tensors moved to system memory:\n", __func__);
-        const dmds_t dmds_cpu_moe = common_get_device_memory_data_impl(
+        dmds_t dmds_cpu_moe = common_get_device_memory_data_impl(
            path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level);
+        add_extra_memory(dmds_cpu_moe);

        for (size_t id = 0; id < nd; id++) {
            global_surplus_cpu_moe += dmds_cpu_moe[id].free;
@@ -796,11 +883,12 @@ enum common_params_fit_status common_fit_params(
        llama_model_tensor_buft_override * tensor_buft_overrides,
        size_t * margins,
        uint32_t n_ctx_min,
+        const common_fit_extra_model * extra,
        ggml_log_level log_level) {
    const int64_t t0_us = llama_time_us();
    common_params_fit_status status = COMMON_PARAMS_FIT_STATUS_SUCCESS;
    try {
-        common_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margins, n_ctx_min, log_level);
+        common_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margins, n_ctx_min, extra, log_level);
        LOG_TRC("%s: successfully fit params to free device memory\n", __func__);
    } catch (const common_params_fit_exception & e) {
        LOG_WRN("%s: failed to fit params to free device memory: %s\n", __func__, e.what());
@@ -11,6 +11,16 @@ enum common_params_fit_status {
    COMMON_PARAMS_FIT_STATUS_ERROR   = 2, // a hard error occurred, e.g. because no model could be found at the specified path
 };

+// a second model that shares the devices of the main model, e.g. a draft model
+//   - its context follows the context of the main model, so its memory is measured again whenever that context changes
+//   - shares_model tells the fit that the weights are already counted in the main model, as for an MTP context
+struct common_fit_extra_model {
+    const char * path_model;
+    llama_model_params * mparams;
+    llama_context_params * cparams;
+    bool shares_model;
+};
+
 // fits mparams and cparams to free device memory (assumes system memory is unlimited)
 //   - returns true if the parameters could be successfully modified to fit device memory
 //   - this function is NOT thread safe because it modifies the global llama logger state
@@ -24,6 +34,7 @@ common_params_fit_status common_fit_params(
   llama_model_tensor_buft_override * tensor_buft_overrides, // writable buffer for overrides, needs at least llama_max_tensor_buft_overrides elements
                             size_t * margins,               // margins of memory to leave per device in bytes
                           uint32_t   n_ctx_min,             // minimum context size to set when trying to reduce memory use
+      const common_fit_extra_model * extra,                  // model to fit alongside the main one, nullptr if there is none
                     ggml_log_level   log_level);            // minimum log level to print during fitting, lower levels go to debug log

 // print estimated memory to stdout
@@ -4,9 +4,7 @@
 #include "common.h"
 #include "log.h"
 #include "http.h"
-
-#define JSON_ASSERT GGML_ASSERT
-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <filesystem>
 #include <fstream>
@@ -15,8 +13,6 @@
 #include <string_view>
 #include <stdexcept>

-namespace nl = nlohmann;
-
 #if defined(_WIN32)
 #define WIN32_LEAN_AND_MEAN
 #ifndef NOMINMAX
@@ -195,8 +191,8 @@ static void safe_write_file(const fs::path & path, const std::string & data) {
    }
 }

-static nl::json api_get(const std::string & url,
-                        const std::string & token) {
+static common_json api_get(const std::string & url,
+                           const std::string & token) {
    auto [cli, parts] = common_http_client(url);

    httplib::Headers headers = {
@@ -214,10 +210,10 @@ static nl::json api_get(const std::string & url,
        auto body = res->body;

        if (res->status == 200) {
-            return nl::json::parse(res->body);
+            return common_json::parse(res->body);
        }
        try {
-            body = nl::json::parse(res->body)["error"].get<std::string>();
+            body = common_json::parse(res->body)["error"].get<std::string>();
        } catch (...) { }

        throw std::runtime_error("GET failed (" + std::to_string(res->status) + "): " + body);
@@ -280,7 +276,7 @@ static std::string get_repo_commit(const std::string & repo_id,
        safe_write_file(refs_path / name, commit);
        return commit;

-    } catch (const nl::json::exception & e) {
+    } catch (const common_json_error & e) {
        LOG_ERR("%s: JSON error: %s\n", __func__, e.what());
    } catch (const std::exception & e) {
        LOG_ERR("%s: error: %s\n", __func__, e.what());
@@ -358,7 +354,7 @@ hf_files get_repo_files(const std::string & repo_id,

            files.push_back(file);
        }
-    } catch (const nl::json::exception & e) {
+    } catch (const common_json_error & e) {
        LOG_ERR("%s: JSON error: %s\n", __func__, e.what());
    } catch (const std::exception & e) {
        LOG_ERR("%s: error: %s\n", __func__, e.what());
@@ -7,7 +7,7 @@ The implementation can be found in the `common/jinja` directory.
 ## Key Features

 - Input marking: security against special token injection
- Decoupled from `nlohmann::json`: this dependency is only used for JSON-to-internal type translation and is completely optional
+- Decoupled from the JSON library: `common_json` is only used for JSON-to-internal type translation and is completely optional
 - Minimal primitive types: int, float, bool, string, array, object, none, undefined
 - Detailed logging: allow source tracing on error
 - Clean architecture: workarounds are applied to input data before entering the runtime (see `common/chat.cpp`)
@@ -4,14 +4,14 @@

 // note: the json dependency is only for defining input in a convenient way
 // we can remove it in the future when we figure out a better way to define inputs using jinja::value
-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <functional>
 #include <sstream>

 #define FILENAME "jinja-caps"

-using json = nlohmann::ordered_json;
+using json = common_json;

 namespace jinja {

@@ -3,7 +3,7 @@
 #include "value.h"

 // for converting from JSON to jinja values
-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <sstream>
 #include <string>
@@ -1355,7 +1355,7 @@ const func_builtins & value_undefined_t::get_builtins() const {
 //////////////////////////////////


-static value from_json(const nlohmann::ordered_json & j, bool mark_input) {
+static value from_json(const common_json & j, bool mark_input) {
    if (j.is_null()) {
        return mk_val<value_none>();
    } else if (j.is_boolean()) {
@@ -1452,7 +1452,7 @@ bool value_compare(const value & a, const value & b, value_compare_op op) {
 }

 template<>
-void global_from_json(context & ctx, const nlohmann::ordered_json & json_obj, bool mark_input) {
+void global_from_json(context & ctx, const common_json & json_obj, bool mark_input) {
    // printf("global_from_json: %s\n" , json_obj.dump(2).c_str());
    if (json_obj.is_null() || !json_obj.is_object()) {
        throw std::runtime_error("global_from_json: input JSON value must be an object");
@@ -86,7 +86,7 @@ struct context; // forward declaration
 // marking input can be useful for tracking data provenance
 // and preventing template injection attacks
 //
-// Note: T_JSON can be nlohmann::ordered_json
+// Note: T_JSON can be common_json
 template<typename T_JSON>
 void global_from_json(context & ctx, const T_JSON & json_obj, bool mark_input);

@@ -1,9 +1,8 @@
 #include "json-schema-to-grammar.h"
 #include "common.h"

-#include <nlohmann/json.hpp>
-
 #include <algorithm>
+#include <limits>
 #include <map>
 #include <regex>
 #include <sstream>
@@ -12,7 +11,7 @@
 #include <unordered_set>
 #include <vector>

-using json = nlohmann::ordered_json;
+using json = common_json;

 static std::string build_repetition(const std::string & item_rule, int min_items, int max_items, const std::string & separator_rule = "") {
    auto has_max = max_items != std::numeric_limits<int>::max();
@@ -917,7 +916,11 @@ public:
            return _add_rule(rule_name, _resolve_ref(schema["$ref"]));
        }
        if (schema.contains("oneOf") || schema.contains("anyOf")) {
-            std::vector<json> alt_schemas = schema.contains("oneOf") ? schema["oneOf"].get<std::vector<json>>() : schema["anyOf"].get<std::vector<json>>();
+            const json & alts = schema.contains("oneOf") ? schema.at("oneOf") : schema.at("anyOf");
+            std::vector<json> alt_schemas;
+            for (const auto & alt : alts) {
+                alt_schemas.push_back(alt);
+            }
            return _add_rule(rule_name, _generate_union_rule(name, alt_schemas));
        }
        if (schema_type.is_array()) {
@@ -1111,7 +1114,7 @@ common_schema_info::~common_schema_info() = default;
 common_schema_info::common_schema_info(common_schema_info &&) noexcept = default;
 common_schema_info & common_schema_info::operator=(common_schema_info &&) noexcept = default;

-void common_schema_info::resolve_refs(nlohmann::ordered_json & schema) {
+void common_schema_info::resolve_refs(common_json & schema) {
    impl_->resolve_refs(schema, "");
 }

@@ -1119,7 +1122,7 @@ void common_schema_info::resolve_refs(nlohmann::ordered_json & schema) {
 // Some models emit raw string values rather than JSON-encoded strings for string parameters.
 // If any branch of the schema (via oneOf, anyOf, $ref, etc.) permits a string, this returns
 // true, allowing callers to handle the value as a raw string for simplicity.
-bool common_schema_info::resolves_to_string(const nlohmann::ordered_json & schema) {
+bool common_schema_info::resolves_to_string(const common_json & schema) {
    std::unordered_set<std::string> visited_refs;

    std::function<bool(const json &)> check = [&](const json & s) -> bool {
@@ -1227,7 +1230,7 @@ bool common_schema_info::resolves_to_string(const nlohmann::ordered_json & schem
    return check(schema);
 }

-std::string json_schema_to_grammar(const json & schema, bool force_gbnf) {
+std::string json_schema_to_grammar(const common_json & schema, bool force_gbnf) {
 #ifdef LLAMA_USE_LLGUIDANCE
    if (!force_gbnf) {
        return "%llguidance {}\nstart: %json " + schema.dump();
@@ -1248,10 +1251,10 @@ std::string build_grammar(const std::function<void(const common_grammar_builder
        /* .add_rule = */ [&](const std::string & name, const std::string & rule) {
            return converter._add_rule(name, rule);
        },
-        /* .add_schema = */ [&](const std::string & name, const nlohmann::ordered_json & schema) {
+        /* .add_schema = */ [&](const std::string & name, const common_json & schema) {
            return converter.visit(schema, name == "root" ? "" : name);
        },
-        /* .resolve_refs = */ [&](nlohmann::ordered_json & schema) {
+        /* .resolve_refs = */ [&](common_json & schema) {
            converter.resolve_refs(schema, "");
        }
    };
@@ -1,12 +1,12 @@
 #pragma once

-#include <nlohmann/json_fwd.hpp>
+#include "json.h"

 #include <functional>
 #include <memory>
 #include <string>

-std::string json_schema_to_grammar(const nlohmann::ordered_json & schema,
+std::string json_schema_to_grammar(const common_json & schema,
                                   bool force_gbnf = false);

 class common_schema_converter;
@@ -24,14 +24,14 @@ class common_schema_info {
    common_schema_info(common_schema_info &&) noexcept;
    common_schema_info & operator=(common_schema_info &&) noexcept;

-    void resolve_refs(nlohmann::ordered_json & schema);
-    bool resolves_to_string(const nlohmann::ordered_json & schema);
+    void resolve_refs(common_json & schema);
+    bool resolves_to_string(const common_json & schema);
 };

 struct common_grammar_builder {
    std::function<std::string(const std::string &, const std::string &)> add_rule;
-    std::function<std::string(const std::string &, const nlohmann::ordered_json &)> add_schema;
-    std::function<void(nlohmann::ordered_json &)> resolve_refs;
+    std::function<std::string(const std::string &, const common_json &)> add_schema;
+    std::function<void(common_json &)> resolve_refs;
 };

 struct common_grammar_options {
@@ -0,0 +1,433 @@
+#include "json.h"
+
+#include "ggml.h"
+
+#define JSON_ASSERT GGML_ASSERT
+#include <nlohmann/json.hpp>
+
+#include <iterator>
+#include <new>
+#include <set>
+#include <unordered_map>
+#include <vector>
+
+using nlohmann::ordered_json;
+
+// a common_json is the backing value, so any value of a tree can be used as a common_json
+static_assert(sizeof(ordered_json)  <= sizeof(common_json),  "common_json storage is too small");
+static_assert(alignof(ordered_json) <= alignof(common_json), "common_json alignment is too weak");
+
+// runs fn and gives every error of the backing library as a common_json_error
+template <typename F>
+static decltype(auto) guard(F && fn) {
+    try {
+        return fn();
+    } catch (const ordered_json::exception & e) {
+        throw common_json_error(e.what());
+    }
+}
+
+static ordered_json & as_json(common_json * self) {
+    return *reinterpret_cast<ordered_json *>(self);
+}
+
+static const ordered_json & as_json(const common_json * self) {
+    return *reinterpret_cast<const ordered_json *>(self);
+}
+
+static common_json & as_common(ordered_json & json) {
+    return *reinterpret_cast<common_json *>(&json);
+}
+
+static const common_json & as_common(const ordered_json & json) {
+    return *reinterpret_cast<const common_json *>(&json);
+}
+
+static ordered_json to_json(const common_json_value & val) {
+    switch (val.type) {
+        case common_json_value::VAL_NULL:   return nullptr;
+        case common_json_value::VAL_BOOL:   return val.val_bool;
+        case common_json_value::VAL_INT:    return val.val_int;
+        case common_json_value::VAL_UINT:   return val.val_uint;
+        case common_json_value::VAL_DOUBLE: return val.val_double;
+        case common_json_value::VAL_STRING: return val.val_string;
+        case common_json_value::VAL_JSON:
+            // one owner means no one else can see this tree, so it is safe to move it out
+            // note: this makes a value single use, same as the json_ref of the backing library
+            if (val.val_json.use_count() == 1) {
+                return std::move(as_json(val.val_json.get()));
+            }
+            return as_json(val.val_json.get());
+    }
+
+    return nullptr;
+}
+
+common_json_value::common_json_value(const char * val) {
+    if (val) {
+        type       = VAL_STRING;
+        val_string = val;
+    } else {
+        type = VAL_NULL;
+    }
+}
+
+common_json_value::common_json_value(const common_json & val) :
+    type(VAL_JSON), val_json(std::make_shared<common_json>(val)) {}
+
+common_json_value::common_json_value(common_json && val) :
+    type(VAL_JSON), val_json(std::make_shared<common_json>(std::move(val))) {}
+
+// the ctors and get<T>() below are explicit specializations, giving strong symbols
+// an explicit instantiation is a weak symbol, dropped by some LTO builds (clang-cl)
+template <typename T>
+static std::shared_ptr<common_json> set_json(const std::set<T> & vals) {
+    common_json out = common_json::array();
+
+    for (const auto & val : vals) {
+        out.push_back(val);
+    }
+
+    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__> & vals) : type(VAL_JSON), val_json(set_json(vals)) {}
+
+COMMON_JSON_SET(int)
+COMMON_JSON_SET(std::string)
+
+#undef COMMON_JSON_SET
+
+template <typename T>
+static std::shared_ptr<common_json> map_json(const T & vals) {
+    common_json out = common_json::object();
+
+    for (const auto & val : vals) {
+        out.set({ val.first, val.second });
+    }
+
+    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__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
+
+COMMON_JSON_MAP(bool)
+COMMON_JSON_MAP(std::string)
+
+#undef COMMON_JSON_MAP
+
+// an unordered map value is usable only for the types below
+#define COMMON_JSON_UMAP(...) template <> common_json_value::common_json_value(const std::unordered_map<std::string, __VA_ARGS__> & vals) : type(VAL_JSON), val_json(map_json(vals)) {}
+
+COMMON_JSON_UMAP(size_t)
+
+#undef COMMON_JSON_UMAP
+
+template <typename T>
+static std::shared_ptr<common_json> vec_json(const std::vector<T> & vals) {
+    common_json out = common_json::array();
+
+    for (const auto & val : vals) {
+        out.push_back(val);
+    }
+
+    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__> & vals) : type(VAL_JSON), val_json(vec_json(vals)) {}
+
+COMMON_JSON_VEC(int)
+COMMON_JSON_VEC(unsigned char)
+COMMON_JSON_VEC(unsigned int)
+COMMON_JSON_VEC(long)
+COMMON_JSON_VEC(unsigned long)
+COMMON_JSON_VEC(long long)
+COMMON_JSON_VEC(unsigned long long)
+COMMON_JSON_VEC(float)
+COMMON_JSON_VEC(double)
+COMMON_JSON_VEC(std::string)
+COMMON_JSON_VEC(std::vector<float>)
+COMMON_JSON_VEC(common_json)
+
+#undef COMMON_JSON_VEC
+
+common_json_value::common_json_value(std::initializer_list<common_json_item> items) :
+    type(VAL_JSON), val_json(std::make_shared<common_json>(items)) {}
+
+// null, same as the backing library
+// operator[] turns it into an object, push_back() into an array
+common_json::common_json() {
+    new (storage) ordered_json();
+}
+
+common_json::common_json(const common_json & other) {
+    new (storage) ordered_json(as_json(&other));
+}
+
+common_json::common_json(common_json && other) noexcept {
+    new (storage) ordered_json(std::move(as_json(&other)));
+}
+
+common_json::common_json(std::initializer_list<common_json_item> items) {
+    new (storage) ordered_json(ordered_json::object());
+
+    for (const auto & item : items) {
+        set(item);
+    }
+}
+
+common_json::common_json(const common_json_value & val) {
+    new (storage) ordered_json(to_json(val));
+}
+
+common_json::common_json(std::nullptr_t) {
+    new (storage) ordered_json(nullptr);
+}
+
+common_json & common_json::operator=(common_json other) noexcept {
+    as_json(this).swap(as_json(&other));
+
+    return *this;
+}
+
+common_json::~common_json() {
+    as_json(this).~basic_json();
+}
+
+common_json common_json::parse(const std::string & text) {
+    try {
+        // the assignment moves the parsed tree in, it does not copy
+        common_json out;
+        as_json(&out) = ordered_json::parse(text);
+        return out;
+    } catch (const std::exception & e) {
+        throw common_json_error(e.what());
+    }
+}
+
+common_json common_json::parse_no_throw(const std::string & text) {
+    common_json out;
+    as_json(&out) = ordered_json::parse(text, nullptr, false);
+    return out;
+}
+
+bool common_json::is_discarded() const {
+    return as_json(this).is_discarded();
+}
+
+common_json common_json::array() {
+    common_json out;
+    as_json(&out) = ordered_json::array();
+    return out;
+}
+
+common_json common_json::array(std::initializer_list<common_json_value> vals) {
+    common_json out;
+    ordered_json & arr = as_json(&out);
+    arr = ordered_json::array();
+
+    for (const auto & val : vals) {
+        arr.push_back(to_json(val));
+    }
+
+    return out;
+}
+
+common_json common_json::object() {
+    common_json out;
+    as_json(&out) = ordered_json::object();
+    return out;
+}
+
+common_json common_json::object(std::initializer_list<common_json_item> items) {
+    return common_json(items);
+}
+
+common_json common_json::make(const common_json_value & val) {
+    return common_json(val);
+}
+
+bool common_json::is_null()           const { return as_json(this).is_null(); }
+bool common_json::is_object()         const { return as_json(this).is_object(); }
+bool common_json::is_array()          const { return as_json(this).is_array(); }
+bool common_json::is_string()         const { return as_json(this).is_string(); }
+bool common_json::is_boolean()        const { return as_json(this).is_boolean(); }
+bool common_json::is_number()         const { return as_json(this).is_number(); }
+bool common_json::is_number_integer() const { return as_json(this).is_number_integer(); }
+bool common_json::is_number_float()   const { return as_json(this).is_number_float(); }
+
+bool   common_json::empty() const { return as_json(this).empty(); }
+size_t common_json::size()  const { return as_json(this).size(); }
+
+bool common_json::contains(const std::string & key) const {
+    return as_json(this).contains(key);
+}
+
+bool common_json::operator==(const common_json_value & val) const {
+    // compare a tree in place, to_json() would copy it
+    if (val.type == common_json_value::VAL_JSON) {
+        return as_json(this) == as_json(val.val_json.get());
+    }
+    return as_json(this) == to_json(val);
+}
+
+bool common_json::operator!=(const common_json_value & val) const {
+    return !(*this == val);
+}
+
+common_json       & common_json::at(const std::string & key)       { return guard([&]() -> common_json       & { return as_common(as_json(this).at(key)); }); }
+const common_json & common_json::at(const std::string & key) const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(key)); }); }
+common_json       & common_json::at(size_t idx)                    { return guard([&]() -> common_json       & { return as_common(as_json(this).at(idx)); }); }
+const common_json & common_json::at(size_t idx)              const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(idx)); }); }
+
+common_json       & common_json::operator[](const std::string & key)       { return guard([&]() -> common_json       & { return as_common(as_json(this)[key]); }); }
+const common_json & common_json::operator[](const std::string & key) const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(key)); }); }
+common_json       & common_json::operator[](size_t idx)                    { return guard([&]() -> common_json       & { return as_common(as_json(this)[idx]); }); }
+const common_json & common_json::operator[](size_t idx)              const { return guard([&]() -> const common_json & { return as_common(as_json(this).at(idx)); }); }
+
+common_json       & common_json::front()       { return as_common(as_json(this).front()); }
+const common_json & common_json::front() const { return as_common(as_json(this).front()); }
+common_json       & common_json::back()        { return as_common(as_json(this).back()); }
+const common_json & common_json::back()  const { return as_common(as_json(this).back()); }
+
+void common_json::clear() {
+    as_json(this).clear();
+}
+
+void common_json::erase(const std::string & key) {
+    guard([&] { as_json(this).erase(key); });
+}
+
+void common_json::erase(size_t idx) {
+    guard([&] { as_json(this).erase(idx); });
+}
+
+void common_json::assign(const common_json_value & val) {
+    as_json(this) = to_json(val);
+}
+
+void common_json::set(const common_json_item & item) {
+    guard([&] { as_json(this)[item.key] = to_json(item.val); });
+}
+
+void common_json::push_back(const common_json_value & val) {
+    guard([&] { as_json(this).push_back(to_json(val)); });
+}
+
+void common_json::push_back(std::initializer_list<common_json_item> items) {
+    common_json val(items);
+
+    guard([&] { as_json(this).push_back(std::move(as_json(&val))); });
+}
+
+size_t common_json::count(const std::string & key) const {
+    return as_json(this).count(key);
+}
+
+void common_json::insert(const common_json & vals) {
+    guard([&] {
+        ordered_json & self = as_json(this);
+
+        self.insert(self.end(), as_json(&vals).begin(), as_json(&vals).end());
+    });
+}
+
+std::string common_json::dump(int indent) const {
+    return guard([&] { return as_json(this).dump(indent); });
+}
+
+std::string common_json::dump_safe(int indent) const {
+    return as_json(this).dump(indent, ' ', false, ordered_json::error_handler_t::replace);
+}
+
+// an array is indexed directly, an object needs a walk from the start
+common_json & common_json::iterator::operator*() const {
+    return guard([&]() -> common_json & {
+        ordered_json & j = as_json(node);
+
+        if (j.is_object()) {
+            return as_common(std::next(j.begin(), idx).value());
+        }
+        if (j.is_array()) {
+            return as_common(j[idx]);
+        }
+
+        // a plain value gives itself once, same as the backing library
+        return *node;
+    });
+}
+
+std::string common_json::iterator::key() const {
+    return guard([&] { return std::next(as_json(node).begin(), idx).key(); });
+}
+
+common_json::iterator common_json::begin() const {
+    return iterator(const_cast<common_json *>(this), 0);
+}
+
+common_json::iterator common_json::end() const {
+    return iterator(const_cast<common_json *>(this), size());
+}
+
+// the keys follow the backing library: the index for an array, "" for a plain value
+common_json::items_view::entry common_json::items_view::iterator::operator*() const {
+    return guard([&]() -> entry {
+        ordered_json & j = as_json(node);
+
+        if (j.is_object()) {
+            auto it = std::next(j.begin(), idx);
+
+            return { it.key(), as_common(it.value()) };
+        }
+        if (j.is_array()) {
+            return { std::to_string(idx), as_common(j[idx]) };
+        }
+
+        return { std::string(), *node };
+    });
+}
+
+common_json::items_view common_json::items() const {
+    return items_view(const_cast<common_json *>(this), size());
+}
+
+// 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;
+}
+
+// get<T>() is usable only for the types below
+
+#define COMMON_JSON_GET(...) template <> __VA_ARGS__ common_json::get<__VA_ARGS__>() const { return guard([&] { return as_json(this).get<__VA_ARGS__>(); }); }
+
+COMMON_JSON_GET(bool)
+COMMON_JSON_GET(int)
+COMMON_JSON_GET(unsigned int)
+COMMON_JSON_GET(long)
+COMMON_JSON_GET(unsigned long)
+COMMON_JSON_GET(long long)
+COMMON_JSON_GET(unsigned long long)
+COMMON_JSON_GET(float)
+COMMON_JSON_GET(double)
+COMMON_JSON_GET(std::string)
+COMMON_JSON_GET(std::vector<float>)
+COMMON_JSON_GET(std::vector<std::string>)
+COMMON_JSON_GET(std::set<std::string>)
+COMMON_JSON_GET(std::vector<int>)
+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);
+}
@@ -0,0 +1,352 @@
+#pragma once
+
+#include <cstddef>
+#include <cstdint>
+#include <initializer_list>
+#include <iterator>
+#include <map>
+#include <memory>
+#include <set>
+#include <stdexcept>
+#include <string>
+#include <string_view>
+#include <type_traits>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+// common_json, a thin wrapper around vendor json library
+// the underlay library is pimpl, we are using nlohmann::json for now
+//
+// many features of the library are deliberately left out, to keep this interface small and generic and to keep compile time down
+//
+// some main differences compared to nlohmann::json :
+// - object keys keep the order in which they are added
+// - errors are always throw as common_json_error
+// - obj.push_back({key, val}) is intentionally unsupported to avoid confusion with push_back on a vector; write it as obj[key] = val for clarity
+// - a braced pair in value position does not build, e.g. {"key", {"a", "b"}}; write array({"a", "b"}) where nlohmann made an array
+//
+// in doubt, search the code base for an existing usage example; do not add anything to this header unless absolutely necessary
+
+class common_json;
+
+// common_json_value holds a list of these, and each of them holds a value, so one must come first
+struct common_json_item;
+
+struct common_json_error : std::runtime_error {
+    using std::runtime_error::runtime_error;
+};
+
+// one value, tagged so that this header stays free of the backing library
+// note: a value that holds a tree is single use, the second use gives null
+struct common_json_value {
+    enum value_type {
+        VAL_NULL,
+        VAL_BOOL,
+        VAL_INT,
+        VAL_UINT,
+        VAL_DOUBLE,
+        VAL_STRING,
+        VAL_JSON,
+    };
+
+    value_type type = VAL_NULL;
+
+    union {
+        bool     val_bool;
+        int64_t  val_int;
+        uint64_t val_uint = 0;
+        double   val_double;
+    };
+
+    std::string                  val_string;
+    std::shared_ptr<common_json> val_json;
+
+    common_json_value(std::nullptr_t = nullptr) : type(VAL_NULL) {}
+    common_json_value(bool val) : type(VAL_BOOL), val_bool(val) {}
+    common_json_value(std::string val) : type(VAL_STRING), val_string(std::move(val)) {}
+    // without this a string_view lands on the common_json ctor below and recurses
+    common_json_value(std::string_view val) : type(VAL_STRING), val_string(val) {}
+    common_json_value(const char * val);
+    common_json_value(const common_json & val);
+    common_json_value(common_json && val);
+    // only for the types instantiated in json.cpp, the rest fails at link time
+    template <typename T> common_json_value(const std::vector<T> & vals);
+    // a set becomes an array, in the set's own order
+    template <typename T> common_json_value(const std::set<T> & vals);
+    // a map becomes an object, keyed in the map's own order
+    template <typename T> common_json_value(const std::map<std::string, T> & vals);
+    template <typename T> common_json_value(const std::unordered_map<std::string, T> & vals);
+
+    // nested object, e.g. {"fn", {{"name", "x"}}}
+    // note: a nested pair {"a", "b"} does not build, use common_json::array({"a", "b"}) for an array
+    common_json_value(std::initializer_list<common_json_item> items);
+
+    template <typename T, typename std::enable_if<std::is_integral<T>::value && !std::is_same<T, bool>::value, int>::type = 0>
+    common_json_value(T val) : type(std::is_signed<T>::value ? VAL_INT : VAL_UINT) {
+        if (std::is_signed<T>::value) {
+            val_int = (int64_t) val;
+        } else {
+            val_uint = (uint64_t) val;
+        }
+    }
+
+    template <typename T, typename std::enable_if<std::is_floating_point<T>::value, int>::type = 0>
+    common_json_value(T val) : type(VAL_DOUBLE), val_double((double) val) {}
+};
+
+struct common_json_item {
+    std::string       key;
+    common_json_value val;
+
+    template <typename T>
+    common_json_item(std::string key, T && val) :
+        key(std::move(key)), val(std::forward<T>(val)) {}
+
+    // a braced list cannot deduce T, so it needs its own overload
+    common_json_item(std::string key, std::initializer_list<common_json_item> items) :
+        key(std::move(key)), val(items) {}
+};
+
+// the types common_json_value holds on its own
+// anything else reaches its common_json ctor and recurses forever
+template <typename T> struct common_json_is_value : std::integral_constant<bool,
+    std::is_arithmetic<T>::value ||
+    std::is_same<T, std::nullptr_t>::value ||
+    std::is_same<T, std::string>::value ||
+    std::is_same<T, std::string_view>::value ||
+    std::is_same<T, char *>::value ||
+    std::is_same<T, const char *>::value ||
+    std::is_same<T, common_json>::value> {};
+
+template <typename T, typename A>
+struct common_json_is_value<std::vector<T, A>> : std::true_type {};
+
+template <typename T, typename C, typename A>
+struct common_json_is_value<std::set<T, C, A>> : std::true_type {};
+
+template <typename V, typename C, typename A>
+struct common_json_is_value<std::map<std::string, V, C, A>> : std::true_type {};
+
+template <typename V, typename H, typename E, typename A>
+struct common_json_is_value<std::unordered_map<std::string, V, H, E, A>> : std::true_type {};
+
+class common_json {
+  public:
+    common_json();
+    common_json(const common_json & other);
+    common_json(common_json && other) noexcept;
+    common_json(std::initializer_list<common_json_item> items);
+    common_json(const common_json_value & val);
+
+    // direct, a value would need two conversions in a row
+    common_json(std::nullptr_t);
+
+    // one step, so that "abc" or a vector can go straight into a common_json
+    template <typename T, typename std::enable_if<!std::is_same<typename std::decay<T>::type, common_json>::value &&
+                                                  !std::is_same<typename std::decay<T>::type, common_json_value>::value, int>::type = 0>
+    common_json(T && val) : common_json(common_json_value(std::forward<T>(val))) {
+        static_assert(common_json_is_value<typename std::decay<T>::type>::value,
+                      "no common_json_value ctor holds this type, add one instead of letting it recurse");
+    }
+
+    // by value, same as the backing library
+    // the right side is copied before the left side can invalidate it, e.g. msg["a"] = msg.at("b")
+    common_json & operator=(common_json other) noexcept;
+
+    ~common_json();
+
+    // throws common_json_error if the text is not valid JSON
+    static common_json parse(const std::string & text);
+
+    // gives a discarded value instead of throwing, check it with is_discarded()
+    static common_json parse_no_throw(const std::string & text);
+
+    bool is_discarded() const;
+
+    static common_json array();
+    static common_json array(std::initializer_list<common_json_value> vals);
+    static common_json object();
+    static common_json object(std::initializer_list<common_json_item> items);
+
+    // holds a single value, e.g. make("abc").dump() gives "\"abc\""
+    static common_json make(const common_json_value & val);
+
+    bool is_null()    const;
+    bool is_object()  const;
+    bool is_array()   const;
+    bool is_string()  const;
+    bool is_boolean() const;
+    bool is_number()  const;
+    bool is_number_integer() const;
+    bool is_number_float()   const;
+
+    bool   empty() const;
+    size_t size()  const;
+
+    bool contains(const std::string & key) const;
+
+    bool operator==(const common_json_value & val) const;
+    bool operator!=(const common_json_value & val) const;
+
+    // at() throws common_json_error if the key is missing, operator[] adds a null value instead
+    // note: a const operator[] cannot add, it throws like at()
+    common_json       & at(const std::string & key);
+    const common_json & at(const std::string & key) const;
+    common_json       & at(size_t idx);
+    const common_json & at(size_t idx) const;
+
+    common_json       & operator[](const std::string & key);
+    const common_json & operator[](const std::string & key) const;
+    common_json       & operator[](const char * key)       { return (*this)[std::string(key)]; }
+    const common_json & operator[](const char * key) const { return (*this)[std::string(key)]; }
+    common_json       & operator[](int idx)       { return (*this)[to_idx(idx)]; }
+    const common_json & operator[](int idx) const { return (*this)[to_idx(idx)]; }
+    common_json       & operator[](size_t idx);
+    const common_json & operator[](size_t idx) const;
+
+    common_json       & front();
+    const common_json & front() const;
+    common_json       & back();
+    const common_json & back()  const;
+
+    void clear();
+
+    void erase(const std::string & key);
+    void erase(size_t idx);
+
+    // only for the types instantiated in json.cpp, the rest fails at link time
+    template <typename T> T get() const;
+
+    // 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;
+
+    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;
+
+    // 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 {
+        return contains(key) ? at(key) : def;
+    }
+
+    void assign(const common_json_value & val);
+    void set(const common_json_item & item);
+    void push_back(const common_json_value & val);
+
+    // appends one object, e.g. push_back({{"a", 1}})
+    void push_back(std::initializer_list<common_json_item> items);
+
+    // 1 if the key is there, 0 if not
+    size_t count(const std::string & key) const;
+
+    // appends every value of another array; inserting an array into itself throws
+    void insert(const common_json & vals);
+
+    // a common_json goes through the copy assignment above, everything else becomes a value
+    template <typename T, typename std::enable_if<!std::is_same<typename std::decay<T>::type, common_json>::value, int>::type = 0>
+    common_json & operator=(T && val) {
+        assign(common_json_value(std::forward<T>(val)));
+        return *this;
+    }
+
+    std::string dump(int indent = -1) const;
+
+    // same as dump(), but bad UTF-8 gets replaced instead of throwing
+    std::string dump_safe(int indent = -1) const;
+
+    // walks an array by index, or an object in insertion order
+    // a plain value gives itself once, same as the backing library
+    class iterator {
+      public:
+        using iterator_category = std::forward_iterator_tag;
+        using value_type        = common_json;
+        using difference_type   = std::ptrdiff_t;
+        using pointer           = common_json *;
+        using reference         = common_json &;
+
+        iterator(common_json * node, size_t idx) : node(node), idx(idx) {}
+
+        common_json & operator*() const;
+        common_json & value()     const { return **this; }
+        std::string   key()       const;
+
+        iterator & operator++() {
+            idx++;
+            return *this;
+        }
+
+        bool operator!=(const iterator & other) const { return idx != other.idx; }
+        bool operator==(const iterator & other) const { return idx == other.idx; }
+
+      private:
+        common_json * node;
+        size_t        idx;
+    };
+
+    iterator begin() const;
+    iterator end()   const;
+
+    // allows: for (const auto & [key, val] : obj.items())
+    class items_view {
+      public:
+        // the members are public, so an entry also works with structured bindings
+        struct entry {
+            std::string   k;
+            common_json & v;
+
+            const std::string & key()   const { return k; }
+            common_json &       value() const { return v; }
+        };
+
+        items_view(common_json * node, size_t n) : node(node), n(n) {}
+
+        class iterator {
+          public:
+            iterator(common_json * node, size_t idx) : node(node), idx(idx) {}
+
+            entry operator*() const;
+
+            iterator & operator++() {
+                idx++;
+                return *this;
+            }
+
+            bool operator!=(const iterator & other) const { return idx != other.idx; }
+
+          private:
+            common_json * node;
+            size_t        idx;
+        };
+
+        iterator begin() const { return iterator(node, 0); }
+        iterator end()   const { return iterator(node, n); }
+
+      private:
+        common_json * node;
+        size_t        n;
+    };
+
+    items_view items() const;
+
+  private:
+    // a negative index must not turn into a huge size_t
+    static size_t to_idx(int idx) {
+        if (idx < 0) {
+            throw common_json_error("negative array index");
+        }
+        return (size_t) idx;
+    }
+
+    // the backing value is built here, json.cpp checks that it fits
+    // it cannot be a pointer: a value inside a tree would then not be a common_json
+    // at() could then only give back a copy instead of a real reference
+    alignas(8) unsigned char storage[32];
+};
+
+using common_json_entry = common_json::items_view::entry;
@@ -10,7 +10,6 @@
 #include <initializer_list>
 #include <map>
 #include <memory>
-#include <nlohmann/json.hpp>
 #include <regex>
 #include <set>
 #include <stdexcept>
@@ -1120,8 +1119,8 @@ common_peg_parser common_peg_parser_builder::chars(const std::string & classes,
    return wrap(arena_.add_parser(common_peg_chars_parser{classes, ranges, negated, min, max}));
 }

-common_peg_parser common_peg_parser_builder::schema(const common_peg_parser & p, const std::string & name, const nlohmann::ordered_json & schema, bool raw) {
-    return wrap(arena_.add_parser(common_peg_schema_parser{p.id(), name, std::make_shared<nlohmann::ordered_json>(schema), raw}));
+common_peg_parser common_peg_parser_builder::schema(const common_peg_parser & p, const std::string & name, const common_json & schema, bool raw) {
+    return wrap(arena_.add_parser(common_peg_schema_parser{p.id(), name, std::make_shared<common_json>(schema), raw}));
 }

 common_peg_parser common_peg_parser_builder::rule(const std::string & name, const common_peg_parser & p, bool trigger) {
@@ -1805,8 +1804,8 @@ void common_peg_arena::build_grammar(const common_grammar_builder & builder, boo
    }
 }

-static nlohmann::json serialize_parser_variant(const common_peg_parser_variant & variant) {
-    using json = nlohmann::json;
+static common_json serialize_parser_variant(const common_peg_parser_variant & variant) {
+    using json = common_json;

    return std::visit([](const auto & p) -> json {
        using T = std::decay_t<decltype(p)>;
@@ -1860,7 +1859,7 @@ static nlohmann::json serialize_parser_variant(const common_peg_parser_variant &
                {"type", "schema"},
                {"child", p.child},
                {"name", p.name},
-                {"schema", p.schema ? *p.schema : nullptr},
+                {"schema", p.schema ? *p.schema : json(nullptr)},
                {"raw", p.raw}
            };
        } else if constexpr (std::is_same_v<T, common_peg_rule_parser>) {
@@ -1888,19 +1887,19 @@ static nlohmann::json serialize_parser_variant(const common_peg_parser_variant &
    }, variant);
 }

-nlohmann::json common_peg_arena::to_json() const {
-    auto parsers = nlohmann::json::array();
+common_json common_peg_arena::to_json() const {
+    auto parsers = common_json::array();
    for (const auto & parser : parsers_) {
        parsers.push_back(serialize_parser_variant(parser));
    }
-    return nlohmann::json{
+    return common_json{
        {"parsers", parsers},
        {"rules", rules_},
        {"root", root_}
    };
 }

-static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json & j) {
+static common_peg_parser_variant deserialize_parser_variant(const common_json & j) {
    if (!j.contains("type") || !j["type"].is_string()) {
        throw std::runtime_error("Parser variant JSON missing or invalid 'type' field");
    }
@@ -1969,9 +1968,9 @@ static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json
        }
        common_peg_chars_parser parser;
        parser.pattern = j["pattern"];
-        parser.negated = j["negated"];
-        parser.min_count = j["min_count"];
-        parser.max_count = j["max_count"];
+        parser.negated = j["negated"].get<bool>();
+        parser.min_count = j["min_count"].get<int>();
+        parser.max_count = j["max_count"].get<int>();
        for (const auto & range_json : j["ranges"]) {
            if (!range_json.contains("start") || !range_json.contains("end")) {
                throw std::runtime_error("char_range missing 'start' or 'end' field");
@@ -2007,7 +2006,7 @@ static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json
        parser.child = j["child"].get<common_peg_parser_id>();
        parser.name = j["name"];
        if (!j["schema"].is_null()) {
-            parser.schema = std::make_shared<nlohmann::ordered_json>(j["schema"]);
+            parser.schema = std::make_shared<common_json>(j["schema"]);
        }
        parser.raw = j["raw"].get<bool>();
        return parser;
@@ -2069,7 +2068,7 @@ static common_peg_parser_variant deserialize_parser_variant(const nlohmann::json
    throw std::runtime_error("Unknown parser type: " + type);
 }

-common_peg_arena common_peg_arena::from_json(const nlohmann::json & j) {
+common_peg_arena common_peg_arena::from_json(const common_json & j) {
    if (!j.contains("parsers") || !j["parsers"].is_array()) {
        throw std::runtime_error("JSON missing or invalid 'parsers' array");
    }
@@ -2109,7 +2108,7 @@ std::string common_peg_arena::save() const {
 }

 void common_peg_arena::load(const std::string & data) {
-    *this = from_json(nlohmann::json::parse(data));
+    *this = from_json(common_json::parse(data));
 }

 common_peg_arena build_peg_parser(const std::function<common_peg_parser(common_peg_parser_builder & builder)> & fn) {
@@ -1,6 +1,6 @@
 #pragma once

-#include <nlohmann/json_fwd.hpp>
+#include "json.h"

 #include <memory>
 #include <set>
@@ -245,7 +245,7 @@ struct common_peg_until_parser {
 struct common_peg_schema_parser {
    common_peg_parser_id child;
    std::string name;
-    std::shared_ptr<nlohmann::ordered_json> schema;
+    std::shared_ptr<common_json> schema;

    // Indicates if the GBNF should accept a raw string that matches the schema.
    bool raw;
@@ -332,8 +332,8 @@ class common_peg_arena {

    std::string dump(common_peg_parser_id id) const;

-    nlohmann::json to_json() const;
-    static common_peg_arena from_json(const nlohmann::json & j);
+    common_json to_json() const;
+    static common_peg_arena from_json(const common_json & j);

    std::string save() const;
    void load(const std::string & data);
@@ -490,7 +490,7 @@ class common_peg_parser_builder {

    // Wraps a parser with JSON schema metadata for grammar generation.
    // Used internally to convert JSON schemas to GBNF grammar rules.
-    common_peg_parser schema(const common_peg_parser & p, const std::string & name, const nlohmann::ordered_json & schema, bool raw = false);
+    common_peg_parser schema(const common_peg_parser & p, const std::string & name, const common_json & schema, bool raw = false);

    // Creates a named rule, stores it in the grammar, and returns a ref.
    // If trigger=true, marks this rule as an entry point for lazy grammar generation.
@@ -2388,6 +2388,9 @@ common_speculative_init_result::common_speculative_init_result(
        cparams.ctx_type = LLAMA_CONTEXT_TYPE_MTP;
    }

+    // the draft context holds as many tokens per sequence as the target context
+    cparams.n_ctx = llama_n_ctx(ctx_tgt);
+
    // note: for small models maybe we can set this to the maximum possible draft from all speculative types
    //       the extra memory for small models is likely negligible?
    cparams.n_rs_seq  = 0;
@@ -58,6 +58,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
    "DSparkDraftModel": "qwen",
    "DSparkSpeculator": "qwen",
    "Lfm2DSparkDraftModel": "qwen",
+    "LingDSparkModel": "qwen",
    "DeepseekV4ForCausalLM": "deepseek",
    "DeepseekV4DSparkModel": "deepseek",
    "DistilBertForMaskedLM": "bert",
@@ -709,7 +709,13 @@ class DFlashModel(Qwen3Model):
        yield from super().modify_tensors(data_torch, name, bid)


-@ModelBase.register("Qwen3DSparkModel", "DSparkDraftModel", "DSparkSpeculator", "Lfm2DSparkDraftModel")
+@ModelBase.register(
+    "Qwen3DSparkModel",
+    "DSparkDraftModel",
+    "DSparkSpeculator",
+    "Lfm2DSparkDraftModel",
+    "LingDSparkModel",
+)
@ModelBase.example("satgeze/Qwen3.6-27B-DSpark")
 class DSparkModel(DFlashModel):
    # DSpark = DFlash + a semi-autoregressive Markov head.
@@ -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: {
@@ -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:
@@ -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);
+}
@@ -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);
@@ -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) {
@@ -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]);
+}
@@ -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"}});
@@ -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(
@@ -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",
@@ -1038,6 +1038,7 @@ bool llm_arch_supports_rs_rollback(const llm_arch & arch) {
        case LLM_ARCH_NEMOTRON_H_MOE:
        case LLM_ARCH_LFM2:
        case LLM_ARCH_LFM2MOE:
+        case LLM_ARCH_BAILINGMOE3:
            return true;
        default:
            return false;
@@ -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);
    }
@@ -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);
@@ -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;
@@ -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);
    }
 }

@@ -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) {
@@ -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);
@@ -1,6 +1,8 @@
 #include "models.h"
 #include "llama-memory-recurrent.h"

+#include <algorithm>
+
 void llama_model_bailingmoe3::load_arch_hparams(llama_model_loader & ml) {
    ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS,      hparams.f_norm_rms_eps);
    ml.get_key(LLM_KV_ATTENTION_KEY_LENGTH_MLA,         hparams.n_embd_head_k_mla_impl);
@@ -179,7 +181,9 @@ static ggml_tensor * bailingmoe3_causal_conv1d(
        int64_t n_seq_tokens,
        int64_t n_seqs,
        int64_t n_tokens,
-        int64_t cache_head) {
+        int64_t cache_head,
+        uint32_t mem_size,
+        uint32_t n_rs_seq) {
    const int64_t d_inner = head_dim * n_head;
    const int64_t conv_state_size = (d_conv - 1) * d_inner;
    const int64_t total_state_size = 3 * conv_state_size;
@@ -193,13 +197,18 @@ static ggml_tensor * bailingmoe3_causal_conv1d(
    x_proj = ggml_reshape_3d(ctx0, x_proj, d_inner, n_seq_tokens, n_seqs);
    ggml_tensor * conv_x = ggml_concat(ctx0, conv_state, ggml_transpose(ctx0, x_proj), 0);

-    ggml_tensor * last_conv_x = ggml_view_3d(ctx0, conv_x, d_conv - 1, d_inner, n_seqs,
-            conv_x->nb[1], conv_x->nb[2], n_seq_tokens * conv_x->nb[0]);
-    ggml_build_forward_expand(gf, ggml_cpy(ctx0, last_conv_x,
-            ggml_view_3d(ctx0, conv_states_all, d_conv - 1, d_inner, n_seqs,
-                (d_conv - 1) * ggml_element_size(conv_states_all),
-                total_state_size * ggml_element_size(conv_states_all),
-                (cache_head * total_state_size + qkv * conv_state_size) * ggml_element_size(conv_states_all))));
+    const int64_t K = (int64_t) n_rs_seq + 1;
+    const int64_t n_written = std::min<int64_t>(n_seq_tokens, K);
+
+    for (int64_t slot = 0; slot < n_written; ++slot) {
+        ggml_tensor * conv_snap = ggml_view_3d(ctx0, conv_x, d_conv - 1, d_inner, n_seqs,
+                conv_x->nb[1], conv_x->nb[2], (conv_x->ne[0] - (d_conv - 1) - slot) * conv_x->nb[0]);
+        ggml_build_forward_expand(gf, ggml_cpy(ctx0, conv_snap,
+                ggml_view_3d(ctx0, conv_states_all, d_conv - 1, d_inner, n_seqs,
+                    (d_conv - 1) * ggml_element_size(conv_states_all),
+                    total_state_size * ggml_element_size(conv_states_all),
+                    ((slot * mem_size + cache_head) * total_state_size + qkv * conv_state_size) * ggml_element_size(conv_states_all))));
+    }

    ggml_tensor * conv_weight = ggml_reshape_2d(ctx0, conv_w, d_conv, d_inner);
    ggml_tensor * out = ggml_ssm_conv(ctx0, conv_x, conv_weight);
@@ -237,6 +246,8 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph
    GGML_ASSERT(ubatch.n_tokens == n_seq_tokens * n_seqs);

    for (int il = 0; il < n_layer; ++il) {
+        res->t_layer_inp[il] = inpL;
+
        const auto & layer = model.layers[il];
        ggml_tensor * inpSA = inpL;
        ggml_tensor * cur = build_norm(inpL, layer.attn_norm, nullptr, LLM_NORM_RMS, il);
@@ -245,18 +256,19 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph
        if (hparams.is_recr(il)) {
            const auto * mctx_cur = inp_rs->mctx;
            const auto cache_head = mctx_cur->get_head();
+            const auto mem_size = mctx_cur->get_size();
            ggml_tensor * conv_states_all = mctx_cur->get_r_l(il);
            ggml_tensor * conv_state_all = build_rs(inp_rs, conv_states_all, hparams.n_embd_r(), n_seqs);

            ggml_tensor * q = bailingmoe3_causal_conv1d(
                    gf, ctx0, conv_states_all, conv_state_all, 0, cur, layer.wq, layer.ssm_q_conv,
-                    d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head);
+                    d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head, mem_size, cparams.n_rs_seq);
            ggml_tensor * k = bailingmoe3_causal_conv1d(
                    gf, ctx0, conv_states_all, conv_state_all, 1, cur, layer.wk, layer.ssm_k_conv,
-                    d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head);
+                    d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head, mem_size, cparams.n_rs_seq);
            ggml_tensor * v = bailingmoe3_causal_conv1d(
                    gf, ctx0, conv_states_all, conv_state_all, 2, cur, layer.wv, layer.ssm_v_conv,
-                    d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head);
+                    d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head, mem_size, cparams.n_rs_seq);

            ggml_tensor * gate = ggml_mul_mat(ctx0, layer.ssm_f_a, cur);
            gate = ggml_add(ctx0, gate, layer.ssm_dt_b);
@@ -276,11 +288,8 @@ llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph
            ggml_tensor * state = build_rs(inp_rs, states_all, hparams.n_embd_s(), n_seqs);
            state = ggml_reshape_4d(ctx0, state, head_dim, head_dim, n_head, n_seqs);

-            auto result = build_delta_net(q, k, v, gate, beta, state, il);
-            ggml_tensor * out = ggml_cont(ctx0, result.first);
-            ggml_build_forward_expand(gf, ggml_cpy(ctx0, result.second,
-                    ggml_view_1d(ctx0, states_all, hparams.n_embd_s() * n_seqs,
-                        cache_head * hparams.n_embd_s() * ggml_element_size(states_all))));
+            ggml_tensor * out = ggml_cont(ctx0, build_recurrent_attn(
+                    inp_rs, states_all, q, k, v, gate, beta, state, il));

            ggml_tensor * out_gate = ggml_mul_mat(ctx0, layer.ssm_g_a, cur);
            out_gate = ggml_reshape_3d(ctx0, out_gate, head_dim, n_head, n_tokens);
@@ -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);
@@ -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)
@@ -8,7 +8,7 @@ void test_json_serialization(testing &t) {
    auto json_serialized = original.to_json().dump();

    t.test("compare before/after", [&](testing &t) {
-        auto deserialized = common_peg_arena::from_json(nlohmann::json::parse(json_serialized));
+        auto deserialized = common_peg_arena::from_json(common_json::parse(json_serialized));

        // Test complex JSON
        std::string input = R"({"name": "test", "values": [1, 2, 3], "nested": {"a": true}})";
@@ -23,6 +23,6 @@ void test_json_serialization(testing &t) {
    });

    t.bench("deserialize", [&]() {
-        auto deserialized = common_peg_arena::from_json(nlohmann::json::parse(json_serialized));
+        auto deserialized = common_peg_arena::from_json(common_json::parse(json_serialized));
    }, 100);
 }
@@ -1,7 +1,7 @@
 #pragma once

 // Common includes for all test files
-#include <nlohmann/json.hpp>
+#include "json.h"
 #include <string>
 #include <vector>

@@ -11,9 +11,9 @@
 #include "simple-tokenize.h"

 struct bench_tool_call {
-    std::string            id;
-    std::string            name;
-    nlohmann::ordered_json args;
+    std::string id;
+    std::string name;
+    common_json args;
 };

 // Test function declarations
@@ -11,9 +11,9 @@
 #include <regex>
 #include <string>

-#include "nlohmann/json.hpp"
+#include "json.h"

-using json = nlohmann::ordered_json;
+using json = common_json;

 static json create_tools();
 static void test_example_native(testing & t);
@@ -63,10 +63,10 @@ static json create_tools() {
                          { { "type", "string" }, { "description", "The city and state, e.g. San Francisco, CA" } } },
                        { "unit",
                          { { "type", "string" },
-                            { "enum", { "celsius", "fahrenheit" } },
+                            { "enum", json::array({ "celsius", "fahrenheit" }) },
                            { "description",
                              "The temperature unit to use. Infer this from the users location." } } } } },
-                    { "required", { "location", "unit" } },
+                    { "required", json::array({ "location", "unit" }) },
                } },
          }                      }
    };
@@ -86,14 +86,14 @@ static json create_tools() {
                          { { "type", "string" }, { "description", "The city and state, e.g. San Francisco, CA" } } },
                        { "unit",
                          { { "type", "string" },
-                            { "enum", { "celsius", "fahrenheit" } },
+                            { "enum", json::array({ "celsius", "fahrenheit" }) },
                            { "description", "The temperature unit to use. Infer this from the users location." } } },
                        { "days",
                          { { "type", "integer" },
                            { "description", "Number of days to forecast (1-10)" },
                            { "minimum", 1 },
                            { "maximum", 10 } } } } },
-                    { "required", { "location", "unit" } },
+                    { "required", json::array({ "location", "unit" }) },
                } },
          }                      }
    };
@@ -114,9 +114,9 @@ static json create_tools() {
                        { "default", 5 } } },
                    { "category",
                      { { "type", "string" },
-                        { "enum", { "api", "troubleshooting", "billing", "general" } },
+                        { "enum", json::array({ "api", "troubleshooting", "billing", "general" }) },
                        { "description", "Filter search by specific category." } } } } },
-                { "required", { "query", "category" } },
+                { "required", json::array({ "query", "category" }) },
                { "additionalProperties", false } } },
            { "strict", true } } }
    };
@@ -341,7 +341,7 @@ static void test_example_native(testing & t) {
                { { "invoice_number", { { "type", "string" } } },
                  { "amount", { { "type", "number" } } },
                  { "due_date", { { "type", "string" } } } } },
-              { "required", { "invoice_number", "amount", "due_date" } } },
+              { "required", json::array({ "invoice_number", "amount", "due_date" }) } },
         /* .parallel_tool_calls =  */ false,
         /* .generation_prompt =    */ "<think>",
         /* .input =                */
@@ -406,7 +406,7 @@ static void test_example_qwen3_coder(testing & t) {

            std::set<std::string> required_properties;
            if (function.contains("required")) {
-                function.at("required").get_to(required_properties);
+                required_properties = function.at("required").get<std::set<std::string>>();
            }

            std::vector<common_peg_parser> arg_parsers;
@@ -661,8 +661,8 @@ void test_command7_parser_compare(testing & t) {
        "5. Provide a detailed cost breakdown that includes accommodation, transportation, meals, and entry fees "
        "to attractions.";

-    std::vector<std::tuple<std::string, std::string, nlohmann::json>> tool_calls = {
-        { "call_0", "plan_trip", nlohmann::json::parse(R"({
+    std::vector<std::tuple<std::string, std::string, common_json>> tool_calls = {
+        { "call_0", "plan_trip", common_json::parse(R"({
            "destination": "Japan",
            "duration": 14,
            "budget": 4000,
@@ -686,16 +686,16 @@ void test_command7_parser_compare(testing & t) {
    if (!tool_calls.empty()) {
        tokens.emplace_back("<|START_ACTION|>");

-        auto json = nlohmann::json::array();
+        auto json = common_json::array();
        for (const auto & tc : tool_calls) {
-            auto tc_json            = nlohmann::json::object();
+            auto tc_json            = common_json::object();
            tc_json["tool_call_id"] = std::get<0>(tc);
            tc_json["tool_name"]    = std::get<1>(tc);
            tc_json["parameters"]   = std::get<2>(tc);
            json.push_back(tc_json);
        }

-        auto tokenized = simple_tokenize(json.dump(-1, ' ', true));
+        auto tokenized = simple_tokenize(json.dump(-1));
        tokens.insert(tokens.end(), tokenized.begin(), tokenized.end());

        tokens.emplace_back("<|END_ACTION|>");
@@ -737,7 +737,7 @@ static void test_prefix_tool_names(testing & t) {
                      {
                          { "arg1", { { "type", "integer" } } },
                      } },
-                    { "required", { "arg1" } },
+                    { "required", json::array({ "arg1" }) },
                } },
          } }
    };
@@ -757,7 +757,7 @@ static void test_prefix_tool_names(testing & t) {
                          { "arg1", { { "type", "integer" } } },
                          { "arg2", { { "type", "integer" } } },
                      } },
-                    { "required", { "arg1" } },
+                    { "required", json::array({ "arg1" }) },
                } },
          } }
    };
@@ -7,7 +7,7 @@
 #include <fstream>
 #include <filesystem>

-#include <nlohmann/json.hpp>
+#include "json.h"

 #undef NDEBUG
 #include <cassert>
@@ -20,7 +20,7 @@
 #include "jinja/lexer.h"
 #include "jinja/caps.h"

-using json = nlohmann::ordered_json;
+using json = common_json;

 static int main_automated_tests(void);

@@ -304,8 +304,8 @@ void run_single(const std::string& contents, json input, bool use_common, bool d
        if (input.contains("eos_token")) {
            eos_token = input["eos_token"].get<std::string>();
        }
-        nlohmann::ordered_json msgs_json = input["messages"];
-        nlohmann::ordered_json tools_json = input["tools"];
+        common_json msgs_json = input["messages"];
+        common_json tools_json = input["tools"];
        auto messages = common_chat_msgs_parse_oaicompat(msgs_json);
        auto tools = common_chat_tools_parse_oaicompat(tools_json);
        auto output = format_using_common(contents, bos_token, eos_token, messages, tools);
@@ -19,12 +19,12 @@
 #include <fstream>
 #include <functional>
 #include <iostream>
-#include <nlohmann/json.hpp>
+#include "json.h"
 #include <set>
 #include <stdexcept>
 #include <string>

-using json = nlohmann::ordered_json;
+using json = common_json;

 static std::ostream & operator<<(std::ostream & os, const common_chat_msg_diff & diff) {
    os << "{ content_delta: " << diff.content_delta << "; ";
@@ -7,13 +7,13 @@
 #include "../src/unicode.h"
 #include "../src/llama-grammar.h"

-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <cassert>
 #include <string>
 #include <vector>

-using json = nlohmann::ordered_json;
+using json = common_json;

 static llama_grammar * build_grammar_with_root(const std::string & grammar_str, const char * grammar_root) {
    return llama_grammar_init_impl(nullptr, grammar_str.c_str(), grammar_root, false, nullptr, 0, nullptr, 0);
@@ -3,7 +3,7 @@
 #include <random>
 #include <cstdlib>

-#include <nlohmann/json.hpp>
+#include "json.h"
 #include "subproc.h"

 #include "jinja/runtime.h"
@@ -14,7 +14,7 @@

 #include "testing.h"

-using json = nlohmann::ordered_json;
+using json = common_json;

 static void test_template(testing & t, const std::string & name, const std::string & tmpl, const json & vars, const std::string & expect);

@@ -240,7 +240,7 @@ static void test_conditionals(testing & t) {

    test_template(t, "is undefined key falsy",
        "{{ 'yes' if not y['x'] else 'no' }}",
-        {{"y", {{}}}},
+        {{"y", json::array({nullptr})}},
        "yes"
    );

@@ -282,7 +282,7 @@ static void test_conditionals(testing & t) {

    test_template(t, "is non-empty object truthy",
        "{{ 'yes' if y else 'no' }}",
-        {{"y", {"x", false}}},
+        {{"y", json::array({"x", false})}},
        "yes"
    );

@@ -6,7 +6,7 @@

 #include "../src/llama-grammar.h"

-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <cassert>
 #include <fstream>
@@ -1442,7 +1442,7 @@ static void test_resolves_to_string() {
    auto test = [](const std::string & name, const std::string & schema_str, bool expected) {
        fprintf(stderr, "- %s\n", name.c_str());
        common_schema_info info;
-        auto schema = nlohmann::ordered_json::parse(schema_str);
+        auto schema = common_json::parse(schema_str);
        info.resolve_refs(schema);
        bool result = info.resolves_to_string(schema);
        if (result != expected) {
@@ -1517,7 +1517,7 @@ int main() {

    test_all("C++", [](const TestCase & tc) {
        try {
-            tc.verify(json_schema_to_grammar(nlohmann::ordered_json::parse(tc.schema), true));
+            tc.verify(json_schema_to_grammar(common_json::parse(tc.schema), true));
            tc.verify_status(SUCCESS);
        } catch (const std::invalid_argument & ex) {
            fprintf(stderr, "Error: %s\n", ex.what());
@@ -1531,7 +1531,7 @@ int main() {
        auto run = [](const TestCase & tc) {
            fprintf(stderr, "- %s\n", tc.name.c_str());
            try {
-                tc.verify(json_schema_to_grammar(nlohmann::ordered_json::parse(tc.schema), true));
+                tc.verify(json_schema_to_grammar(common_json::parse(tc.schema), true));
                tc.verify_status(SUCCESS);
            } catch (const std::invalid_argument & ex) {
                fprintf(stderr, "Error: %s\n", ex.what());
@@ -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
@@ -9,7 +9,7 @@
 #include "http.h"
 #include "log.h"

-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <algorithm>
 #include <cstdio>
@@ -55,7 +55,7 @@ static const char * COMMIT = "aaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaaa";
 static void serve_repos(httplib::Server & server) {
    server.Get(R"(/api/models/(.+)/refs)", [](const httplib::Request & req, httplib::Response & res) {
        if (g_repos.count(req.matches[1])) {
-            res.set_content(nlohmann::json{{"branches", {{{"name", "main"}, {"targetCommit", COMMIT}}}}}.dump(),
+            res.set_content(common_json{{"branches", common_json::array({ common_json{{"name", "main"}, {"targetCommit", COMMIT}} })}}.dump(),
                            "application/json");
        } else {
            res.status = 404;
@@ -66,7 +66,7 @@ static void serve_repos(httplib::Server & server) {
            res.status = 404;
            return;
        }
-        auto files = nlohmann::json::array();
+        auto files = common_json::array();
        size_t i = 0;
        for (const auto & p : g_repos[req.matches[1]]) {
            char oid[41];
@@ -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;
 }
@@ -6,8 +6,7 @@
 #include "log.h"
 #include "console.h"

-#define JSON_ASSERT GGML_ASSERT
-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <algorithm>
 #include <cctype>
@@ -16,7 +15,7 @@
 #include <map>
 #include <set>

-using json = nlohmann::ordered_json;
+using json = common_json;

 struct cli_context_impl {
    json messages      = json::array();
@@ -73,7 +72,7 @@ static std::string format_error_message(const json & err) {

 // err is the raw response body of a failed request; it may or may not be JSON
 static std::string format_error_message(const std::string & err) {
-    json parsed = json::parse(err, nullptr, false);
+    json parsed = json::parse_no_throw(err);
    if (!parsed.is_discarded()) {
        return format_error_message(parsed);
    }
@@ -157,7 +156,7 @@ bool cli_context::init() {
            if (!list_and_ask_models()) {
                return false;
            }
-        } catch (const json::parse_error & e) {
+        } catch (const common_json_error & e) {
            ui::show_error(e.what());
            ui::show_message("This might be caused by an incorrect server-base endpoint URL");
            return false;
@@ -364,7 +363,7 @@ bool cli_context::generate_completion(generated_content & content_out, cli_timin
    ui::assistant_turn a;

    std::string err = client.post_sse("/v1/chat/completions", body.dump(), should_stop, [&](const std::string & payload) {
-        json chunk = json::parse(payload, nullptr, false);
+        json chunk = json::parse_no_throw(payload);
        if (chunk.is_discarded()) {
            return;
        }
@@ -33,6 +33,7 @@ int llama_fit_params(int argc, char ** argv) {
    if (!params.fit_params_print) {
        const common_params_fit_status status = common_fit_params(params.model.path.c_str(), &mparams, &cparams,
                params.tensor_split, params.tensor_buft_overrides.data(), params.fit_params_target.data(), params.fit_params_min_ctx,
+                nullptr,
                params.verbosity >= LOG_LEVEL_DEBUG ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR);
        if (status != COMMON_PARAMS_FIT_STATUS_SUCCESS) {
            LOG_ERR("%s: failed to fit CLI arguments to free memory, exiting...\n", __func__);
@@ -2294,6 +2294,7 @@ int llama_bench(int argc, char ** argv) {
                fit_overrides.data(),
                margins.data(),
                inst.fit_min_ctx,
+                nullptr,
                params.verbose ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR);
       }

@@ -137,9 +137,15 @@ struct clip_graph {
            int il,
            ggml_tensor * sinks = nullptr) const;

-    // implementation of the 2D RoPE without adding a new op in ggml
-    // this is not efficient (use double the memory), but works on all backends
-    // TODO: there was a more efficient which relies on ggml_view and ggml_rope_ext_inplace, but the rope inplace does not work well with non-contiguous tensors ; we should fix that and revert back to the original implementation in https://github.com/ggml-org/llama.cpp/pull/13065
+    // implementation of the 2D RoPE using two ggml_rope_ext calls
+    //
+    // unlike GGML_ROPE_TYPE_VISION which forces NEOX ordering, this rotates adjacent pairs (normal ordering)
+    //
+    // example:
+    //  given a single head with size = 8 --> [00000000]
+    //  dims [0, 4) rotate with pos_a, dims [4, 8) rotate with pos_b --> [aaaabbbb]
+    //  interleave_freq = false --> both halves use the same inv_freq set (like GGML_ROPE_TYPE_VISION)
+    //  interleave_freq = true  --> first half uses even inv_freq, second half uses odd inv_freq (used by pixtral)
    ggml_tensor * build_rope_2d(
        ggml_context * ctx0,
        ggml_tensor * cur,
@@ -819,8 +819,6 @@ ggml_tensor * clip_graph::build_attn(
 }

 // implementation of the 2D RoPE without adding a new op in ggml
-// this is not efficient (use double the memory), but works on all backends
-// TODO: there was a more efficient which relies on ggml_view and ggml_rope_ext_inplace, but the rope inplace does not work well with non-contiguous tensors ; we should fix that and revert back to the original implementation in https://github.com/ggml-org/llama.cpp/pull/13065
 ggml_tensor * clip_graph::build_rope_2d(
    ggml_context * ctx0,
    ggml_tensor * cur,
@@ -829,9 +827,7 @@ ggml_tensor * clip_graph::build_rope_2d(
    const float freq_base,
    const bool interleave_freq
 ) {
-    const int64_t n_dim  = cur->ne[0];
-    const int64_t n_head = cur->ne[1];
-    const int64_t n_pos  = cur->ne[2];
+    const int64_t n_dim = cur->ne[0];

    // for example, if we have cur tensor of shape (n_dim=8, n_head, n_pos)
    // we will have a list of 4 inv_freq: 1e-0, 1e-1, 1e-2, 1e-3
@@ -845,46 +841,30 @@ ggml_tensor * clip_graph::build_rope_2d(
                                ? std::pow(freq_base, (float)-2/n_dim)
                                : 1.0;

-    // first half
-    ggml_tensor * first;
-    {
-        first = ggml_view_3d(ctx0, cur,
-            n_dim/2, n_head, n_pos,
-            cur->nb[1],
-            cur->nb[2],
-            0);
-        first = ggml_rope_ext(
-            ctx0,
-            first,
-            pos_a,      // positions
-            nullptr,    // freq factors
-            n_dim/2,    // n_dims
-            0, 0, freq_base,
-            1.0f, 0.0f, 1.0f, 0.0f, 0.0f
-        );
-    }
+    // first half, dims [0, n_dim/2)
+    cur = ggml_rope_ext(
+        ctx0,
+        cur,
+        pos_a,      // positions
+        nullptr,    // freq factors
+        n_dim/2,    // n_dims
+        0, 0, freq_base,
+        1.0f, 0.0f, 1.0f, 0.0f, 0.0f
+    );

-    // second half
-    ggml_tensor * second;
-    {
-        second = ggml_view_3d(ctx0, cur,
-            n_dim/2, n_head, n_pos,
-            cur->nb[1],
-            cur->nb[2],
-            n_dim/2 * ggml_element_size(cur));
-        second = ggml_rope_ext(
-            ctx0,
-            second,
-            pos_b,      // positions
-            nullptr,    // freq factors
-            n_dim/2,    // n_dims
-            0, 0, freq_base,
-            freq_scale_odd,
-            0.0f, 1.0f, 0.0f, 0.0f
-        );
-    }
+    // second half, dims [n_dim/2, n_dim)
+    cur = ggml_rope_ext(
+        ctx0,
+        cur,
+        pos_b,      // positions
+        nullptr,    // freq factors
+        n_dim/2,    // n_dims
+        0, 0, freq_base,
+        freq_scale_odd,
+        0.0f, 1.0f, 0.0f, 0.0f
+    );
+    cur = ggml_rope_set_offset(cur, n_dim/2);

-    cur = ggml_concat(ctx0, first, second, 0);
    return cur;
 }

@@ -44,51 +44,31 @@ ggml_cgraph * clip_graph_gemma4v::build() {

    // similar to build_rope_2d, but use neox ordering
    auto add_pos = [&](ggml_tensor * cur, const clip_layer &) {
-        const int64_t n_dim  = cur->ne[0];
-        const int64_t n_head = cur->ne[1];
-        const int64_t n_pos  = cur->ne[2];
+        const int64_t n_dim = cur->ne[0];

-        // first half
-        ggml_tensor * first;
-        {
-            first = ggml_view_4d(ctx0, cur,
-                n_dim/2, n_head, n_pos, n_batch,
-                cur->nb[1],
-                cur->nb[2],
-                cur->nb[3],
-                0);
-            first = ggml_rope_ext(
-                ctx0,
-                first,
-                pos_x,      // positions
-                nullptr,    // freq factors
-                n_dim/2,    // n_dims
-                GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
-                1.0f, 0.0f, 1.0f, 0.0f, 0.0f
-            );
-        }
+        // first half, dims [0, n_dim/2)
+        cur = ggml_rope_ext(
+            ctx0,
+            cur,
+            pos_x,      // positions
+            nullptr,    // freq factors
+            n_dim/2,    // n_dims
+            GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
+            1.0f, 0.0f, 1.0f, 0.0f, 0.0f
+        );

-        // second half
-        ggml_tensor * second;
-        {
-            second = ggml_view_4d(ctx0, cur,
-                n_dim/2, n_head, n_pos, n_batch,
-                cur->nb[1],
-                cur->nb[2],
-                cur->nb[3],
-                n_dim/2 * ggml_element_size(cur));
-            second = ggml_rope_ext(
-                ctx0,
-                second,
-                pos_y,      // positions
-                nullptr,    // freq factors
-                n_dim/2,    // n_dims
-                GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
-                1.0f, 0.0f, 1.0f, 0.0f, 0.0f
-            );
-        }
+        // second half, dims [n_dim/2, n_dim)
+        cur = ggml_rope_ext(
+            ctx0,
+            cur,
+            pos_y,      // positions
+            nullptr,    // freq factors
+            n_dim/2,    // n_dims
+            GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta,
+            1.0f, 0.0f, 1.0f, 0.0f, 0.0f
+        );
+        cur = ggml_rope_set_offset(cur, n_dim/2);

-        cur = ggml_concat(ctx0, first, second, 0);
        return cur;
    };

@@ -2,30 +2,22 @@

 ggml_tensor * clip_graph_minimax_m3::apply_rope(
        ggml_tensor * x, ggml_tensor * pos_h, ggml_tensor * pos_w) {
-    const int64_t Hn  = x->ne[1];
-    const int64_t P   = x->ne[2];
-    const size_t  es  = ggml_element_size(x);
-    const int     dh  = (int) x->ne[0];
-    const int     axd = 2 * ((2 * (dh / 2) / 3) / 2);
+    const int dh  = (int) x->ne[0];
+    const int axd = 2 * ((2 * (dh / 2) / 3) / 2);

-    GGML_ASSERT(x->nb[0] == es);
    GGML_ASSERT(3 * axd <= dh);

    const float th  = hparams.rope_theta;

    // layout of x is [t, h, w, pad]
    // t is unrotated, h and w are rotated, pad is unrotated
-    // note: everything from n_dims onward untouched, so w and pad are rotated in one call.
-    auto sl = [&](int off, int n) {
-        return ggml_cont(ctx0, ggml_view_3d(ctx0, x, n, Hn, P, x->nb[1], x->nb[2], (size_t) off * es));
-    };
-    ggml_tensor * t = sl(0,       axd);
-    ggml_tensor * h = sl(axd,     axd);
-    ggml_tensor * w = sl(2 * axd, dh - 2 * axd); // w + pad
+    x = ggml_rope_ext(ctx0, x, pos_h, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
+    x = ggml_rope_set_offset(x, axd);

-    h = ggml_rope_ext(ctx0, h, pos_h, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
-    w = ggml_rope_ext(ctx0, w, pos_w, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
-    return ggml_concat(ctx0, ggml_concat(ctx0, t, h, 0), w, 0);
+    x = ggml_rope_ext(ctx0, x, pos_w, nullptr, axd, GGML_ROPE_TYPE_NEOX, 0, th, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f);
+    x = ggml_rope_set_offset(x, 2 * axd);
+
+    return x;
 }

 ggml_cgraph * clip_graph_minimax_m3::build() {
@@ -5,7 +5,7 @@
 #include "gguf.h"
 #include "jinja/runtime.h"
 #include "log.h"
-#include "nlohmann/json.hpp"
+#include "json.h"
 #include "peg-parser.h"

 #include <fstream>
@@ -15,7 +15,7 @@
 #include <sstream>
 #include <string>

-using json = nlohmann::ordered_json;
+using json = common_json;

 enum class output_mode {
    ANALYSIS,  // Only output analysis results (default)
@@ -11,9 +11,9 @@
 #include <vector>
 #include <algorithm>

-#include "nlohmann/json.hpp"
+#include "json.h"

-using json = nlohmann::ordered_json;
+using json = common_json;

 // ANSI color codes - using 256-color palette for brighter colors (all bold)
 #define ANSI_RESET       "\033[0m"
@@ -153,7 +153,7 @@ json server_chat_convert_responses_to_chatcmpl(const json & response_body) {
                        prev_msg["content"] = json::array();
                    }
                    auto & prev_content = prev_msg["content"];
-                    prev_content.insert(prev_content.end(), chatcmpl_content.begin(), chatcmpl_content.end());
+                    prev_content.insert(chatcmpl_content);
                } else {
                    item.erase("status");
                    item.erase("type");
@@ -6,9 +6,7 @@
 #include "server-common.h"
 #include "server-http.h"

-#include <nlohmann/json_fwd.hpp>
-
-using json = nlohmann::ordered_json;
+#include "json.h"

 // Convert OpenAI Responses API format to OpenAI Chat Completions API format
 json server_chat_convert_responses_to_chatcmpl(const json & body);
@@ -1540,7 +1540,7 @@ std::vector<llama_token_data> get_token_probabilities(llama_context * ctx, int i
 }

 std::string safe_json_to_str(const json & data) {
-    return data.dump(-1, ' ', false, json::error_handler_t::replace);
+    return data.dump_safe();
 }

 // TODO: reuse llama_detokenize
@@ -6,8 +6,7 @@
 #include "chat.h"
 #include "mtmd.h"

-#define JSON_ASSERT GGML_ASSERT
-#include <nlohmann/json.hpp>
+#include "json.h"

 #include <atomic>
 #include <chrono>
@@ -19,7 +18,7 @@
 #include <string>
 #include <vector>

-using json = nlohmann::ordered_json;
+using json = common_json;

 #define SLT_DBG(slot, fmt, ...) LOG_DBG("slot %12.*s: id %2d | task %d | " fmt, 12, __func__, (slot).id, ((slot).task ? (slot).task->id : -1), __VA_ARGS__)
 #define SLT_TRC(slot, fmt, ...) LOG_TRC("slot %12.*s: id %2d | task %d | " fmt, 12, __func__, (slot).id, ((slot).task ? (slot).task->id : -1), __VA_ARGS__)
@@ -42,9 +41,9 @@ static T json_value(const json & body, const std::string & key, const T & defaul
    // Fallback null to default value
    if (body.contains(key) && !body.at(key).is_null()) {
        try {
-            return body.at(key);
-        } catch (NLOHMANN_JSON_NAMESPACE::detail::type_error const & err) {
-            LOG_WRN("Wrong type supplied for parameter '%s'. Expected '%s', using default value: %s\n", key.c_str(), json(default_value).type_name(), err.what());
+            return body.at(key).get<T>();
+        } catch (const common_json_error & err) {
+            LOG_WRN("Wrong type supplied for parameter '%s', using default value: %s\n", key.c_str(), err.what());
            return default_value;
        }
    } else {
@@ -35,8 +35,6 @@
 #include <windows.h>
 #endif

-using json = nlohmann::ordered_json;
-
 constexpr int HTTP_POLLING_SECONDS = 1;

 static common_speculative_output_limits server_output_limits(const common_params & params) {
@@ -657,14 +655,14 @@ struct server_slot {
            res["n_prompt_tokens_processed"] = stats.n_prompt_processed;
            res["n_prompt_tokens_cache"]     = stats.n_prompt_cached;
            res["params"] = ptask->params.to_json(only_metrics);
-            res["next_token"] = {
+            res["next_token"] = json::array({
                {
                    {"has_next_token", has_next_token},
                    {"has_new_line",   has_new_line},
                    {"n_remain",       n_remaining()},
                    {"n_decoded",      stats.n_gen},
                }
-            };
+            });

            if (!only_metrics) {
                res["prompt"] = ptask->tokens.detokenize(ctx_tgt, true);
@@ -1040,62 +1038,7 @@ private:
            }
        }

-        // optionally reserve VRAM for the draft / MTP context before fitting the target model
-        if (params_base.fit_params) {
-            if (has_spec) {
-                // MTP draft context lives on the target model, only context+compute are new
-                bool measure_model_bytes = has_draft;
-
-                common_params params_dft = common_base_params_to_speculative(params_base);
-
-                auto mparams_dft = common_model_params_to_llama(params_dft);
-                auto cparams_dft = common_context_params_to_llama(params_dft);
-                if (spec_mtp) {
-                    cparams_dft.ctx_type = LLAMA_CONTEXT_TYPE_MTP;
-                }
-                cparams_dft.n_rs_seq = 0;
-
-                std::vector<ggml_backend_dev_t> devs;
-                uint32_t hp_ngl = 0;
-                uint32_t hp_nct = 0;
-                uint32_t hp_nex = 0;
-                try {
-                    auto dmd = common_get_device_memory_data(
-                        params_dft.model.path.c_str(), &mparams_dft, &cparams_dft,
-                        devs, hp_ngl, hp_nct, hp_nex, GGML_LOG_LEVEL_ERROR);
-
-                    GGML_ASSERT(!params_base.fit_params_target.empty());
-                    size_t total = 0;
-
-                    std::vector<ggml_backend_dev_t> tgt_devices = params.devices;
-
-                    if (tgt_devices.empty()) {
-                        for(size_t i = 0; i < ggml_backend_dev_count(); ++i) {
-                           tgt_devices.push_back(ggml_backend_dev_get(i));
-                        }
-                    }
-
-                    for (size_t j = 0; j < devs.size(); ++j) {
-                        const size_t bytes = (measure_model_bytes ? dmd[j].model : 0) + dmd[j].context + dmd[j].compute;
-                        total += bytes;
-                        for (size_t i = 0; i < tgt_devices.size(); i++) {
-                            if (tgt_devices[i] == devs[j]) {
-                                SRV_DBG("[spec] adding %.2f MiB to fit_params_target for device %s\n",
-                                        bytes / (1024.0 * 1024.0), ggml_backend_dev_name(devs[j]));
-                                params_base.fit_params_target[i] += bytes;
-                                break;
-                            }
-                        }
-                    }
-                    SRV_TRC("[spec] estimated memory usage of %s is %.2f MiB\n",
-                            has_draft ? "draft model" : "MTP context",
-                            total / (1024.0 * 1024.0));
-                } catch (const std::exception & e) {
-                    SRV_WRN("[spec] failed to measure %s memory: %s\n",
-                            has_draft ? "draft model" : "MTP context", e.what());
-                }
-            }
-        }
+        // note: the draft / MTP context is fitted together with the target model, see common_fit_extra_model

        // attach a progress callback
        {
@@ -4220,7 +4163,8 @@ std::unique_ptr<server_res_generator> server_routes::handle_completions_impl(
        // tasks.reserve(inputs.size()); // TODO: this is inaccurate due to child tasks

        // message delimiters for checkpointing
-        auto delimiters = common_chat_msg_delimiters_parse(json_value(data, "message_delimiters", json::array()));
+        json delims = json_value(data, "message_delimiters", json::array());
+        auto delimiters = common_chat_msg_delimiters_parse(delims);
        delimiters.tokenize(ctx_server.vocab);

        for (size_t i = 0; i < inputs.size(); i++) {
@@ -4483,8 +4427,8 @@ static json get_res_model_info(const server_context_meta & meta) {
 static json get_res_models(const server_context_meta & meta) {
    // note: do NOT use ctx_server here, otherwise it's not possible to use this during sleep

-    return {
-        {"models", {
+    return json{
+        {"models", json::array({
            {
                {"name",  meta.model_name},
                {"model", meta.model_name},
@@ -4493,23 +4437,23 @@ static json get_res_models(const server_context_meta & meta) {
                {"digest", ""}, // dummy value, llama.cpp does not support managing model file's hash
                {"type", "model"},
                {"description", ""},
-                {"tags", {""}},
-                {"capabilities", meta.has_mtmd ? json({"completion","multimodal"}) : json({"completion"})},
+                {"tags", json::array({""})},
+                {"capabilities", meta.has_mtmd ? json::array({"completion","multimodal"}) : json::array({"completion"})},
                {"parameters", ""},
                {"details", {
                    {"parent_model", ""},
                    {"format", "gguf"},
                    {"family", ""},
-                    {"families", {""}},
+                    {"families", json::array({""})},
                    {"parameter_size", ""},
                    {"quantization_level", ""}
                }}
            }
-        }},
+        })},
        {"object", "list"},
-        {"data", {
+        {"data", json::array({
            get_res_model_info(meta),
-        }}
+        })}
    };
 }

@@ -5045,7 +4989,7 @@ void server_routes::init_routes() {

        std::string content;
        if (body.count("tokens") != 0) {
-            const llama_tokens tokens = body.at("tokens");
+            const llama_tokens tokens = body.at("tokens").get<llama_tokens>();
            content = tokens_to_str(ctx_server.vocab, tokens);
        }

@@ -5352,7 +5296,7 @@ std::unique_ptr<server_res_generator> server_routes::handle_embeddings_impl(cons

    int embd_normalize = params.embd_normalize;
    if (body.count("embd_normalize") != 0) {
-        embd_normalize = body.at("embd_normalize");
+        embd_normalize = body.at("embd_normalize").get<int>();
        if (meta->pooling_type == LLAMA_POOLING_TYPE_NONE) {
            SRV_DBG("embd_normalize is not supported by pooling type %d, ignoring it\n", meta->pooling_type);
        }
@@ -4,7 +4,7 @@
 #include "server-task.h"
 #include "server-queue.h"

-#include <nlohmann/json_fwd.hpp>
+#include "json.h"

 #include <cstddef>
 #include <memory>
@@ -2462,7 +2462,7 @@ server_http_proxy::server_http_proxy(
    bool has_files = !files.empty();

    if (has_files) {
-        json form_fields = json::parse(body, nullptr, false);
+        json form_fields = json::parse_no_throw(body);
        if (!form_fields.is_discarded()) {
            auto boundary = generate_multipart_boundary();
            effective_body = build_multipart_body(form_fields, files, boundary);
@@ -503,7 +503,7 @@ std::vector<std::unique_ptr<field>> make_llama_cmpl_schema(const common_params &
        ->set_handler([&](field_eval_context & ctx, const json & data) {
            const auto & samplers = data.at("samplers");
            if (samplers.is_array()) {
-                ctx.params.sampling.samplers = common_sampler_types_from_names(samplers);
+                ctx.params.sampling.samplers = common_sampler_types_from_names(samplers.get<std::vector<std::string>>());
            } else if (samplers.is_string()) {
                ctx.params.sampling.samplers = common_sampler_types_from_chars(samplers.get<std::string>());
            }
@@ -580,8 +580,7 @@ static void handle_with_catch(const char * name, std::function<void()> func) {

 // treat a null value as absent so clients can send null to request the server default
 static bool has_value(const json & data, const char * n) {
-    auto it = data.find(n);
-    return it != data.end() && !it->is_null();
+    return data.contains(n) && !data.at(n).is_null();
 }

 template <typename T>
@@ -12,8 +12,6 @@

 #include <sstream>

-using json = nlohmann::ordered_json;
-
 //
 // task_params
 //
@@ -304,7 +302,7 @@ json completion_token_output::probs_vector_to_json(const std::vector<completion_
 }

 float completion_token_output::logarithm(float x) {
-    // nlohmann::json converts -inf to null, so we need to prevent that
+    // the JSON library converts -inf to null, so we need to prevent that
    return x == 0.0f ? std::numeric_limits<float>::lowest() : std::log(x);
 }

@@ -407,7 +405,7 @@ json server_task_result_cmpl_final::to_json_oaicompat() {
        res["__verbose"] = to_json_non_oaicompat();
    }
    if (stats.is_set()) {
-        res.push_back({"timings", stats.to_json()});
+        res["timings"] = stats.to_json();
    }

    return res;
@@ -455,7 +453,7 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat() {
        res["__verbose"] = to_json_non_oaicompat();
    }
    if (stats.is_set()) {
-        res.push_back({"timings", stats.to_json()});
+        res["timings"] = stats.to_json();
    }

    return res;
@@ -516,7 +514,7 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat_stream() {
    }

    if (stats.is_set()) {
-        deltas.back().push_back({"timings", stats.to_json()});
+        deltas.back()["timings"] = stats.to_json();
    }

    // extra fields for debugging purposes
@@ -709,7 +707,7 @@ json server_task_result_cmpl_final::to_json_oaicompat_resp_stream() {
    });

    if (stats.is_set()) {
-        server_sent_events.back().at("data").push_back({"timings", stats.to_json()});
+        server_sent_events.back().at("data")["timings"] = stats.to_json();
    }

    return server_sent_events;
@@ -1061,10 +1059,10 @@ json server_task_result_cmpl_partial::to_json_non_oaicompat() {
    };
    // populate the timings object when needed (usually for the last response or with timings_per_token enabled)
    if (stats.is_set()) {
-        res.push_back({"timings", stats.to_json()});
+        res["timings"] = stats.to_json();
    }
    if (is_progress) {
-        res.push_back({"prompt_progress", progress.to_json()});
+        res["prompt_progress"] = progress.to_json();
    }
    if (!prob_output.probs.empty()) {
        res["completion_probabilities"] = completion_token_output::probs_vector_to_json({prob_output}, post_sampling_probs);
@@ -1101,10 +1099,10 @@ json server_task_result_cmpl_partial::to_json_oaicompat() {
        res["__verbose"] = to_json_non_oaicompat();
    }
    if (stats.is_set()) {
-        res.push_back({"timings", stats.to_json()});
+        res["timings"] = stats.to_json();
    }
    if (is_progress) {
-        res.push_back({"prompt_progress", progress.to_json()});
+        res["prompt_progress"] = progress.to_json();
    }

    return res;
@@ -1155,10 +1153,10 @@ json server_task_result_cmpl_partial::to_json_oaicompat_chat() {
        }

        if (stats.is_set()) {
-            last_json.push_back({"timings", stats.to_json()});
+            last_json["timings"] = stats.to_json();
        }
        if (is_progress) {
-            last_json.push_back({"prompt_progress", progress.to_json()});
+            last_json["prompt_progress"] = progress.to_json();
        }
    }

@@ -1305,10 +1303,10 @@ json server_task_result_cmpl_partial::to_json_oaicompat_resp() {
    if (!events.empty()) {
        json & data = events.back().at("data");
        if (stats.is_set()) {
-            data.push_back({"timings", stats.to_json()});
+            data["timings"] = stats.to_json();
        }
        if (is_progress) {
-            data.push_back({"prompt_progress", progress.to_json()});
+            data["prompt_progress"] = progress.to_json();
        }
    }

@@ -11,7 +11,6 @@
 // TODO: prevent including the whole server-common.h as we only use server_tokens
 #include "server-common.h"

-using json = nlohmann::ordered_json;

 enum server_task_type {
    SERVER_TASK_TYPE_COMPLETION,
@@ -2156,7 +2156,7 @@ void server_tools::setup(const std::vector<std::string> & enabled_tools,
                res->status = 200;
                res->data   = safe_json_to_str(result);
            }
-        } catch (const json::exception & e) {
+        } catch (const common_json_error & e) {
            res->status = 400;
            res->data   = safe_json_to_str(format_error_response(e.what(), ERROR_TYPE_INVALID_REQUEST));
        } catch (const std::invalid_argument & e) {
@@ -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)}
@@ -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}
@@ -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;

@@ -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}
@@ -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)
 			);
 		}
 	}
@@ -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}
@@ -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>
@@ -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>
@@ -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>
@@ -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>
@@ -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.
@@ -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>
@@ -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>
@@ -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)
 *
@@ -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"
@@ -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,
@@ -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>
@@ -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;
@@ -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';
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Aman Gupta	b0539c43ed	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>	2026-08-23 13:57:49 +03:00
Gaurav Garg	d3371929bb	[Tensor parallel] Fix meta tensor split state propagation (#27574 ) * ggml : fix meta tensor split state propagation * Add test-llama-archs to CI	2026-08-23 18:49:12 +08:00
Aleksander Grygier	8144f3192e	ui: Chat Conversation Tabbed navigation (#27263 ) * ui : add browser-style conversation tabs store Track open conversation tabs in order, persisted to localStorage and pruned against the loaded conversation list on init. The chat layout syncs the route's tab on every navigation, so any way of reaching a conversation opens a tab for it. * ui : add temporary new-chat tabs New-chat tabs are unsaved conversations carrying a temporary id used directly as the route (#/chat/<id>). They live in memory and are only persisted to the database - keeping the same id so the route and tab stay stable - when the first message is sent. Deleting one drops it without confirmation, and deleting conversations now closes their tabs. * ui : render conversation tab bar in chat layout Desktop-only tab bar above the chat screen, one tab per open conversation or new-chat tab. The active tab follows the route id; clicking navigates, middle-click or the close button closes (switching to the left neighbor), and a trailing + starts a new chat. Tabs appear only on chat-id routes; the bare #/ new-chat view has none. The bare route stays put unless a prompt/model deep-link routes it to a new-chat tab. * ui : route new-chat entry points through tabs The sidebar New chat item, Cmd+Shift+O, the search page and the arrow-key fallback now open a new-chat tab instead of navigating to the ?new_chat URL, which is removed. New chat is no longer a special route but a tab like any other conversation. * ui : track sidebar expanded state in a shared ui store Move the desktop sidebar expanded/collapsed state out of deviceStore into a dedicated uiStore so the chat tab bar can react to it. Assisted-by: pi * chat : add opt-in conversation tabs setting Add a Display setting that turns browser-style conversation tabs on or off, enabled by default. Assisted-by: pi * chat : add browser-style conversation tabs with a new-chat screen Track open conversations as tabs above the chat, one per open chat, plus a single New chat tab for the bare `#/` route. New chat is just the `#/` screen - no temporary conversations - and its tab is dropped when navigating away. Sending the first message creates a real conversation and opens a tab for it. Assisted-by: pi * chat : turn tab bar into a horizontally scrollable carousel Make the tab bar a horizontally scrollable carousel with edge scroll buttons and active-tab centering, and align its styling with the sidebar. Assisted-by: pi * chat : restyle the scroll-to-bottom button to match tab styling Assisted-by: pi * chat : add close-tab keyboard shortcut Assisted-by: pi * chat : soften tab bar fade and dim inactive tabs Assisted-by: pi * feat: Add stop button to tabs * refactor: Componentize * ui : fix carousel scrollability detection Observe the content wrapper as well as the container, since adding overflowing items does not change the container's own box size. Also expose an onScrollableChange callback. Assisted-by: pi * ui : add unified ScrollCarousel component Single carousel component with top/center variants, gap and scroll options, and hover-revealed chevrons. Rename the HorizontalScrollCarousel accessibility story accordingly. Assisted-by: pi * ui : migrate carousels to ScrollCarousel Switch the settings mobile header, attachments list, thumbnail strip, and MCP resources to the unified component, and drop HorizontalScrollCarousel. Assisted-by: pi * ui : improve chat tabs carousel UX Scroll newly added tabs into view, fade overflowing tabs at the edges, and hide the New chat button while a new-chat tab is open. Assisted-by: pi * refactor: Naming * chat : add keyboard shortcut to jump between conversation tabs Shift+Cmd/Ctrl+Left/Right cycles the open tabs, mirroring the existing Shift+Cmd/Ctrl+Up/Down conversation navigation. Assisted-by: pi * chat : make the whole tab item act as a link The full tab is now a link instead of only the inner label button, while the stop and close buttons stay interactive by swallowing their clicks. Assisted-by: pi * chat : adjust tab bar width and use a shared offset variable Widen the tab bar for the expanded sidebar and rename the tab bar height variable to --chat-tabs-offset with a smaller value so the chat screen min-height accounts for the overlay without overshooting. Assisted-by: pi * chat : account for the tab bar offset in the assistant min-height Subtract the tab bar offset when it is shown so the last assistant message does not overflow the available viewport space. Assisted-by: pi * refactor: Post-review fixes * ui : restore deep links on the chat start page - handle ?model selection, with ?load=true eager router loading - ?q now creates a conversation, sends the prompt, and clears the params - show the not-available-model dialog for unknown models - never block mount on the conversation list Assisted-by: pi * ui : fix tab item link nesting and centralize tab constants - the tab anchor covers the whole item while stop/close stay siblings, so interactive elements are never nested inside the anchor - cmd/ctrl/middle clicks are left to the browser (new window) - extract the tab labels, the active-tab data attribute, and the sidebar-offset max widths into constants Assisted-by: pi * ui : tidy scroll carousel hook and keep mobile header arrows on - drop the dead scrollLeft/scrollRight helpers and the unused onScrollableChange/scrollBy props - init the carousel once instead of inside a derived - restore items-start on the center variant - always show the settings header arrows on touch Assisted-by: pi * ui : keep the new-chat tab across reloads and fall back on close - the new-chat sentinel is no longer pruned on init, so reloading on the bare new-chat route keeps the tab the user is on - closing the active conversation falls back to the new-chat screen when Conversation tabs are off Assisted-by: pi * ui : don't block startup on the conversation list - prune persisted tabs after the list loads in the background instead of awaiting it during init - openNewChat now returns void; its return value was never read Assisted-by: pi * ui: fix routing nits * chore: Update doc comments * refactor: Mark fire-and-forget openNewChat calls as `void` * chat: fix the deep-linked prompt, the tab width and the tab shortcuts The chat start page creates the conversation and hands the prompt over to the chat route, which still sees it in the query string. Sending it on both sides queues the second copy as a pending message, which shows up as a stray user bubble once the answer lands and vanishes on reload since it never reaches the database. The tab bar takes the max width of the collapsed sidebar while it is expanded, and the other way round. The tab list is pruned against a snapshot of the loaded conversations, so a conversation created while that list is still loading loses its tab even though the route just opened it. The active tab then falls out of the list and the cycling shortcut jumps to an edge on every keypress instead of moving one tab over. Tabs synced from the route are kept as they are, only the persisted ones are pruned. The rich chat input claims ctrl or alt with shift and an arrow for its badge-aware word jump, which now belongs to the tab cycling shortcut. Holding shift hands the key combination over, the plain word jump is unchanged. The close-tab shortcut consumes the event before checking whether the setting is on, and the logo background loses its importance flag. --------- Co-authored-by: Pascal <admin@serveurperso.com>	2026-08-23 10:46:49 +02:00
Alessandro de Oliveira Faria (A.K.A.CABELO)	6657ded4fa	vendor : update subprocess.h (#27409 )	2026-08-23 10:38:29 +03:00
Aman Karki	29ea9412a6	cuda : add POOL_1D support (#27573 ) * cuda : add POOL_1D support * fix: add missing trailing newline for editorconfig compliance	2026-08-23 10:37:32 +03:00
Xuan-Son Nguyen	70adb1b4ce	common: json.h: fix clang lto (#27575 )	2026-08-23 01:11:10 +02:00
Safi Ullah	3f545becce	vulkan : added the PAD_REFLECT_1D operation (#26586 ) * vulkan : added PAD_REFLECT_1D operation Implemented the GGML_OP_PAD_REFLECT_1D operation for the Vulkan backend Changes: - pad_reflect_1d.comp: implemented the GLSL compute shader with reflection logic - vulkan-shaders-gen.cpp: register the shader for SPIR-V compilation - ggml-vulkan.cpp: pushed constants struct, pipeline creation, supports_op, dispatch function, compute switch and debug validation Tested the PAD_REFLECT_1D on Intel Iris Xe (Vulkan 1.4, Mesa 25.2.8): Correctness: PAD_REFLECT_1D(type=f32,ne_a=[512,34,2,1],pad_0=10,pad_1=9) = Pass PAD_REFLECT_1D(type=f32,ne_a=[3000,384,4,1],pad_0=10,pad_1=9) = Pass 2/2 tests passed - All test are passed Performance: ne_a=[512,34,2,1] -> 5.38 us/run, 24.55 GB/s ne_a=[3000,80,1,1] -> 30.09 us/run, 59.62 GB/s ne_a=[3000,384,4,1] -> 158.31 us/run, 54.39 GB/s * Update ggml/src/ggml-vulkan/vulkan-shaders/pad_reflect_1d.comp Co-authored-by: Jeff Bolz <jbolz@nvidia.com> --------- Co-authored-by: Jeff Bolz <jbolz@nvidia.com>	2026-08-22 14:42:20 -05:00
Xuan-Son Nguyen	b21e4de745	mtmd: use ggml_rope_set_offset (#27521 ) * mtmd: use ggml_rope_set_offset * add comment	2026-08-22 16:33:47 +02:00
Xuan-Son Nguyen	d9f918d2d0	common: add json.h abstraction (#27511 ) * add common/json * migrate common * adapt jinja * migrate server * big wip * migrate tests * wip * revert some excessive changes * wip * wip 2 * revert redundant changes * fix server crash * various fixes * fix ci * harden a bit * clean up * rm json-shim * add some comments * rm redundant decl	2026-08-22 16:28:28 +02:00
Xuan-Son Nguyen	2fb989b9e7	fit: also take into account n_streams (#27496 ) * fit: also take into account n_streams * server: make the draft context follow the target context With a non-unified KV cache the target context now holds n_ctx_train tokens per sequence, while the draft context was still created with n_ctx = 0 and fell back to n_ctx_train / n_streams per sequence. A slot filled beyond that point makes the draft batch fail to decode, and the server answers 500 on the request. The draft context now takes its size from the target context, so both hold the same number of tokens per sequence. Contexts that share their cells with the target no longer need the kv_size override. The memory reserved for the draft model before fitting is measured at the largest context the target can take, since the draft context grows with the target and a fixed byte margin cannot express that. * fit: take an optional second model into account Illustrates the alternative discussed on the draft context fix. The memory of a draft or MTP context is currently handed to the fit as a fixed byte margin, which cannot express a memory that grows with the context the fit is still deciding on. common_fit_params now takes an optional second model that shares the devices of the main one. Its context follows the main context and its memory is measured again whenever that context changes, so the reduce path stays exact instead of conservative. A model that cannot be measured on its own, such as a shared cell MTP context, is skipped with a warning and the main model is fitted alone. This drops the reservation block in the server, which no longer has to probe the trained context size of the target to guess an upper bound. --------- Co-authored-by: Pascal <admin@serveurperso.com>	2026-08-22 16:16:06 +02:00
Xuan-Son Nguyen	9fee29e943	arg: remove -no-cnv from cli [no ci] (#27542 ) * arg: remove -no-cnv from cli * clarify about not adding exccesive test cases	2026-08-22 15:53:56 +02:00
Mario Limonciello	e85caa81ea	ci : Restore ROCm job for Ubuntu (#27399 ) * Revert "ci : disable ubuntu-rocm (#26969)" This reverts commit `9558fa44c9`. * ci: set ccache compiler_check=content for ROCm build The ROCm toolchain is pip-installed fresh on every run, so the clang binary's mtime changes each time. With ccache's default compiler_check=mtime that invalidates the whole cache and warm builds only reached ~70% hits. Hash the compiler contents instead so the cache survives toolchain reinstalls. * Update ccache size to 1GB We're waivering with so many architectures built, we need a bigger ccache limit. * merge fix --------- Co-authored-by: Jim Wu <ywu@xilinx.com> Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>	2026-08-22 13:28:30 +03:00
Tiwei Bie	2115b73d8e	model : support DSpark for bailingmoe3 (#27508 )	2026-08-22 12:19:48 +03:00