ggml-cpu : fix unused variables in arm arch (#31180)

test-llama-archs included src/llama-arch.h and src/llama-model-saver.h,
which prevented building on Windows with shared libraries since the
used symbols are not exported from the DLL.

- declare llama_arch_* wrappers around the internal llm_ arch query
  functions in the llama-ext.h staging header, together with
  llama_arch_kv_name() and llama_arch_supports_save()
- slim llama_model_saver down to the private implementation of
  llama_model_save_to_file(): the fixture-writer half used by the test
  moves into tests/test-llama-archs.cpp as a local kv_writer
- replace the roundtrip saver usage with the public
  llama_model_save_to_file() writing to a temp path (also removes the
  tmpfile() fd leak and its Windows workaround)
- move test-llama-archs, test-generate-models, test-llama-context and
  test-fusion out of the NOT WIN32 OR NOT BUILD_SHARED_LIBS guard

Assisted-by: pi:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (2x DGX Spark (RDMA, TP=2))

.github/workflows/build-apple.yml

@@ -107,62 +107,6 @@ jobs:
           # ref: https://github.com/ggml-org/llama.cpp/pull/19802#issuecomment-4013704023
           ctest -L main -E "test-llama-archs|test-llama-context|test-speculative-rejection-sampling" --verbose --timeout 900
 
-  macos-latest-x64:
-    runs-on: macos-15-intel
-
-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
-
-      - name: ccache
-        uses: ggml-org/ccache-action@v1.2.24
-        with:
-          restore: false
-          save: false
-
-      - name: ccache-buckets-restore
-        uses: ./.github/actions/ccache-buckets
-        env:
-          HF_TOKEN: ${{ secrets.HF_TOKEN_CI }}
-        with:
-          key: apple-x64
-          folder: llama.cpp
-          hf_bucket: ggml-org/cache
-
-      - name: Build
-        id: cmake_build
-        run: |
-          sysctl -a
-          # Metal is disabled due to intermittent failures with Github runners not having a GPU:
-          # https://github.com/ggml-org/llama.cpp/actions/runs/8635935781/job/23674807267#step:5:2313
-          cmake -B build \
-            -DCMAKE_BUILD_RPATH="@loader_path" \
-            -DLLAMA_FATAL_WARNINGS=ON \
-            -DLLAMA_BUILD_BORINGSSL=ON \
-            -DGGML_METAL=OFF \
-            -DGGML_RPC=ON \
-            -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3
-          time cmake --build build --config Release -j $(sysctl -n hw.logicalcpu)
-
-      - name: ccache-buckets-save
-        if: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
-        uses: ./.github/actions/ccache-buckets
-        env:
-          HF_TOKEN: ${{ secrets.HF_TOKEN_CACHE_OUTPUT }}
-        with:
-          key: apple-x64
-          folder: llama.cpp
-          evict-old-files: 1d
-          hf_bucket: ggml-org/cache
-          save: true
-
-      - name: Test
-        id: cmake_test
-        run: |
-          cd build
-          ctest -L main --verbose --timeout 900
-
   macos-latest-ios-xcode:
     runs-on: macos-latest
 

ci/run.sh

@@ -460,10 +460,10 @@ function gg_run_qwen3_0_6b {
     #       https://github.com/ggml-org/llama.cpp/actions/runs/39639985732
     if [ -z ${GG_BUILD_WEBGPU} ]; then
         (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa off --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa on  --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
         (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa off                ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
-        (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa on                 ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
     fi
+    (time ./bin/test-llama-context --model ${model_q4_0} -ngl 10 -c 1024 -fa on  --no-op-offload) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
+    (time ./bin/test-llama-context --model ${model_q4_0} -ngl 99 -c 1024 -fa on                 ) 2>&1 | tee -a $OUT/${ci}-test-llama-context.log
 
     check_ppl "f16"  "$(cat $OUT/${ci}-tg-f16.log  | grep "^\[1\]")"
     if [ -z ${GG_BUILD_NO_BF16} ]; then

common/common.cpp

@@ -742,50 +742,56 @@ static inline bool glob_class_match(const char c, const char * pattern, const ch
 }
 
 // simple glob: * matches non-/ chars, ** matches anything including /, [] matches character class
-static inline bool glob_match(const char * pattern, const char * str) {
-    if (*pattern == '\0') {
-        return *str == '\0';
-    }
-    if (pattern[0] == '*' && pattern[1] == '*') {
-        const char * p = pattern + 2;
-        if (glob_match(p, str)) return true;
-        if (*str != '\0') return glob_match(pattern, str + 1);
+// only wildcards recurse, too many wildcards: no match
+static bool glob_match(const char * pattern, const char * str, int depth = 0) {
+    if (depth >= COMMON_RECURSIVE_LIMIT) {
         return false;
     }
-    if (*pattern == '*') {
-        const char * p = pattern + 1;
-        for (; *str != '\0' && *str != '/'; str++) {
-            if (glob_match(p, str)) return true;
-        }
-        return glob_match(p, str);
-    }
-    if (*pattern == '?' && *str != '\0' && *str != '/') {
-        return glob_match(pattern + 1, str + 1);
-    }
-    if (*pattern == '[') {
-        const char * class_end = pattern + 1;
-        // If first character after '[' is ']' or '-', treat it as literal
-        if (*class_end == ']' || *class_end == '-') {
-            class_end++;
-        }
-        while (*class_end != '\0' && *class_end != ']') {
-            class_end++;
-        }
-        if (*class_end == ']') {
-            if (*str == '\0') return false;
-            bool matched = glob_class_match(*str, pattern + 1, class_end);
-            return matched && glob_match(class_end + 1, str + 1);
-        } else {
-            if (*str == '[') {
-                return glob_match(pattern + 1, str + 1);
+    while (*pattern != '\0') {
+        if (pattern[0] == '*' && pattern[1] == '*') {
+            const char * p = pattern + 2;
+            for (;; str++) {
+                if (glob_match(p, str, depth + 1)) return true;
+                if (*str == '\0') return false;
             }
-            return false;
         }
+        if (*pattern == '*') {
+            const char * p = pattern + 1;
+            for (; *str != '\0' && *str != '/'; str++) {
+                if (glob_match(p, str, depth + 1)) return true;
+            }
+            return glob_match(p, str, depth + 1);
+        }
+        if (*pattern == '?' && *str != '\0' && *str != '/') {
+            pattern++;
+            str++;
+            continue;
+        }
+        if (*pattern == '[') {
+            const char * class_end = pattern + 1;
+            // If first character after '[' is ']' or '-', treat it as literal
+            if (*class_end == ']' || *class_end == '-') {
+                class_end++;
+            }
+            while (*class_end != '\0' && *class_end != ']') {
+                class_end++;
+            }
+            if (*class_end == ']') {
+                if (*str == '\0' || !glob_class_match(*str, pattern + 1, class_end)) return false;
+                pattern = class_end + 1;
+                str++;
+                continue;
+            }
+            if (*str != '[') return false;
+            pattern++;
+            str++;
+            continue;
+        }
+        if (*pattern != *str) return false;
+        pattern++;
+        str++;
     }
-    if (*pattern == *str) {
-        return glob_match(pattern + 1, str + 1);
-    }
-    return false;
+    return *str == '\0';
 }
 
 bool glob_match(const std::string & pattern, const std::string & str) {

common/common.h

@@ -31,6 +31,8 @@
 #define DIRECTORY_SEPARATOR '/'
 #endif // _WIN32
 
+#define COMMON_RECURSIVE_LIMIT 128
+
 #define COM_DBG(fmt, ...) LOG_DBG("cmn  %12.*s: " fmt, 12, __func__, __VA_ARGS__)
 #define COM_TRC(fmt, ...) LOG_TRC("cmn  %12.*s: " fmt, 12, __func__, __VA_ARGS__)
 #define COM_INF(fmt, ...) LOG_INF("cmn  %12.*s: " fmt, 12, __func__, __VA_ARGS__)

common/jinja/parser.cpp

@@ -21,6 +21,17 @@ static bool is_type(const statement_ptr & ptr) {
 class parser {
     const std::vector<token> & tokens;
     size_t current = 0;
+    int depth = 0;
+
+    struct depth_guard {
+        parser & p;
+        depth_guard(parser & p) : p(p) {
+            if (++p.depth > JINJA_RECURSIVE_LIMIT) {
+                throw parser_exception("Parser Error: template is nested more than " + std::to_string(JINJA_RECURSIVE_LIMIT) + " levels deep", p.source, p.peek().pos);
+            }
+        }
+        ~depth_guard() { p.depth--; }
+    };
 
     std::string source; // for error reporting
 
@@ -94,6 +105,7 @@ private:
     }
 
     statement_ptr parse_any() {
+        depth_guard guard(*this);
         size_t start_pos = current;
         switch (peek().t) {
             case token::comment:
@@ -254,6 +266,7 @@ private:
             size_t pos0 = current;
             ++current; // consume {%
             ++current; // consume 'elif'
+            depth_guard guard(*this);
             alternate.push_back(parse_if_statement(pos0)); // nested If
         } else if (is_statement({"else"})) {
             ++current; // consume {%
@@ -330,6 +343,7 @@ private:
     }
 
     statement_ptr parse_if_expression() {
+        depth_guard guard(*this);
         auto a = parse_logical_or_expression();
         if (is_identifier("if")) {
             // Ternary expression
@@ -373,6 +387,7 @@ private:
     statement_ptr parse_logical_negation_expression() {
         // Try parse unary operators
         if (is_identifier("not")) {
+            depth_guard guard(*this);
             size_t start_pos = current;
             auto op = next();
             return mk_stmt<unary_expression>(start_pos, op, parse_logical_negation_expression());
@@ -458,6 +473,7 @@ private:
 
     statement_ptr parse_unary_expression() {
         if (is(token::unary_operator)) {
+            depth_guard guard(*this);
             size_t start_pos = current;
             auto op = next();
             return mk_stmt<unary_expression>(start_pos, op, parse_unary_expression());
@@ -477,9 +493,11 @@ private:
         size_t start_pos = current;
         auto expr = mk_stmt<call_expression>(start_pos, std::move(callee), parse_args());
         auto member = parse_member_expression(std::move(expr)); // foo.x().y
-        return is(token::open_paren)
-            ? parse_call_expression(std::move(member)) // foo.x()()
-            : std::move(member);
+        if (!is(token::open_paren)) {
+            return member;
+        }
+        depth_guard guard(*this);
+        return parse_call_expression(std::move(member)); // foo.x()()
     }
 
     statements parse_args() {

common/jinja/runtime.cpp

@@ -3,6 +3,7 @@
 #include "value.h"
 #include "utils.h"
 
+#include <exception>
 #include <string>
 #include <vector>
 #include <memory>
@@ -52,32 +53,39 @@ static void ensure_key_type_allowed(const value & val) {
 
 // execute with error handling
 value statement::execute(context & ctx) const {
+    std::exception_ptr pending;
+    std::string message;
     try {
         return execute_impl(ctx);
     } catch (const continue_statement::signal & /* ex */) {
-        throw;
+        pending = std::current_exception();
     } catch (const break_statement::signal & /* ex */) {
-        throw;
+        pending = std::current_exception();
     } catch (const rethrown_exception & /* ex */) {
-        throw;
+        pending = std::current_exception();
     } catch (const not_implemented_exception & /* ex */) {
-        throw;
+        pending = std::current_exception();
     } catch (const std::exception & e) {
         const std::string & source = *ctx.src;
         if (source.empty()) {
             std::ostringstream oss;
             oss << "\nError executing " << type() << " at position " << pos << ": " << e.what();
-            throw rethrown_exception(oss.str());
+            message = oss.str();
         } else {
             std::ostringstream oss;
             oss << "\n------------\n";
             oss << "While executing " << type() << " at " << get_line_col(source, pos) << " in source:\n";
             oss << peak_source(source, pos) << "\n";
             oss << "Error: " << e.what();
-            // throw as another exception to avoid repeated formatting
-            throw rethrown_exception(oss.str());
+            message = oss.str();
         }
     }
+    // throw outside the catch: MSVC keeps the unwound frames on the stack while a catch block runs
+    if (pending) {
+        std::rethrow_exception(pending);
+    }
+    // throw as another exception to avoid repeated formatting
+    throw rethrown_exception(message);
 }
 
 value identifier::execute_impl(context & ctx) const {
@@ -747,7 +755,11 @@ value macro_statement::execute_impl(context & ctx) const {
     std::string name = cast_stmt<identifier>(this->name)->val;
 
     const func_handler func = [this, name](const func_args & args) -> value {
+        if (args.ctx.macro_depth >= JINJA_RECURSIVE_LIMIT) {
+            throw std::runtime_error("macro calls are nested more than " + std::to_string(JINJA_RECURSIVE_LIMIT) + " levels deep");
+        }
         context macro_ctx(args.ctx); // new scope for macro execution
+        macro_ctx.macro_depth = args.ctx.macro_depth + 1;
 
         bind_parameters(name, this->args, args, macro_ctx);
 
@@ -778,7 +790,11 @@ value call_statement::execute_impl(context & ctx) const {
     context caller_ctx(ctx); // new scope for caller execution
 
     const func_handler func = [this, caller_ctx = std::move(caller_ctx)](const func_args & args) -> value {
+        if (args.ctx.macro_depth >= JINJA_RECURSIVE_LIMIT) {
+            throw std::runtime_error("macro calls are nested more than " + std::to_string(JINJA_RECURSIVE_LIMIT) + " levels deep");
+        }
         context block_ctx(caller_ctx); // new scope for block execution
+        block_ctx.macro_depth = args.ctx.macro_depth + 1;
 
         bind_parameters("caller", this->caller_args, args, block_ctx);
 

common/jinja/runtime.h

@@ -58,6 +58,8 @@ struct context {
 
     bool is_get_stats = false; // whether to collect stats
 
+    int macro_depth = 0;
+
     visitor_fn visitor;
 
     // src is optional, used for error reporting
@@ -78,6 +80,7 @@ struct context {
         src(parent.src),
         current_time(parent.current_time),
         is_get_stats(parent.is_get_stats),
+        macro_depth(parent.macro_depth),
         // inherit variables (for example, when entering a new scope)
         env(mk_val<value_object>(*parent.env)) {}
 
@@ -710,7 +713,7 @@ struct rethrown_exception : public std::exception {
 
 //////////////////////
 
-static void gather_string_parts_recursive(const value & val, value_string & parts) {
+static void gather_string_parts_recursive(const value & val, value_string & parts, int depth = 0) {
     // TODO: probably allow print value_none as "None" string? currently this breaks some templates
     if (is_val<value_string>(val)) {
         const auto & str_val = cast_val<value_string>(val)->val_str;
@@ -718,10 +721,12 @@ static void gather_string_parts_recursive(const value & val, value_string & part
     } else if (is_val<value_int>(val) || is_val<value_float>(val) || is_val<value_bool>(val)) {
         std::string str_val = val->as_string().str();
         parts->val_str.append(str_val);
+    } else if (is_val<value_array>(val) && depth >= JINJA_RECURSIVE_LIMIT) {
+        throw std::runtime_error("value is nested more than " + std::to_string(JINJA_RECURSIVE_LIMIT) + " levels deep");
     } else if (is_val<value_array>(val)) {
         auto items = cast_val<value_array>(val)->as_array();
         for (const auto & item : items) {
-            gather_string_parts_recursive(item, parts);
+            gather_string_parts_recursive(item, parts, depth + 1);
         }
     }
 }

common/jinja/utils.h

@@ -6,6 +6,8 @@
 #include <cstdint>
 #include <cstring>
 
+#define JINJA_RECURSIVE_LIMIT 128
+
 namespace jinja {
 
 static void string_replace_all(std::string & s, const std::string & search, const std::string & replace) {

common/jinja/value.cpp

@@ -1441,8 +1441,6 @@ const func_builtins & value_undefined_t::get_builtins() const {
 
 
 // the runtime walks values with recursion, a deeper input overflows the stack
-static constexpr int FROM_JSON_MAX_DEPTH = 128;
-
 static value from_json(const common_json & j, bool mark_input, int depth = 0) {
     if (j.is_null()) {
         return mk_val<value_none>();
@@ -1458,8 +1456,8 @@ static value from_json(const common_json & j, bool mark_input, int depth = 0) {
             str->mark_input();
         }
         return str;
-    } else if ((j.is_array() || j.is_object()) && depth >= FROM_JSON_MAX_DEPTH) {
-        throw std::invalid_argument("JSON is nested more than " + std::to_string(FROM_JSON_MAX_DEPTH) + " levels deep");
+    } else if ((j.is_array() || j.is_object()) && depth >= JINJA_RECURSIVE_LIMIT) {
+        throw std::invalid_argument("JSON is nested more than " + std::to_string(JINJA_RECURSIVE_LIMIT) + " levels deep");
     } else if (j.is_array()) {
         auto arr = mk_val<value_array>();
         for (const auto & item : j) {
@@ -1664,6 +1662,8 @@ static void value_to_json_internal(std::ostringstream & oss, const value & val,
             }
         }
         oss << "\"";
+    } else if ((is_val<value_array>(val) || is_val<value_object>(val)) && curr_lvl >= JINJA_RECURSIVE_LIMIT) {
+        throw std::runtime_error("value is nested more than " + std::to_string(JINJA_RECURSIVE_LIMIT) + " levels deep");
     } else if (is_val<value_array>(val)) {
         const auto & arr = val->as_array();
         oss << "[";
@@ -1729,17 +1729,20 @@ std::string value_to_string_repr(const value & val) {
 }
 
 // stats utility
-void value_t::stats_t::mark_used(value & val, bool deep) {
+void value_t::stats_t::mark_used(value & val, bool deep, int depth) {
     val->stats.used = true;
+    if (deep && (is_val<value_array>(val) || is_val<value_object>(val)) && depth >= JINJA_RECURSIVE_LIMIT) {
+        throw std::runtime_error("value is nested more than " + std::to_string(JINJA_RECURSIVE_LIMIT) + " levels deep");
+    }
     if (deep) {
         if (is_val<value_array>(val)) {
             for (auto & item : val->val_arr) {
-                mark_used(item, deep);
+                mark_used(item, deep, depth + 1);
             }
         } else if (is_val<value_object>(val)) {
             for (auto & pair : val->val_obj) {
-                mark_used(pair.first, deep);
-                mark_used(pair.second, deep);
+                mark_used(pair.first, deep, depth + 1);
+                mark_used(pair.second, deep, depth + 1);
             }
         }
     }

common/jinja/value.h

@@ -119,7 +119,7 @@ struct value_t {
         // ops can be builtin calls or operators: "array_access", "object_access"
         std::set<std::string> ops;
         // utility to recursively mark value and its children as used
-        static void mark_used(value & val, bool deep = false);
+        static void mark_used(value & val, bool deep = false, int depth = 0);
     } stats;
 
     value_t() = default;

common/json-schema-to-grammar.cpp

@@ -272,15 +272,12 @@ static bool is_reserved_name(const std::string & name) {
     return RESERVED_NAMES.find(name) != RESERVED_NAMES.end();
 }
 
-static std::regex INVALID_RULE_CHARS_RE("[^a-zA-Z0-9-]+");
 static std::regex GRAMMAR_LITERAL_ESCAPE_RE("[\r\n\"\\\\]");
 static std::regex GRAMMAR_RANGE_LITERAL_ESCAPE_RE("[\r\n\"\\]\\-\\\\]");
 static std::unordered_map<char, std::string> GRAMMAR_LITERAL_ESCAPES = {
     {'\r', "\\r"}, {'\n', "\\n"}, {'"', "\\\""}, {'-', "\\-"}, {']', "\\]"}, {'\\', "\\\\"}
 };
 
-static const int MAX_PATTERN_DEPTH = 100;
-
 static std::unordered_set<char> NON_LITERAL_SET = {'|', '.', '(', ')', '[', ']', '{', '}', '*', '+', '?', '^', '$'};
 static std::unordered_set<char> ESCAPED_IN_REGEXPS_BUT_NOT_IN_LITERALS = {'^', '$', '.', '[', ']', '(', ')', '|', '{', '}', '*', '+', '?'};
 
@@ -312,6 +309,22 @@ static std::string format_literal(const std::string & literal) {
 
 std::string gbnf_format_literal(const std::string & literal) { return format_literal(literal); }
 
+// same as std::regex_replace with [^a-zA-Z0-9-]+ -> "-", no std::regex: libstdc++ recurses per char
+std::string gbnf_format_rule_name(const std::string & name) {
+    std::string out;
+    bool in_invalid = false;
+    for (char c : name) {
+        if ((c >= 'a' && c <= 'z') || (c >= 'A' && c <= 'Z') || (c >= '0' && c <= '9') || c == '-') {
+            out += c;
+            in_invalid = false;
+        } else if (!in_invalid) {
+            out += '-';
+            in_invalid = true;
+        }
+    }
+    return out;
+}
+
 static size_t gbnf_escape_length(const std::string & pattern, size_t pos) {
     if (pos + 1 >= pattern.length() || pattern[pos] != '\\') {
         return 0;
@@ -347,6 +360,17 @@ private:
     std::unordered_set<std::string> _refs_being_resolved;
     std::vector<std::string> _errors;
     std::vector<std::string> _warnings;
+    int _depth = 0;
+
+    struct depth_guard {
+        int & depth;
+        depth_guard(int & depth) : depth(depth) {
+            if (++depth > COMMON_RECURSIVE_LIMIT) {
+                throw std::invalid_argument("JSON schema is nested more than " + std::to_string(COMMON_RECURSIVE_LIMIT) + " levels deep");
+            }
+        }
+        ~depth_guard() { depth--; }
+    };
 
     template <typename T>
     static const T & as(const common_chat_schema & node) {
@@ -354,7 +378,7 @@ private:
     }
 
     std::string _add_rule(const std::string & name, const std::string & rule) {
-        std::string esc_name = regex_replace(name, INVALID_RULE_CHARS_RE, "-");
+        std::string esc_name = gbnf_format_rule_name(name);
         if (_rules.find(esc_name) == _rules.end() || _rules[esc_name] == rule) {
             _rules[esc_name] = rule;
             return esc_name;
@@ -482,7 +506,7 @@ private:
                         }
                     }
                     paren_depth++;
-                    if (paren_depth > MAX_PATTERN_DEPTH) {
+                    if (paren_depth > COMMON_RECURSIVE_LIMIT) {
                         throw unsupported_pattern("pattern nesting too deep");
                     }
                     seq.emplace_back("(" + to_rule(transform()) + ")", false);
@@ -644,7 +668,11 @@ private:
         std::string char_rule = _add_primitive("char", PRIMITIVE_RULES.at("char"));
         std::ostringstream out;
         out << "[\"] ( ";
-        std::function<void(size_t)> visit = [&](size_t idx) {
+        // one level per char of the longest name
+        std::function<void(size_t, int)> visit = [&](size_t idx, int depth) {
+            if (depth >= COMMON_RECURSIVE_LIMIT) {
+                throw std::invalid_argument("JSON schema property name is longer than " + std::to_string(COMMON_RECURSIVE_LIMIT) + " chars");
+            }
             const auto & node = trie.nodes[idx];
             std::string rejects;
             auto first = true;
@@ -659,7 +687,7 @@ private:
                 out << "[" << c << "]";
                 if (!trie.nodes[child].children.empty()) {
                     out << " (";
-                    visit(child);
+                    visit(child, depth + 1);
                     out << ")";
                 } else {
                     out << " " << char_rule << "+";
@@ -669,7 +697,7 @@ private:
                 out << " | [^\"" << rejects << "] " << char_rule << "*";
             }
         };
-        visit(0);
+        visit(0, 0);
 
         out << " )";
         if (trie.nodes[0].pattern < 0) {
@@ -682,8 +710,7 @@ private:
     std::string _resolve_ref(const common_chat_schema_ref & schema) {
         auto it = schema.ref.find('#');
         std::string ref_fragment = it != std::string::npos ? schema.ref.substr(it + 1) : schema.ref;
-        static const std::regex nonalphanumeric_regex(R"([^a-zA-Z0-9-]+)");
-        std::string ref_name = "ref" + std::regex_replace(ref_fragment, nonalphanumeric_regex, "-");
+        std::string ref_name = "ref" + gbnf_format_rule_name(ref_fragment);
         if (_rules.find(ref_name) == _rules.end() && _refs_being_resolved.find(schema.ref) == _refs_being_resolved.end()) {
             if (!schema.target) {
                 _errors.push_back("Unresolved $ref " + schema.ref);
@@ -833,10 +860,13 @@ public:
         std::unordered_set<std::string> required;
         std::vector<std::pair<std::string, const common_chat_schema *>> properties;
         std::map<std::string, size_t> enum_values;
-        std::function<void(const common_chat_schema &, bool)> add_component = [&](const common_chat_schema & comp, bool is_required) {
+        std::function<void(const common_chat_schema &, bool, int)> add_component = [&](const common_chat_schema & comp, bool is_required, int depth) {
+            if (depth >= COMMON_RECURSIVE_LIMIT) {
+                throw std::invalid_argument("JSON schema is nested more than " + std::to_string(COMMON_RECURSIVE_LIMIT) + " levels deep");
+            }
             if (comp.kind() == common_chat_schema::KIND_REF) {
                 if (const auto * target = as<common_chat_schema_ref>(comp).target) {
-                    add_component(*target, is_required);
+                    add_component(*target, is_required, depth + 1);
                 }
             } else if (comp.kind() == common_chat_schema::KIND_OBJECT) {
                 for (const auto & prop : as<common_chat_schema_object>(comp).properties) {
@@ -854,10 +884,10 @@ public:
         for (const auto & child : schema.children) {
             if (child->kind() == common_chat_schema::KIND_ANY_OF) {
                 for (const auto & alt : as<common_chat_schema_any_of>(*child).children) {
-                    add_component(*alt, false);
+                    add_component(*alt, false, 0);
                 }
             } else {
-                add_component(*child, true);
+                add_component(*child, true, 0);
             }
         }
         if (!enum_values.empty()) {
@@ -875,6 +905,8 @@ public:
     }
 
     std::string visit(const common_chat_schema & schema, const std::string & name) {
+        // the $ref targets are visited too, a $ref chain is not bounded by the builder
+        depth_guard guard(_depth);
         std::string rule_name = is_reserved_name(name) ? name + "-" : name.empty() ? "root" : name;
         std::string sub_name  = name + (name.empty() ? "" : "-");
 

common/json-schema-to-grammar.h

@@ -19,5 +19,6 @@ struct common_grammar_options {
 };
 
 std::string gbnf_format_literal(const std::string & literal);
+std::string gbnf_format_rule_name(const std::string & name);
 
 std::string build_grammar(const std::function<void(const common_grammar_builder &)> & cb, const common_grammar_options & options = {});

common/json-schema.cpp

@@ -9,6 +9,18 @@
 #include <utility>
 #include <vector>
 
+namespace {
+struct depth_guard {
+    int & depth;
+    depth_guard(int & depth) : depth(depth) {
+        if (++depth > COMMON_RECURSIVE_LIMIT) {
+            throw std::invalid_argument("JSON schema is nested more than " + std::to_string(COMMON_RECURSIVE_LIMIT) + " levels deep");
+        }
+    }
+    ~depth_guard() { depth--; }
+};
+} // namespace
+
 class common_chat_schema_builder {
     const common_json &      root_;
     common_chat_schema_document & doc_;
@@ -224,18 +236,11 @@ class common_chat_schema_builder {
         return node;
     }
 
-    // the builder and the code that walks the built schema recurse once per level, a deeper schema overflows the stack
-    static constexpr int max_depth = 128;
-    int                  depth_    = 0;
+    int depth_ = 0;
 
     common_chat_schema_ptr build_node(const common_json & schema, const std::string & path) {
-        if (depth_ >= max_depth) {
-            throw std::invalid_argument("JSON schema is nested more than " + std::to_string(max_depth) + " levels deep");
-        }
-        depth_++;
-        auto node = build_node_impl(schema, path);
-        depth_--;
-        return node;
+        depth_guard guard(depth_);
+        return build_node_impl(schema, path);
     }
 
     common_chat_schema_ptr build_node_impl(const common_json & schema, const std::string & path) {
@@ -379,7 +384,9 @@ static common_chat_schema::value_type json_type(const common_json & value) {
     return common_chat_schema::TYPE_OBJECT;
 }
 
-static common_chat_schema::type_set value_types_impl(const common_chat_schema & s, std::unordered_set<const common_chat_schema *> & visited) {
+// depth counts the $ref, anyOf and allOf hops, a $ref chain is not bounded by the builder
+static common_chat_schema::type_set value_types_impl(const common_chat_schema & s, std::unordered_set<const common_chat_schema *> & visited, int & depth) {
+    depth_guard guard(depth);
     switch (s.kind()) {
         case common_chat_schema::KIND_ANY:
             return common_chat_schema::type_set::all();
@@ -413,21 +420,21 @@ static common_chat_schema::type_set value_types_impl(const common_chat_schema &
                 // a cycle contributes no type, to be safe
                 return {};
             }
-            auto types = value_types_impl(*target, visited);
+            auto types = value_types_impl(*target, visited, depth);
             visited.erase(target);
             return types;
         }
         case common_chat_schema::KIND_ANY_OF: {
             common_chat_schema::type_set types;
             for (const auto & child : static_cast<const common_chat_schema_any_of &>(s).children) {
-                types |= value_types_impl(*child, visited);
+                types |= value_types_impl(*child, visited, depth);
             }
             return types;
         }
         case common_chat_schema::KIND_ALL_OF: {
             auto types = common_chat_schema::type_set::all();
             for (const auto & child : static_cast<const common_chat_schema_all_of &>(s).children) {
-                types &= value_types_impl(*child, visited);
+                types &= value_types_impl(*child, visited, depth);
             }
             return types;
         }
@@ -437,10 +444,12 @@ static common_chat_schema::type_set value_types_impl(const common_chat_schema &
 
 common_chat_schema::type_set common_chat_schema::value_types() const {
     std::unordered_set<const common_chat_schema *> visited;
-    return value_types_impl(*this, visited);
+    int depth = 0;
+    return value_types_impl(*this, visited, depth);
 }
 
-static bool may_be_string_impl(const common_chat_schema & s, std::unordered_set<const common_chat_schema *> & visited) {
+static bool may_be_string_impl(const common_chat_schema & s, std::unordered_set<const common_chat_schema *> & visited, int & depth) {
+    depth_guard guard(depth);
     switch (s.kind()) {
         case common_chat_schema::KIND_STRING:
             return true;
@@ -459,13 +468,13 @@ static bool may_be_string_impl(const common_chat_schema & s, std::unordered_set<
             if (!target || !visited.insert(target).second) {
                 return false;
             }
-            bool result = may_be_string_impl(*target, visited);
+            bool result = may_be_string_impl(*target, visited, depth);
             visited.erase(target);
             return result;
         }
         case common_chat_schema::KIND_ANY_OF:
             for (const auto & child : static_cast<const common_chat_schema_any_of &>(s).children) {
-                if (may_be_string_impl(*child, visited)) {
+                if (may_be_string_impl(*child, visited, depth)) {
                     return true;
                 }
             }
@@ -477,7 +486,7 @@ static bool may_be_string_impl(const common_chat_schema & s, std::unordered_set<
                 if (child->kind() == common_chat_schema::KIND_ANY) {
                     continue;
                 }
-                if (!may_be_string_impl(*child, visited)) {
+                if (!may_be_string_impl(*child, visited, depth)) {
                     return false;
                 }
                 any_string = true;
@@ -491,7 +500,8 @@ static bool may_be_string_impl(const common_chat_schema & s, std::unordered_set<
 
 bool common_chat_schema::may_be_string() const {
     std::unordered_set<const common_chat_schema *> visited;
-    return may_be_string_impl(*this, visited);
+    int depth = 0;
+    return may_be_string_impl(*this, visited, depth);
 }
 
 const char * common_chat_schema::kind_name(node_kind kind) {

common/peg-parser.cpp

@@ -10,10 +10,12 @@
 #include <initializer_list>
 #include <map>
 #include <memory>
-#include <regex>
 #include <set>
 #include <stdexcept>
 
+// The parser can run with a lower limit than COMMON_RECURSIVE_LIMIT
+static constexpr int RULE_RECURSIVE_LIMIT = 64;
+
 // Trick to catch missing branches
 template <typename T>
 inline constexpr bool is_always_false_v = false;
@@ -186,13 +188,17 @@ std::string common_peg_ast_node::sanitized_text() const {
 }
 
 void common_peg_ast_arena::visit(common_peg_ast_id id, const common_peg_ast_visitor & visitor) const {
-    if (id == COMMON_PEG_INVALID_AST_ID) {
-        return;
-    }
-    const auto & node = get(id);
-    visitor(node);
-    for (const auto & child : node.children) {
-        visit(child, visitor);
+    // no recursion: the AST is several nodes deep per level of the input
+    std::vector<common_peg_ast_id> stack = { id };
+    while (!stack.empty()) {
+        const auto cur = stack.back();
+        stack.pop_back();
+        if (cur == COMMON_PEG_INVALID_AST_ID) {
+            continue;
+        }
+        const auto & node = get(cur);
+        visitor(node);
+        stack.insert(stack.end(), node.children.rbegin(), node.children.rend());
     }
 }
 
@@ -736,8 +742,15 @@ struct parser_executor {
     }
 
     common_peg_parse_result operator()(const common_peg_rule_parser & p) {
+        // only rules recurse, a recursive rule is entered once per input level
+        if (ctx.rule_depth[p.child] >= RULE_RECURSIVE_LIMIT) {
+            throw common_peg_recursion_error("input is nested more than " + std::to_string(RULE_RECURSIVE_LIMIT) + " levels deep");
+        }
+
         // Parse the child
+        ctx.rule_depth[p.child]++;
         auto result = arena.parse(p.child, ctx, start_pos);
+        ctx.rule_depth[p.child]--;
 
         if (!result.fail()) {
             std::string_view text;
@@ -819,6 +832,10 @@ common_peg_parse_result common_peg_arena::parse(common_peg_parse_context & ctx,
     if (root_ == COMMON_PEG_INVALID_PARSER_ID) {
         throw std::runtime_error("No root parser set");
     }
+
+    // Initialize rule depth for all parsers
+    ctx.rule_depth.assign(parsers_.size(), 0);
+
     return parse(root_, ctx, start);
 }
 
@@ -1069,11 +1086,6 @@ common_peg_parser operator|(const std::string & str, const common_peg_parser & p
     return operator|(str.c_str(), p);
 }
 
-static std::string rule_name(const std::string & name) {
-    static const std::regex invalid_rule_chars_re("[^a-zA-Z0-9-]+");
-    return std::regex_replace(name, invalid_rule_chars_re, "-");
-}
-
 common_peg_parser_builder::common_peg_parser_builder() {}
 
 common_peg_parser common_peg_parser_builder::sequence(const std::vector<common_peg_parser_id> & parsers) {
@@ -1155,14 +1167,14 @@ common_peg_parser common_peg_parser_builder::schema(const common_peg_parser & p,
 }
 
 common_peg_parser common_peg_parser_builder::rule(const std::string & name, const common_peg_parser & p, bool trigger) {
-    auto clean_name = rule_name(name);
+    auto clean_name = gbnf_format_rule_name(name);
     auto rule_id = arena_.add_parser(common_peg_rule_parser{clean_name, p.id(), trigger});
     arena_.add_rule(clean_name, rule_id);
     return ref(clean_name);
 }
 
 common_peg_parser common_peg_parser_builder::rule(const std::string & name, const std::function<common_peg_parser()> & builder_fn, bool trigger) {
-    auto clean_name = rule_name(name);
+    auto clean_name = gbnf_format_rule_name(name);
     if (arena_.has_rule(clean_name)) {
         return ref(clean_name);
     }

common/peg-parser.h

@@ -6,6 +6,7 @@
 
 #include <memory>
 #include <set>
+#include <stdexcept>
 #include <unordered_map>
 #include <string>
 #include <string_view>
@@ -182,6 +183,10 @@ inline common_peg_parse_flags operator~(common_peg_parse_flags a) {
     return static_cast<common_peg_parse_flags>(~int(a));
 }
 
+struct common_peg_recursion_error : std::invalid_argument {
+    using std::invalid_argument::invalid_argument;
+};
+
 struct common_peg_parse_context {
     std::string input;               // [h,   e,  l,  l,  o,  _,  w,  o,  r,  l,  d]
     std::vector<llama_token> tokens; // [id, -1, -1, -1, -1, id, -1, -1, -1, -1, -1]
@@ -189,6 +194,7 @@ struct common_peg_parse_context {
     common_peg_ast_arena ast;
 
     int parse_depth;
+    std::vector<int> rule_depth; // index: rule child id
 
     common_peg_parse_context(common_peg_parse_flags flags = COMMON_PEG_PARSE_FLAG_NONE)
         : flags(flags), parse_depth(0) {}

conversion/deepseek.py

@@ -803,9 +803,6 @@ class DeepseekV4Model(TextModel):
         self._dsv4_mxfp4_generated = True
         return ()
 
-    def _format_dsv4_tensor_name(self, key: gguf.MODEL_TENSOR, bid: int | None, suffix: str = ".weight") -> str:
-        return self.format_tensor_name(key, bid, suffix)
-
     def _map_dsv4_tensor_name(self, name: str, bid: int | None) -> tuple[gguf.MODEL_TENSOR, str]:
         root_map: dict[str, tuple[gguf.MODEL_TENSOR, str]] = {
             "embed.weight": (gguf.MODEL_TENSOR.TOKEN_EMBD, ".weight"),
@@ -879,17 +876,17 @@ class DeepseekV4Model(TextModel):
 
     def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
         if re.match(r"layers\.\d+\.ffn\.experts\.\d+\.w[123]\.(weight|scale)$", name):
-            return []
+            return
 
         # hash layers route text tokens via tid2eid and image tokens via bias_vl; gate.bias is unused
         if name.endswith(".ffn.gate.bias") and bid is not None and bid < self.hparams["num_hash_layers"]:
-            return []
+            return
 
         tensor_key, suffix = self._map_dsv4_tensor_name(name, bid)
         if tensor_key == gguf.MODEL_TENSOR.FFN_GATE_TID2EID:
-            return []
+            return
 
-        return [(self._format_dsv4_tensor_name(tensor_key, bid, suffix), data_torch)]
+        yield self.format_tensor_name(tensor_key, bid, suffix), data_torch
 
     def tensor_force_quant(self, name: str, new_name: str, bid: int | None, n_dims: int) -> gguf.GGMLQuantizationType | bool:
         del bid  # unused

conversion/deepseek41.py

@@ -145,12 +145,6 @@ class DeepseekV41Model(DeepseekV4Model):
         # num_hidden_layers lives in text_config, which is merged only after indexing; MTP stages are skipped anyway
         return TextModel.index_tensors(self, remote_hf_model_id=remote_hf_model_id)
 
-    @classmethod
-    def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
-        if item[0].startswith("vision."):
-            return None
-        return super().filter_tensors(item)
-
     def set_gguf_parameters(self):
         super().set_gguf_parameters()
         hp = self.hparams
@@ -282,12 +276,13 @@ class DeepseekV41Model(DeepseekV4Model):
 
     def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
         if name.endswith(".engram.k_weight"):
-            return []  # folded into engram_gate
+            return  # folded into engram_gate
         if name.endswith(".engram.q_weight"):
             # Engram gate weight = q_weight * k_weight (bf16 * bf16 is exact in f32)
             k = self.model_tensors[name.removesuffix("q_weight") + "k_weight"]
-            return [(self.format_tensor_name(gguf.MODEL_TENSOR.ENGRAM_GATE, bid), data_torch.float() * k().float())]
-        return super().modify_tensors(data_torch, name, bid)
+            yield self.format_tensor_name(gguf.MODEL_TENSOR.ENGRAM_GATE, bid), data_torch.float() * k().float()
+            return
+        yield from super().modify_tensors(data_torch, name, bid)
 
     def tensor_force_quant(self, name: str, new_name: str, bid: int | None, n_dims: int) -> gguf.GGMLQuantizationType | bool:
         if new_name.endswith((".engram_gate.weight", ".ffn_gate_inp.weight")) or name in self._dsv4_f32_tensors:
@@ -411,13 +406,15 @@ class DeepseekV41DSparkModel(DeepseekV41Model):
 
     def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
         if name.endswith(".ffn.gate.bias_vl"):
-            return []
+            return
         if name in self._DSPARK_ROOT:
-            return [(self.format_tensor_name(self._DSPARK_ROOT[name], None), data_torch)]
+            yield self.format_tensor_name(self._DSPARK_ROOT[name], None), data_torch
+            return
         m = re.match(r"layers\.\d+\.(main_proj\.weight|main_norm\.weight|norm\.weight)$", name)
         if m is not None:
-            return [(self.format_tensor_name(self._DSPARK_ROOT[m.group(1)], None), data_torch)]
-        return super().modify_tensors(data_torch, name, bid)
+            yield self.format_tensor_name(self._DSPARK_ROOT[m.group(1)], None), data_torch
+            return
+        yield from super().modify_tensors(data_torch, name, bid)
 
     def tensor_force_quant(self, name: str, new_name: str, bid: int | None, n_dims: int) -> gguf.GGMLQuantizationType | bool:
         if new_name == "conf_proj.weight":

conversion/qwen4exp.py

@@ -41,6 +41,8 @@ class Qwen4ExpTextModel(_Qwen35MRopeMixin, _LinearAttentionVReorderBase):
         self._ple_shards: dict[int, str] = {}
         self._ple_row_dim: int | None = None
         self._mtp_fc: dict[str, Tensor] = {}
+        # halves of the merged weights, keyed by (bid, merged tensor), see _merge_shared_input
+        self._merge_parts: dict[tuple[int, gguf.MODEL_TENSOR], dict[int, Tensor]] = {}
 
     @classmethod
     def filter_tensors(cls, item):
@@ -131,44 +133,79 @@ class Qwen4ExpTextModel(_Qwen35MRopeMixin, _LinearAttentionVReorderBase):
         # int64 hash constants must stay exact; 1-D tensors force F32, so use KV
         if name.endswith("ple_embedding.layer_multipliers"):
             self._ple_multipliers = [int(x) for x in data_torch.tolist()]
-            return []
+            return
         if name.endswith("ple_embedding.ngram_heads_offsets"):
             self._ple_head_offsets = [int(x) for x in data_torch.tolist()]
-            return []
+            return
         if name.endswith("ple_embedding.ngram_heads_vocab_sizes"):
             self._ple_head_vocab_sizes = [int(x) for x in data_torch.tolist()]
-            return []
+            return
 
         if ".ngram_embedding.shard_" in name:
-            return self._place_ple_shard(data_torch, name)
+            yield from self._place_ple_shard(data_torch, name)
+            return
 
         # eh_proj([e ; h_s]) = fc_embedding(e) + fc_hidden(h_s) for every hc stream s
         if name.endswith((".nextn_fc_embedding.weight", ".nextn_fc_hidden.weight")):
             self._mtp_fc[name.rsplit(".", 2)[1]] = data_torch
             if len(self._mtp_fc) < 2:
-                return []
+                return
             eh = torch.cat([self._mtp_fc.pop("nextn_fc_embedding"), self._mtp_fc.pop("nextn_fc_hidden")], dim=1)
-            return [(self.format_tensor_name(gguf.MODEL_TENSOR.NEXTN_EH_PROJ, bid, ".weight"), eh)]
+            yield (self.format_tensor_name(gguf.MODEL_TENSOR.NEXTN_EH_PROJ, bid, ".weight"), eh)
+            return
 
         # one projection feeds indexer q and k; split it, as minimax-m3 does
         if ".indexer.index_qk_proj.weight" in name:
             n_q = self.hparams["indexer_n_heads"] * self.hparams["indexer_head_dim"]
-            q = data_torch[:n_q]
-            k = data_torch[n_q:]
-            return [
-                (self.format_tensor_name(gguf.MODEL_TENSOR.INDEXER_Q_PROJ, bid, ".weight"), q),
-                (self.format_tensor_name(gguf.MODEL_TENSOR.INDEXER_K_PROJ, bid, ".weight"), k),
-            ]
+            yield (self.format_tensor_name(gguf.MODEL_TENSOR.INDEXER_Q_PROJ, bid, ".weight"), data_torch[:n_q])
+            yield (self.format_tensor_name(gguf.MODEL_TENSOR.INDEXER_K_PROJ, bid, ".weight"), data_torch[n_q:])
+            return
 
         # Gemma zero-centred gammas the inherited norm.weight rule misses
         if name.endswith((".ple.norm_key.weight", ".ple.norm_query.weight", ".ple.norm_conv.weight",
                           ".indexer.q_layernorm.weight", ".indexer.k_layernorm.weight")):
-            return [(self.map_tensor_name(name), data_torch + 1)]
+            yield (self.map_tensor_name(name), data_torch + 1)
+            return
 
         if name.endswith(".ple.conv1d.weight"):
-            return [(self.map_tensor_name(name), data_torch.squeeze())]
+            yield (self.map_tensor_name(name), data_torch.squeeze())
+            return
 
-        return super().modify_tensors(data_torch, name, bid)
+        for new_name, data in super().modify_tensors(data_torch, name, bid):
+            yield from self._merge_shared_input(new_name, data, bid)
+
+    # matmuls that read the same input, written as one weight [a | b] so the runtime does one matmul:
+    # (merged tensor, (a, b), linear-attention layers only)
+    _SHARED_INPUT_MERGES = (
+        (gguf.MODEL_TENSOR.HC_ATTN_DOWN_INJECT, (gguf.MODEL_TENSOR.HC_ATTN_DOWN, gguf.MODEL_TENSOR.HC_ATTN_INJECT), False),
+        (gguf.MODEL_TENSOR.HC_FFN_DOWN_INJECT,  (gguf.MODEL_TENSOR.HC_FFN_DOWN,  gguf.MODEL_TENSOR.HC_FFN_INJECT),  False),
+        (gguf.MODEL_TENSOR.ATTN_QKV_GATE,       (gguf.MODEL_TENSOR.ATTN_QKV,     gguf.MODEL_TENSOR.ATTN_GATE),      True),
+        (gguf.MODEL_TENSOR.SSM_BETA_ALPHA_CAT,  (gguf.MODEL_TENSOR.SSM_BETA,     gguf.MODEL_TENSOR.SSM_ALPHA),      True),
+    )
+
+    # caches one half of a merged weight until the other half arrives, then yields [a | b]
+    def _merge_shared_input(self, new_name: str, data: Tensor, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        # a pre-quantized checkpoint carries per-tensor scales, which one merged weight cannot hold
+        if bid is None or self.hparams.get("quantization_config"):
+            yield (new_name, data)
+            return
+
+        layer_types = self.hparams.get("layer_types") or []
+        is_linear = bid < len(layer_types) and layer_types[bid] == "linear_attention"
+
+        match = next(((merged, parts.index(part)) for merged, parts, linear_only in self._SHARED_INPUT_MERGES
+                      for part in parts
+                      if (is_linear or not linear_only) and self.match_model_tensor_name(new_name, part, bid)), None)
+        if match is None:
+            yield (new_name, data)
+            return
+
+        merged, idx = match
+        halves = self._merge_parts.setdefault((bid, merged), {})
+        halves[idx] = data
+        if len(halves) == 2:
+            del self._merge_parts[(bid, merged)]
+            yield (self.format_tensor_name(merged, bid), torch.cat([halves[0], halves[1]], dim=0))
 
     # the shards concatenate into a tensor of well over 100 GB
     # use LazyChunkedTensor here, a single shard resident at a time
@@ -214,6 +251,8 @@ class Qwen4ExpTextModel(_Qwen35MRopeMixin, _LinearAttentionVReorderBase):
         super().prepare_tensors()
         if self._mtp_fc:
             raise ValueError(f"MTP projection missing its other half: {sorted(self._mtp_fc)}")
+        if self._merge_parts:
+            raise ValueError(f"merged weights missing a half: {sorted((bid, t.name) for bid, t in self._merge_parts)}")
         n_parts = self.hparams.get("split_ngram_parts", 0)
         if self._ple_shards and len(self._ple_shards) != n_parts:
             raise ValueError(

ggml/src/ggml-cpu/arch/arm/quants.c

@@ -287,6 +287,9 @@ void ggml_vec_dot_q2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const voi
 
     sumf = vaddvq_f32(sumv);
 #else
+    UNUSED(nb);
+    UNUSED(x);
+    UNUSED(y);
     ggml_vec_dot_q2_0_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc);
     return;
 #endif

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

@@ -214,52 +214,36 @@ static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device)
 }
 
 #if defined(GGML_USE_HIP)
-static int ggml_cuda_parse_id(char devName[]) {
+static int ggml_cuda_parse_id(const char * devName) {
     // A list of possible Target IDs can be found under the rocclr/clr repo in device.cpp
     // these values are not stable so this is susceptible to breakage
     // https://github.com/ROCm/clr/blob/amd-staging/rocclr/device/device.cpp
-    int archMajor = 0x0;
-    int archMinor = 0x0;
-    int archNum = GGML_CUDA_CC_OFFSET_AMD;
-    int archLen = strlen(devName);
-    char archName[archLen + 1];
+    // Extract the base architecture model ("gfx908:sramecc+:xnack-" -> "908")
+    std::string archName = devName;
+    archName.erase(0, 3);
+    archName = archName.substr(0, archName.find(':'));
 
-    // strip leading 'gfx' while copying into our buffer
-    if (archLen > 3) {
-        strcpy(archName, &devName[3]);
-        archLen -= 3;
-    }
+    const auto hex = [](const std::string & s) { return (int) strtoul(s.c_str(), nullptr, 16); };
 
-    // trim trailing :xnack- or :sramecc- statuses
-    archLen = strcspn(archName, ":");
-    archName[archLen] = '\0';
+    int cc = 0;
 
     // tease out the version information
-    if (archLen > 8) {
+    if (archName.size() > 8) {
         // versions labeled generic use '-' as delimiter
-        // strip the trailing "-generic" then iterate through what remains
-        if ((strstr(archName, "-generic"))) {
-            archName[archLen - 8] = '\0';
-            char * pch;
-            if ((pch = strtok(archName, "-"))) {
-                archMajor = (int)strtoul(pch, 0, 16);
-                if ((pch = strtok(NULL, "-"))) {
-                    archMinor = 0x10 * (int)strtoul(pch, 0, 16);
-                }
+        if (archName.compare(archName.size() - 8, 8, "-generic") == 0) {
+            archName.resize(archName.size() - 8);
+            const size_t sep = archName.find('-');
+            cc = 0x100 * hex(archName.substr(0, sep));
+            if (sep != std::string::npos) {
+                cc += 0x10 * hex(archName.substr(sep + 1));
             }
         }
-    } else if (archLen >= 3) {
-        // last two digits should be the minor * 0x10 + stepping
-        archMinor = (int)strtoul(&archName[archLen - 2], 0, 16);
-        archName[archLen - 2] = '\0';
-
-        // only the major version remains
-        archMajor = (int)strtoul(archName, 0, 16);
+    } else if (archName.size() >= 3) {
+        // ordinary IDs are the version in hex
+        cc = hex(archName);
     }
-    archNum += archMajor * 0x100;
-    archNum += archMinor;
 
-    return archNum;
+    return GGML_CUDA_CC_OFFSET_AMD + cc;
 }
 #endif // defined(GGML_USE_HIP)
 

ggml/src/ggml-metal/kernels/mul_mm.metal

@@ -323,10 +323,10 @@ kernel void kernel_mul_mm(
         // if no bounds checks on the output are needed, we can directly write to device memory
         device float * C = (device float *) dst +
             (r0 + 32*(sgitg &  1)) + \
-            (r1 + 16*(sgitg >> 1)) * args.ne0 + im*args.ne1*args.ne0;
+            (uint64_t) (r1 + 16*(sgitg >> 1))*args.ne0 + (uint64_t) im*args.ne1*args.ne0;
 
         for (short i = 0; i < 8; i++) {
-            simdgroup_store(mc[i], C + 8*(i%4) + 8*args.ne0*(i/4), args.ne0, 0, false);
+            simdgroup_store(mc[i], C + 8*(i%4) + (uint64_t) 8*args.ne0*(i/4), args.ne0, 0, false);
         }
     } else {
         // block is smaller than 64x32, we should avoid writing data outside of the matrix
@@ -342,7 +342,7 @@ kernel void kernel_mul_mm(
 
         if (sgitg == 0) {
             for (int j = tiitg; j < nr1; j += NR1) {
-                device float  * D  = (device float  *) dst + r0 + (r1 + j)*args.ne0 + im*args.ne1*args.ne0;
+                device float  * D  = (device float  *) dst + r0 + (uint64_t) (r1 + j)*args.ne0 + (uint64_t) im*args.ne1*args.ne0;
                 device float4 * D4 = (device float4 *) D;
 
                 threadgroup float  * C  = temp_str + (j*NR0);
@@ -633,7 +633,7 @@ kernel void kernel_mul_mm_id(
     const int id = ids_i32[idb + lr1];
 
     const short i11 = (id % args.ne20) % args.ne11;
-    const short i12 = (id / args.ne20);
+    const int   i12 = (id / args.ne20);
 
     device const block_q * x = (device const block_q *)(src0 + args.nb01*(r0 + lr0) + offset0) + offset1;
 
@@ -866,9 +866,9 @@ kernel void kernel_mul_mm_id(
         const int id = ids_i32[idb + j];
 
         const short ide = id % args.ne20;
-        const short idt = id / args.ne20;
+        const int   idt = id / args.ne20;
 
-        device float  * D  = (device float  *) dst + r0 + ide*args.ne0 + idt*args.ne1*args.ne0;
+        device float  * D  = (device float  *) dst + r0 + (uint64_t) ide*args.ne0 + (uint64_t) idt*args.ne1*args.ne0;
         device float4 * D4 = (device float4 *) D;
 
         threadgroup float  * C  = (threadgroup float  *) shmem + j*NR0;

ggml/src/ggml-webgpu/ggml-webgpu.cpp

@@ -670,6 +670,8 @@ static void ggml_backend_webgpu_buffer_memset(webgpu_global_context & ctx,
                                               uint32_t                value,
                                               size_t                  offset,
                                               size_t                  size) {
+    std::lock_guard<std::recursive_mutex> lock(ctx->mutex);
+
     std::vector<uint32_t>             params  = { (uint32_t) offset, (uint32_t) size, value };
     std::vector<wgpu::BindGroupEntry> entries = { ggml_webgpu_make_bind_group_entry(0, buf, 0, buf.GetSize()) };
     size_t                            bytes_per_wg =

gguf-py/gguf/constants.py

@@ -920,6 +920,10 @@ class MODEL_TENSOR(IntEnum):
     HC_FFN_DOWN          = auto() # qwen4exp
     HC_FFN_UP            = auto() # qwen4exp
     HC_FFN_INJECT        = auto() # qwen4exp
+    HC_ATTN_DOWN_INJECT  = auto() # qwen4exp, [hc_attn_down | hc_attn_inject]
+    HC_FFN_DOWN_INJECT   = auto() # qwen4exp, [hc_ffn_down | hc_ffn_inject]
+    ATTN_QKV_GATE        = auto() # qwen4exp, [attn_qkv | attn_gate]
+    SSM_BETA_ALPHA_CAT   = auto() # qwen4exp, [ssm_beta | ssm_alpha], unlike the grouped qwen3next ssm_ba
     PLE_KEY              = auto() # qwen4exp
     PLE_VALUE            = auto() # qwen4exp
     PLE_NORM_KEY         = auto() # qwen4exp
@@ -1746,6 +1750,10 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
     MODEL_TENSOR.HC_FFN_DOWN:               "blk.{bid}.hc_ffn_down",          # qwen4exp
     MODEL_TENSOR.HC_FFN_UP:                 "blk.{bid}.hc_ffn_up",            # qwen4exp
     MODEL_TENSOR.HC_FFN_INJECT:             "blk.{bid}.hc_ffn_inject",        # qwen4exp
+    MODEL_TENSOR.HC_ATTN_DOWN_INJECT:       "blk.{bid}.hc_attn_down_inject",  # qwen4exp
+    MODEL_TENSOR.HC_FFN_DOWN_INJECT:        "blk.{bid}.hc_ffn_down_inject",   # qwen4exp
+    MODEL_TENSOR.ATTN_QKV_GATE:             "blk.{bid}.attn_qkv_gate",        # qwen4exp
+    MODEL_TENSOR.SSM_BETA_ALPHA_CAT:        "blk.{bid}.ssm_beta_alpha",       # qwen4exp
     MODEL_TENSOR.PLE_KEY:                   "blk.{bid}.ple_key",              # qwen4exp
     MODEL_TENSOR.PLE_VALUE:                 "blk.{bid}.ple_value",            # qwen4exp
     MODEL_TENSOR.PLE_NORM_KEY:              "blk.{bid}.ple_norm_key",         # qwen4exp
@@ -3309,6 +3317,10 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
         MODEL_TENSOR.HC_FFN_DOWN,
         MODEL_TENSOR.HC_FFN_UP,
         MODEL_TENSOR.HC_FFN_INJECT,
+        MODEL_TENSOR.HC_ATTN_DOWN_INJECT,
+        MODEL_TENSOR.HC_FFN_DOWN_INJECT,
+        MODEL_TENSOR.ATTN_QKV_GATE,
+        MODEL_TENSOR.SSM_BETA_ALPHA_CAT,
         # full attention layers: ATTN_Q holds [q|gate] interleaved per head
         MODEL_TENSOR.ATTN_Q,
         MODEL_TENSOR.ATTN_Q_NORM,

include/llama.h

@@ -1564,6 +1564,10 @@ extern "C" {
                                float   penalty_freq,     // must be finite, 0.0 = disabled
                                float   penalty_present); // must be finite, 0.0 = disabled
 
+    /// maximum number of distinct DRY sequence breakers, duplicates not counted
+    /// processing each breaker scans the entire vocabulary, so the init cost is O(num_breakers x n_vocab)
+    #define LLAMA_DRY_MAX_SEQ_BREAKERS 64
+
     ///  @details DRY sampler, designed by p-e-w, as described in: https://github.com/oobabooga/text-generation-webui/pull/5677, porting Koboldcpp implementation authored by pi6am: https://github.com/LostRuins/koboldcpp/pull/982
     LLAMA_API struct llama_sampler * llama_sampler_init_dry(
             const struct llama_vocab *  vocab,

skills/code-review/SKILL.md

@@ -58,6 +58,7 @@ Mandatory on every review; any finding here is **blocking**. Rule of thumb: GGUF
 - **Reused/reserved buffers:** recheck bounds after a buffer is shrunk or reused; watch `reserve()` then index-by-assumed-size, and header fields read before their length is checked.
 - **Server JSON ints:** clamp client-supplied integers (token/discard counts, offsets) to non-negative and an upper bound before they reach index/pointer arithmetic.
 - **RPC-deserialized fields:** treat every field (type/buffer/data/ne/nb/op_params) as hostile - validate before use. Null/zero buffers skipping validation, attacker data pointers, out-of-range type indices, and negative strides sign-extending past a corner-only assert all give arbitrary read/write.
+- **Recursion on untrusted input:** a function that recurses once per level (or per char) of request JSON, model output, a grammar, a template or a regex pattern overflows the stack - the HTTP threads have 512 KB on macOS, so a few hundred levels are enough. Every such function needs its own depth bound, checked in the function itself, not a pre-walk at a caller (the next caller skips it): libcommon uses `COMMON_RECURSIVE_LIMIT` (`common/common.h`), jinja `JINJA_RECURSIVE_LIMIT` (`common/jinja/utils.h`), the grammar `GRAMMAR_RECURSIVE_LIMIT` (`src/llama-grammar.h`). Count the levels of the input, not the frames, and throw `std::invalid_argument` for request input so the server answers 400. Look for recursion the code does not spell out: mutual recursion, recursive `std::function` lambdas, visitor/virtual dispatch, `std::regex` (libstdc++ recurses per char of the input), and per-char helpers that recurse where a loop would do.
 - **Lifetime/UAF:** flag stored raw pointers to caller/temporary storage, cached pointers to buffers a later free releases, async ops whose source may drop before completion, and structures not invalidated on free/realloc. Null-check conditionally-built or "not required" tensors before dereferencing.
 
 ## Approach and design (when a new component/infra is introduced)

src/llama-arch.cpp

@@ -1,5 +1,6 @@
 #include "llama-arch.h"
 
+#include "llama-ext.h"
 #include "llama-impl.h"
 
 #include <map>
@@ -635,6 +636,10 @@ static const std::map<llm_tensor, const char *> LLM_TENSOR_NAMES = {
     { LLM_TENSOR_HC_FFN_DOWN,                            "blk.%d.hc_ffn_down" },
     { LLM_TENSOR_HC_FFN_UP,                              "blk.%d.hc_ffn_up" },
     { LLM_TENSOR_HC_FFN_INJECT,                          "blk.%d.hc_ffn_inject" },
+    { LLM_TENSOR_HC_ATTN_DOWN_INJECT,                    "blk.%d.hc_attn_down_inject" },
+    { LLM_TENSOR_HC_FFN_DOWN_INJECT,                     "blk.%d.hc_ffn_down_inject" },
+    { LLM_TENSOR_ATTN_QKV_GATE,                          "blk.%d.attn_qkv_gate" },
+    { LLM_TENSOR_SSM_BETA_ALPHA_CAT,                     "blk.%d.ssm_beta_alpha" },
     { LLM_TENSOR_PLE_KEY,                                "blk.%d.ple_key" },
     { LLM_TENSOR_PLE_VALUE,                              "blk.%d.ple_value" },
     { LLM_TENSOR_PLE_NORM_KEY,                           "blk.%d.ple_norm_key" },
@@ -907,6 +912,10 @@ static const std::map<llm_tensor, llm_tensor_info> LLM_TENSOR_INFOS = {
     {LLM_TENSOR_HC_FFN_DOWN,                {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
     {LLM_TENSOR_HC_FFN_UP,                  {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
     {LLM_TENSOR_HC_FFN_INJECT,              {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
+    {LLM_TENSOR_HC_ATTN_DOWN_INJECT,        {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
+    {LLM_TENSOR_HC_FFN_DOWN_INJECT,         {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
+    {LLM_TENSOR_ATTN_QKV_GATE,              {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
+    {LLM_TENSOR_SSM_BETA_ALPHA_CAT,         {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
     {LLM_TENSOR_PLE_KEY,                    {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
     {LLM_TENSOR_PLE_VALUE,                  {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
     {LLM_TENSOR_PLE_NORM_KEY,               {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
@@ -1348,3 +1357,43 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) {
             return true;
     }
 }
+
+bool llama_arch_supports_save(llm_arch arch) {
+    switch (arch) {
+        case LLM_ARCH_SENSENOVA_U1:
+        case LLM_ARCH_GEMMA3N:
+        case LLM_ARCH_BITNET:
+        case LLM_ARCH_T5:
+        case LLM_ARCH_APERTUS:
+        case LLM_ARCH_STEP35:
+        case LLM_ARCH_CLEF: // the head tensors are not saved
+        case LLM_ARCH_VLLM_DECISION2:
+            return false;
+        default:
+            return true;
+    }
+}
+
+std::vector<llm_arch> llama_arch_all() {
+    return llm_arch_all();
+}
+
+const char * llama_arch_name(llm_arch arch) {
+    return llm_arch_name(arch);
+}
+
+llm_arch llama_arch_from_string(const std::string & name) {
+    return llm_arch_from_string(name);
+}
+
+std::string llama_arch_kv_name(llm_arch arch, llm_kv key) {
+    return LLM_KV(arch)(key);
+}
+
+bool llama_arch_supports_sm_tensor(llm_arch arch) {
+    return llm_arch_supports_sm_tensor(arch);
+}
+
+bool llama_arch_supports_mixed_batch(llm_arch arch) {
+    return llm_arch_supports_mixed_batch(arch);
+}

src/llama-arch.h

@@ -679,6 +679,10 @@ enum llm_tensor {
     LLM_TENSOR_HC_FFN_DOWN,    // qwen4exp
     LLM_TENSOR_HC_FFN_UP,      // qwen4exp
     LLM_TENSOR_HC_FFN_INJECT,  // qwen4exp
+    LLM_TENSOR_HC_ATTN_DOWN_INJECT, // qwen4exp
+    LLM_TENSOR_HC_FFN_DOWN_INJECT,  // qwen4exp
+    LLM_TENSOR_ATTN_QKV_GATE,       // qwen4exp
+    LLM_TENSOR_SSM_BETA_ALPHA_CAT,  // qwen4exp
     LLM_TENSOR_PLE_KEY,        // qwen4exp
     LLM_TENSOR_PLE_VALUE,      // qwen4exp
     LLM_TENSOR_PLE_NORM_KEY,   // qwen4exp

src/llama-ext.h

@@ -5,11 +5,36 @@
 // try as much as possible to not include this header in the rest of the codebase
 
 #include "llama.h"
+#include "llama-arch.h"
 
 #include <cstdint>
 #include <map>
+#include <string>
 #include <vector>
 
+//
+// llama architecture queries
+//
+
+// all the LLM architectures known to the library
+LLAMA_API std::vector<llm_arch> llama_arch_all();
+
+// the canonical name of an architecture (e.g. "llama"), or "unknown"
+LLAMA_API const char * llama_arch_name(llm_arch arch);
+
+// the architecture with the given canonical name, or LLM_ARCH_UNKNOWN
+LLAMA_API llm_arch llama_arch_from_string(const std::string & name);
+
+// the arch-resolved GGUF key of a KV id, e.g. LLM_KV_CONTEXT_LENGTH -> "llama.context_length"
+LLAMA_API std::string llama_arch_kv_name(llm_arch arch, llm_kv key);
+
+LLAMA_API bool llama_arch_supports_sm_tensor(llm_arch arch);
+
+LLAMA_API bool llama_arch_supports_mixed_batch(llm_arch arch);
+
+// FIXME temporary function for better error messages
+LLAMA_API bool llama_arch_supports_save(llm_arch arch);
+
 // Reserve a new compute graph. It is valid until the next call to llama_graph_reserve.
 LLAMA_API struct ggml_cgraph * llama_graph_reserve(
         struct llama_context * ctx,

src/llama-grammar.cpp

@@ -8,9 +8,11 @@
 #include <algorithm>
 #include <cstdint>
 #include <set>
+#include <numeric>
 #include <stdexcept>
 
 #define MAX_REPETITION_THRESHOLD 2000
+#define GRAMMAR_RECURSIVE_LIMIT 256
 //
 // helpers
 //
@@ -440,13 +442,14 @@ const char * llama_grammar_parser::parse_alternates(
         const char        * src,
         const std::string & rule_name,
         uint32_t            rule_id,
-        bool                is_nested) {
+        bool                is_nested,
+        uint32_t            depth) {
     llama_grammar_rule rule;
-    const char * pos = parse_sequence(src, rule_name, rule, is_nested);
+    const char * pos = parse_sequence(src, rule_name, rule, is_nested, depth);
     while (*pos == '|') {
         rule.push_back({LLAMA_GRETYPE_ALT, 0});
         pos = parse_space(pos + 1, true);
-        pos = parse_sequence(pos, rule_name, rule, is_nested);
+        pos = parse_sequence(pos, rule_name, rule, is_nested, depth);
     }
     rule.push_back({LLAMA_GRETYPE_END, 0});
     add_rule(rule_id, rule);
@@ -457,7 +460,8 @@ const char * llama_grammar_parser::parse_sequence(
         const char         * src,
         const std::string  & rule_name,
         llama_grammar_rule & rule,
-        bool               is_nested) {
+        bool               is_nested,
+        uint32_t           depth) {
     size_t last_sym_start = rule.size();
     const char * pos = src;
     uint64_t n_prev_rules = 1;
@@ -595,11 +599,14 @@ const char * llama_grammar_parser::parse_sequence(
             n_prev_rules = 1;
             rule.push_back({LLAMA_GRETYPE_RULE_REF, ref_rule_id});
         } else if (*pos == '(') { // grouping
+            if (depth >= GRAMMAR_RECURSIVE_LIMIT) {
+                throw std::runtime_error("grammar recursion depth limit exceeded (max " + std::to_string(GRAMMAR_RECURSIVE_LIMIT) + "), please reduce nesting depth");
+            }
             // parse nested alternates into synthesized rule
             pos = parse_space(pos + 1, true);
             uint32_t n_rules_before = symbol_ids.size();
             uint32_t sub_rule_id = generate_symbol_id(rule_name);
-            pos = parse_alternates(pos, rule_name, sub_rule_id, true);
+            pos = parse_alternates(pos, rule_name, sub_rule_id, true, depth + 1);
             n_prev_rules = std::max(1u, (uint32_t)symbol_ids.size() - n_rules_before);
             last_sym_start = rule.size();
             // output reference to synthesized rule
@@ -941,34 +948,19 @@ static void llama_grammar_advance_stack(
     }
 }
 
-static llama_grammar_candidates llama_grammar_reject_candidates(
-        const llama_grammar_rules      & rules,
-        const llama_grammar_stacks     & stacks,
-        const llama_grammar_candidates & candidates) {
-    GGML_ASSERT(!stacks.empty()); // REVIEW
-
-    if (candidates.empty()) {
-        return {};
-    }
-
-    auto rejects = llama_grammar_reject_candidates_for_stack(rules, stacks.front(), candidates);
-
-    for (size_t i = 1, size = stacks.size(); i < size; ++i) {
-        rejects = llama_grammar_reject_candidates_for_stack(rules, stacks[i], rejects);
-    }
-
-    return rejects;
-}
-
 static bool llama_grammar_detect_left_recursion(
         const llama_grammar_rules & rules,
         size_t rule_index,
         std::vector<bool> * rules_visited,
         std::vector<bool> * rules_in_progress,
-        std::vector<bool> * rules_may_be_empty) {
+        std::vector<bool> * rules_may_be_empty,
+        uint32_t depth = 0) {
     if ((*rules_in_progress)[rule_index]) {
         return true;
     }
+    if (depth > GRAMMAR_RECURSIVE_LIMIT) {
+        throw std::runtime_error("grammar rules are chained more than " + std::to_string(GRAMMAR_RECURSIVE_LIMIT) + " levels deep");
+    }
 
     (*rules_in_progress)[rule_index] = true;
 
@@ -994,7 +986,7 @@ static bool llama_grammar_detect_left_recursion(
     bool recurse_into_nonterminal = true;
     for (size_t i = 0; i < rule.size(); i++) {
         if (rule[i].type == LLAMA_GRETYPE_RULE_REF && recurse_into_nonterminal) {
-            if (llama_grammar_detect_left_recursion(rules, (size_t)rule[i].value, rules_visited, rules_in_progress, rules_may_be_empty)) {
+            if (llama_grammar_detect_left_recursion(rules, (size_t)rule[i].value, rules_visited, rules_in_progress, rules_may_be_empty, depth + 1)) {
                 return true;
             }
             if (!((*rules_may_be_empty)[(size_t)rule[i].value])) {
@@ -1058,56 +1050,63 @@ void llama_grammar_accept(struct llama_grammar * grammar, uint32_t chr) {
     grammar->stacks = std::move(stacks_new);
 }
 
-llama_grammar_candidates llama_grammar_reject_candidates_for_stack(
+// checks the code point at depth of each candidate (an index into the input) against the stack
+// returns true when some match: they go to next, and the stacks that follow to next_stacks
+static bool llama_grammar_reject_step(
         const llama_grammar_rules      & rules,
         const llama_grammar_stack      & stack,
-        const llama_grammar_candidates & candidates) {
-
-    llama_grammar_candidates rejects;
-    rejects.reserve(candidates.size());
+        const llama_grammar_candidates & candidates,
+        size_t                           depth,
+        const std::vector<uint32_t>    & cur,
+              std::vector<uint32_t>    & rejects,
+              std::vector<uint32_t>    & next,
+              llama_grammar_stacks     & next_stacks) {
+    rejects.reserve(cur.size());
 
     if (stack.empty()) {
-        for (const auto & tok : candidates) {
-            if (*tok.code_points != 0 || tok.partial_utf8.n_remain != 0) {
-                rejects.push_back(tok);
+        for (const uint32_t i : cur) {
+            const auto & tok = candidates[i];
+            if (tok.code_points[depth] != 0 || tok.partial_utf8.n_remain != 0) {
+                rejects.push_back(i);
             }
         }
-        return rejects;
+        return false;
     }
 
     const llama_grammar_element * stack_pos = stack.back();
 
     // if the top of the stack is a token rule, then we only need to check the token id
     if (stack_pos->type == LLAMA_GRETYPE_TOKEN || stack_pos->type == LLAMA_GRETYPE_TOKEN_NOT) {
-        for (const auto & tok : candidates) {
-            if (*tok.code_points == 0) {
+        for (const uint32_t i : cur) {
+            const auto & tok = candidates[i];
+            if (tok.code_points[depth] == 0) {
                 // reached the end of a token consumed by char rules, reject iff it ended
                 // in a partial response
                 if (tok.partial_utf8.n_remain != 0) {
-                    rejects.push_back(tok);
+                    rejects.push_back(i);
                 }
             } else if (!llama_grammar_match_token(stack_pos, tok.id)) {
-                rejects.push_back(tok);
+                rejects.push_back(i);
             }
         }
-        return rejects;
+        return false;
     }
 
-    llama_grammar_candidates next_candidates;
-    next_candidates.reserve(candidates.size());
+    next.reserve(cur.size());
 
-    for (const auto & tok : candidates) {
-        if (*tok.code_points == 0) {
+    for (const uint32_t i : cur) {
+        const auto & tok = candidates[i];
+        if (tok.code_points[depth] == 0) {
             // reached end of full codepoints in token, reject iff it ended in a partial sequence
             // that cannot satisfy this position in grammar
             if (tok.partial_utf8.n_remain != 0 &&
                     !llama_grammar_match_partial_char(stack_pos, tok.partial_utf8)) {
-                rejects.push_back(tok);
+                rejects.push_back(i);
             }
-        } else if (llama_grammar_match_char(stack_pos, *tok.code_points).first) {
-            next_candidates.push_back({ tok.index, tok.code_points + 1, tok.partial_utf8, tok.id });
+        } else if (llama_grammar_match_char(stack_pos, tok.code_points[depth]).first) {
+            next.push_back(i);
         } else {
-            rejects.push_back(tok);
+            rejects.push_back(i);
         }
     }
 
@@ -1118,17 +1117,80 @@ llama_grammar_candidates llama_grammar_reject_candidates_for_stack(
     if (!llama_grammar_is_end_of_sequence(stack_pos_after)) {
         stack_after.push_back(stack_pos_after);
     }
-    llama_grammar_stacks next_stacks;
     llama_grammar_advance_stack(rules, stack_after, next_stacks);
 
-    auto next_rejects = llama_grammar_reject_candidates(rules, next_stacks, next_candidates);
-    for (const auto & tok : next_rejects) {
-        rejects.push_back({ tok.index, tok.code_points - 1, tok.partial_utf8, tok.id });
+    return true;
+}
+
+// the candidates that no stack accepts, each stack narrows the rejects of the previous one
+// no recursion: a level per code point, a token from the GGUF can be long
+static llama_grammar_candidates llama_grammar_reject_candidates(
+        const llama_grammar_rules      & rules,
+        const llama_grammar_stacks     & stacks,
+        const llama_grammar_candidates & candidates) {
+    GGML_ASSERT(!stacks.empty()); // REVIEW
+
+    // a level checks the candidates that reach its code point against its stacks in turn
+    struct level {
+        llama_grammar_stacks  stacks;
+        size_t                i;       // the stack in progress
+        std::vector<uint32_t> cur;     // the candidates rejected by stacks[0, i)
+        std::vector<uint32_t> rejects; // the rejects of stacks[i] at this code point
+    };
+
+    std::vector<uint32_t> all(candidates.size());
+    std::iota(all.begin(), all.end(), 0u);
+
+    std::vector<level> levels;
+    levels.push_back({ stacks, 0, std::move(all), {} });
+
+    while (true) {
+        level & lv = levels.back();
+
+        if (lv.i == lv.stacks.size() || lv.cur.empty()) {
+            if (levels.size() == 1) {
+                break;
+            }
+            // the rejects at the next code point belong to stacks[i] of the parent level
+            level & parent = levels[levels.size() - 2];
+            parent.rejects.insert(parent.rejects.end(), lv.cur.begin(), lv.cur.end());
+            levels.pop_back();
+            parent.cur = std::move(parent.rejects);
+            parent.rejects = {};
+            parent.i++;
+            continue;
+        }
+
+        std::vector<uint32_t> rejects;
+        std::vector<uint32_t> next;
+        llama_grammar_stacks  next_stacks;
+
+        if (!llama_grammar_reject_step(rules, lv.stacks[lv.i], candidates, levels.size() - 1, lv.cur, rejects, next, next_stacks)) {
+            lv.cur = std::move(rejects);
+            lv.i++;
+            continue;
+        }
+
+        GGML_ASSERT(!next_stacks.empty());
+        lv.rejects = std::move(rejects);
+        levels.push_back({ std::move(next_stacks), 0, std::move(next), {} });
     }
 
+    llama_grammar_candidates rejects;
+    rejects.reserve(levels[0].cur.size());
+    for (const uint32_t i : levels[0].cur) {
+        rejects.push_back(candidates[i]);
+    }
     return rejects;
 }
 
+llama_grammar_candidates llama_grammar_reject_candidates_for_stack(
+        const llama_grammar_rules      & rules,
+        const llama_grammar_stack      & stack,
+        const llama_grammar_candidates & candidates) {
+    return llama_grammar_reject_candidates(rules, llama_grammar_stacks{ stack }, candidates);
+}
+
 ////////////////////
 
 struct llama_grammar * llama_grammar_init_impl(
@@ -1163,14 +1225,19 @@ struct llama_grammar * llama_grammar_init_impl(
     std::vector<bool> rules_visited(n_rules);
     std::vector<bool> rules_in_progress(n_rules);
     std::vector<bool> rules_may_be_empty(n_rules);
-    for (size_t i = 0; i < n_rules; i++) {
-        if (rules_visited[i]) {
-            continue;
-        }
-        if (llama_grammar_detect_left_recursion(vec_rules, i, &rules_visited, &rules_in_progress, &rules_may_be_empty)) {
-            LLAMA_LOG_ERROR("unsupported grammar, left recursion detected for nonterminal at index %zu", i);
-            return nullptr;
+    try {
+        for (size_t i = 0; i < n_rules; i++) {
+            if (rules_visited[i]) {
+                continue;
+            }
+            if (llama_grammar_detect_left_recursion(vec_rules, i, &rules_visited, &rules_in_progress, &rules_may_be_empty)) {
+                LLAMA_LOG_ERROR("unsupported grammar, left recursion detected for nonterminal at index %zu", i);
+                return nullptr;
+            }
         }
+    } catch (const std::exception & err) {
+        LLAMA_LOG_ERROR("unsupported grammar: %s\n", err.what());
+        return nullptr;
     }
 
     // loop over alternates of start rule to build initial stacks
@@ -1256,14 +1323,19 @@ struct llama_grammar * llama_grammar_init_impl(
     std::vector<bool> rules_visited(n_rules);
     std::vector<bool> rules_in_progress(n_rules);
     std::vector<bool> rules_may_be_empty(n_rules);
-    for (size_t i = 0; i < n_rules; i++) {
-        if (rules_visited[i]) {
-            continue;
-        }
-        if (llama_grammar_detect_left_recursion(vec_rules, i, &rules_visited, &rules_in_progress, &rules_may_be_empty)) {
-            LLAMA_LOG_ERROR("unsupported grammar, left recursion detected for nonterminal at index %zu\n", i);
-            return nullptr;
+    try {
+        for (size_t i = 0; i < n_rules; i++) {
+            if (rules_visited[i]) {
+                continue;
+            }
+            if (llama_grammar_detect_left_recursion(vec_rules, i, &rules_visited, &rules_in_progress, &rules_may_be_empty)) {
+                LLAMA_LOG_ERROR("unsupported grammar, left recursion detected for nonterminal at index %zu\n", i);
+                return nullptr;
+            }
         }
+    } catch (const std::exception & err) {
+        LLAMA_LOG_ERROR("unsupported grammar: %s\n", err.what());
+        return nullptr;
     }
 
     // loop over alternates of start rule to build initial stacks

src/llama-grammar.h

@@ -102,13 +102,15 @@ struct llama_grammar_parser {
             const char        * src,
             const std::string & rule_name,
             uint32_t            rule_id,
-            bool                is_nested);
+            bool                is_nested,
+            uint32_t            depth = 0);
 
     const char * parse_sequence(
             const char         * src,
             const std::string  & rule_name,
             llama_grammar_rule & rule,
-            bool               is_nested);
+            bool               is_nested,
+            uint32_t           depth = 0);
 
     const char * parse_rule(const char * src);
 

src/llama-hparams.h

@@ -61,7 +61,7 @@ struct llama_hparams {
     //       for example: n_embd_out() vs n_embd_out_impl
 
     bool vocab_only;
-    bool no_alloc;
+    bool no_alloc; // TODO: no_alloc is not a model hyperparameter; move it out of llama_hparams
     bool rope_finetuned;
     bool use_par_res;
     bool swin_norm;
@@ -358,6 +358,8 @@ struct llama_hparams {
     llama_token ple_embd_token   = LLAMA_TOKEN_NULL;
     // hash constant sets, stored back to back: 1 shared by all PLE layers (qwen4exp), or 1 per PLE layer in layer order (deepseek41)
     uint32_t ple_n_modules       = 1;
+    // deepseek41 Engram: the token that pads the n-gram history, an index into llama_model::engram_token_map
+    uint32_t engram_pad_id       = 0;
     // the file lists PLE layer indices, so this is never a per-layer gguf array and can hold one bit per layer
     std::bitset<LLAMA_MAX_LAYERS> is_ple_impl;
     // the hash multipliers reach ~2e13 and have to stay 64-bit

src/llama-kv-cache-dsv4.cpp

@@ -428,8 +428,11 @@ static void dsv4_state_read_k_cache(
         dsv4_state_read_tensor_streams(io, kv->get_k_storage(il), kv_size, n_rows_ref, s0, ns);
     }
 
-    // refill the mirrors from their source rows just read
-    std::vector<uint8_t> buf;
+    if (n_rows_ref == 0) {
+        return;
+    }
+
+    // refill the mirrors from their sources; a mirror gets the same rows as its source, so a whole copy is exact
     for (uint32_t il : kv->get_layer_ids()) {
         const int32_t src = kv_src[il];
         if (src == (int32_t) il) {
@@ -437,20 +440,11 @@ static void dsv4_state_read_k_cache(
         }
         ggml_tensor * t_src = kv->get_k_storage(src);
         ggml_tensor * t_dst = kv->get_k_storage(il);
-        GGML_ASSERT(t_src->type == t_dst->type && t_src->ne[0] == t_dst->ne[0]);
+        if (t_src->type != t_dst->type || !ggml_are_same_shape(t_src, t_dst)) {
+            throw std::runtime_error("DSV4 K-cache mirror layout mismatch");
+        }
 
-        const size_t row_size      = ggml_row_size(t_src->type, t_src->ne[0]);
-        const size_t stream_stride = (size_t) kv_size*row_size;
-        const size_t size          = (size_t) n_rows_ref*row_size;
-        if (size == 0) {
-            continue;
-        }
-        buf.resize(size);
-        for (uint32_t s = 0; s < ns; ++s) {
-            const size_t offset = (size_t) (s0 + s)*stream_stride;
-            ggml_backend_tensor_get(t_src, buf.data(), offset, size);
-            ggml_backend_tensor_set(t_dst, buf.data(), offset, size);
-        }
+        ggml_backend_tensor_copy(t_src, t_dst);
     }
 }
 
@@ -946,13 +940,12 @@ llama_dsv4_comp_state::llama_dsv4_comp_state(
             uint32_t        n_rs_seq,
         const char    * name,
         const llama_memory_i::layer_filter_cb & filter) :
+    hparams(model.hparams),
     ratio(ratio),
     state_size(state_size),
     n_embd_state(n_embd_state),
     n_stream(unified ? 1 : n_seq_max),
     n_rs_seq(n_rs_seq) {
-    const llama_hparams & hparams = model.hparams;
-
     struct ggml_backend_buft_comparator {
         bool operator()(const ggml_backend_buffer_type_t & lhs, const ggml_backend_buffer_type_t & rhs) const {
             return strcmp(ggml_backend_buft_name(lhs), ggml_backend_buft_name(rhs)) < 0;
@@ -1027,7 +1020,15 @@ llama_dsv4_comp_state::llama_dsv4_comp_state(
     }
 
     for (auto & [buft, ctx] : ctx_map) {
-        ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft);
+        ggml_backend_buffer_t buf;
+        if (hparams.no_alloc) {
+            buf = ggml_backend_buft_alloc_buffer(buft, /*size =*/ 0); // dummy buffer
+            for (ggml_tensor * t = ggml_get_first_tensor(ctx.get()); t != nullptr; t = ggml_get_next_tensor(ctx.get(), t)) {
+                t->buffer = buf; // set dummy buffer for DSV4 state so that the backend scheduler won't try to allocate it
+            }
+        } else {
+            buf = ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft); // real buffer
+        }
         if (!buf) {
             throw std::runtime_error("failed to allocate buffer for DSV4 compressor state");
         }
@@ -1119,9 +1120,15 @@ uint32_t llama_dsv4_comp_state::get_n_rows() const {
 
 std::map<ggml_backend_buffer_type_t, size_t> llama_dsv4_comp_state::memory_breakdown() const {
     std::map<ggml_backend_buffer_type_t, size_t> ret;
-    for (const auto & [_, buf] : ctxs_bufs) {
+    for (const auto & [ctx, buf] : ctxs_bufs) {
         ggml_backend_buffer_type_t buft = ggml_backend_buffer_get_type(buf.get());
-        ret[buft] += ggml_backend_buffer_get_size(buf.get());
+
+        if (hparams.no_alloc) {
+            GGML_ASSERT(ggml_backend_buffer_get_base(buf.get()) == nullptr);
+            ret[buft] += ggml_backend_alloc_ctx_tensors_from_buft_size(ctx.get(), buft);
+        } else {
+            ret[buft] += ggml_backend_buffer_get_size(buf.get());
+        }
     }
     return ret;
 }
@@ -1339,10 +1346,10 @@ llama_kv_cache_dsv4::llama_kv_cache_dsv4(
     // crosses devices. no mirrors without offload
     const uint32_t n_layer = model.hparams.n_layer();
     comp_read.assign(n_layer, -1);
-    comp_kv_src.resize(n_layer);
+    comp_il_src.resize(n_layer);
     for (uint32_t il = 0; il < n_layer; ++il) {
         const int32_t src = is_v41 ? model.hparams.dsv41_kv_src[il] : (int32_t) il;
-        comp_kv_src[il] = src >= 0 ? src : (int32_t) il;
+        comp_il_src[il] = src >= 0 ? src : (int32_t) il;
         if (model.hparams.dsv4_compress_ratios[il] == 0) {
             continue;
         }
@@ -1353,7 +1360,7 @@ llama_kv_cache_dsv4::llama_kv_cache_dsv4(
         // the first layer of this source on this device holds the mirror
         comp_read[il] = il;
         for (int32_t j = src + 1; j < (int32_t) il; ++j) {
-            if (comp_read[j] == j && comp_kv_src[j] == src && model.dev_layer(j) == model.dev_layer(il)) {
+            if (comp_read[j] == j && comp_il_src[j] == src && model.dev_layer(j) == model.dev_layer(il)) {
                 comp_read[il] = j;
                 break;
             }
@@ -1798,9 +1805,9 @@ void llama_kv_cache_dsv4::state_write(llama_io_write_i & io, llama_seq_id seq_id
         const uint32_t n_rows_lid = seq_id >= 0 ?
             dsv4_state_n_used_k_rows(pos_max, lid_state->get_ratio(), kv_lid->get_size()) : kv_lid->get_size();
 
-        dsv4_state_write_k_cache(io, kv_csa.get(), seq_id, flags, n_rows_csa, comp_kv_src);
-        dsv4_state_write_k_cache(io, kv_hca.get(), seq_id, flags, n_rows_hca, comp_kv_src);
-        dsv4_state_write_k_cache(io, kv_lid.get(), seq_id, flags, n_rows_lid, comp_kv_src);
+        dsv4_state_write_k_cache(io, kv_csa.get(), seq_id, flags, n_rows_csa, comp_il_src);
+        dsv4_state_write_k_cache(io, kv_hca.get(), seq_id, flags, n_rows_hca, comp_il_src);
+        dsv4_state_write_k_cache(io, kv_lid.get(), seq_id, flags, n_rows_lid, comp_il_src);
     }
 
     csa_state->state_write(io, seq_id, flags, rs_idx);
@@ -1838,9 +1845,9 @@ void llama_kv_cache_dsv4::state_read(llama_io_read_i & io, llama_seq_id seq_id,
     if (!partial_only) {
         clear_compressed(seq_id, true);
 
-        dsv4_state_read_k_cache(io, kv_csa.get(), seq_id, flags, comp_kv_src);
-        dsv4_state_read_k_cache(io, kv_hca.get(), seq_id, flags, comp_kv_src);
-        dsv4_state_read_k_cache(io, kv_lid.get(), seq_id, flags, comp_kv_src);
+        dsv4_state_read_k_cache(io, kv_csa.get(), seq_id, flags, comp_il_src);
+        dsv4_state_read_k_cache(io, kv_hca.get(), seq_id, flags, comp_il_src);
+        dsv4_state_read_k_cache(io, kv_lid.get(), seq_id, flags, comp_il_src);
     }
 
     csa_state->state_read(io, seq_id, flags);

src/llama-kv-cache-dsv4.h

@@ -61,6 +61,8 @@ private:
         std::vector<ggml_tensor *> score_stream;
     };
 
+    const llama_hparams & hparams;
+
     const uint32_t ratio;
     const uint32_t state_size;
     const uint32_t n_embd_state;
@@ -168,7 +170,7 @@ private:
     std::vector<uint32_t> rs_idx;
 
     // [n_layer] storage layer -> layer whose rows it holds (itself, or the source of a mirror)
-    std::vector<int32_t> comp_kv_src;
+    std::vector<int32_t> comp_il_src;
 
     // [n_layer] layer -> layer whose compressed-KV storage it reads (itself, its V4.1 kv source or a mirror), -1 = none
     std::vector<int32_t> comp_read;

src/llama-memory-recurrent.cpp

@@ -116,7 +116,15 @@ llama_memory_recurrent::llama_memory_recurrent(
 
     // allocate tensors and initialize the buffers to avoid NaNs in the padding
     for (auto & [buft, ctx] : ctx_map) {
-        ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft);
+        ggml_backend_buffer_t buf;
+        if (hparams.no_alloc) {
+            buf = ggml_backend_buft_alloc_buffer(buft, /*size =*/ 0); // dummy buffer
+            for (ggml_tensor * t = ggml_get_first_tensor(ctx.get()); t != nullptr; t = ggml_get_next_tensor(ctx.get(), t)) {
+                t->buffer = buf; // set dummy buffer for RS cache so that the backend scheduler won't try to allocate it
+            }
+        } else {
+            buf = ggml_backend_alloc_ctx_tensors_from_buft(ctx.get(), buft); // real buffer
+        }
         if (!buf) {
             throw std::runtime_error("failed to allocate buffer for rs cache");
         }
@@ -435,8 +443,15 @@ void llama_memory_recurrent::set_rs_idx(llama_seq_id seq_id, uint32_t idx) {
 
 std::map<ggml_backend_buffer_type_t, size_t> llama_memory_recurrent::memory_breakdown() const {
     std::map<ggml_backend_buffer_type_t, size_t> ret;
-    for (const auto & [_, buf] : ctxs_bufs) {
-        ret[ggml_backend_buffer_get_type(buf.get())] += ggml_backend_buffer_get_size(buf.get());
+    for (const auto & [ctx, buf] : ctxs_bufs) {
+        ggml_backend_buffer_type_t buft = ggml_backend_buffer_get_type(buf.get());
+
+        if (hparams.no_alloc) {
+            GGML_ASSERT(ggml_backend_buffer_get_base(buf.get()) == nullptr);
+            ret[buft] += ggml_backend_alloc_ctx_tensors_from_buft_size(ctx.get(), buft);
+        } else {
+            ret[buft] += ggml_backend_buffer_get_size(buf.get());
+        }
     }
     return ret;
 }

src/llama-model-saver.cpp

@@ -4,44 +4,23 @@
 #include "gguf.h"
 
 #include "llama-arch.h"
+#include "llama-ext.h"
 #include "llama.h"
 #include "llama-hparams.h"
 #include "llama-model.h"
 #include "llama-vocab.h"
-#include "models/models.h"
 
 #include <cstdint>
 #include <string>
 #include <unordered_set>
 
-bool llama_model_saver_supports_arch(llm_arch arch) {
-    switch (arch) {
-        case LLM_ARCH_SENSENOVA_U1:
-        case LLM_ARCH_GEMMA3N:
-        case LLM_ARCH_BITNET:
-        case LLM_ARCH_T5:
-        case LLM_ARCH_APERTUS:
-        case LLM_ARCH_STEP35:
-        case LLM_ARCH_CLEF: // the head tensors are not saved
-        case LLM_ARCH_VLLM_DECISION2:
-            return false;
-        default:
-            return true;
-    }
-}
-
 llama_model_saver::llama_model_saver(const struct llama_model * model) :
-        gguf_ctx(gguf_init_empty()), gguf_ctx_owned(true), model(model), llm_kv(model->arch) {
-    GGML_ASSERT(llama_model_saver_supports_arch(model->arch));
+        gguf_ctx(gguf_init_empty()), model(model), llm_kv(model->arch) {
+    GGML_ASSERT(llama_arch_supports_save(model->arch));
 }
 
-llama_model_saver::llama_model_saver(enum llm_arch arch, struct gguf_context * gguf_ctx) :
-        gguf_ctx(gguf_ctx == nullptr ? gguf_init_empty() : gguf_ctx), gguf_ctx_owned(gguf_ctx == nullptr), model(nullptr), llm_kv(arch) {}
-
 llama_model_saver::~llama_model_saver() {
-    if (gguf_ctx_owned) {
-        gguf_free(gguf_ctx);
-    }
+    gguf_free(gguf_ctx);
 }
 
 void llama_model_saver::add_kv(const enum llm_kv key, const uint32_t value) {
@@ -119,11 +98,6 @@ void llama_model_saver::add_kv(const enum llm_kv key, const Container & value, c
         GGML_ABORT("fatal error");
     }
 }
-// instantiate for external usage:
-template void llama_model_saver::add_kv<std::vector<uint32_t>>(const enum llm_kv, const std::vector<uint32_t> &, const bool);
-template void llama_model_saver::add_kv<std::vector<float>>(const enum llm_kv, const std::vector<float> &, const bool);
-template void llama_model_saver::add_kv<std::vector<uint64_t>>(const enum llm_kv, const std::vector<uint64_t> &, const bool);
-template void llama_model_saver::add_kv<std::vector<int32_t>>(const enum llm_kv, const std::vector<int32_t> &, const bool);
 
 void llama_model_saver::add_kv(const enum llm_kv key, const std::vector<std::string> & value) {
     std::vector<const char *> tmp(value.size());
@@ -342,8 +316,6 @@ void llama_model_saver::add_kv_from_model() {
     add_kv(LLM_KV_HYPER_CONNECTION_LOW_RANK,             hparams.hc_low_rank);
 
     if (model->arch == LLM_ARCH_DEEPSEEK41) {
-        const auto * m41 = dynamic_cast<const llama_model_deepseek41 *>(model);
-        GGML_ASSERT(m41 != nullptr);
         std::vector<uint32_t> kv_src;
         std::vector<uint32_t> idx_src;
         for (uint32_t il = 0; il < hparams.n_layer(); ++il) {
@@ -362,8 +334,8 @@ void llama_model_saver::add_kv_from_model() {
             add_kv(LLM_KV_ATTENTION_INDEXER_CANDIDATE_BLOCK_SIZE,   hparams.dsv41_cand_block_size);
         }
         if (hparams.ple_n_heads > 0) {
-            add_kv(LLM_KV_ENGRAM_PAD_TOKEN_ID, m41->engram_pad_id);
-            add_kv(LLM_KV_ENGRAM_TOKEN_MAP,    m41->engram_token_map);
+            add_kv(LLM_KV_ENGRAM_PAD_TOKEN_ID, hparams.engram_pad_id);
+            add_kv(LLM_KV_ENGRAM_TOKEN_MAP,    model->engram_token_map);
         }
     }
 
@@ -551,7 +523,3 @@ void llama_model_saver::add_tensors_from_model() {
 void llama_model_saver::save(const std::string & path_model) {
     gguf_write_to_file(gguf_ctx, path_model.c_str(), false);
 }
-
-void llama_model_saver::save(FILE * file) {
-    gguf_write_to_file_ptr(gguf_ctx, file, false);
-}

src/llama-model-saver.h

@@ -6,17 +6,12 @@
 
 #include <vector>
 
-// FIXME temporary function for better error messages
-bool llama_model_saver_supports_arch(llm_arch arch);
-
 struct llama_model_saver {
     struct gguf_context * gguf_ctx = nullptr;
-    const bool gguf_ctx_owned;
     const struct llama_model * model;
     const struct LLM_KV llm_kv;
 
     llama_model_saver(const struct llama_model * model);
-    llama_model_saver(enum llm_arch arch, struct gguf_context * gguf_ctx);
     ~llama_model_saver();
 
     void add_kv(enum llm_kv key, uint32_t     value);
@@ -41,5 +36,4 @@ struct llama_model_saver {
     void add_tensors_from_model();
 
     void save(const std::string & path_model);
-    void save(FILE * file);
 };

src/llama-model.h

@@ -623,6 +623,10 @@ struct llama_layer {
     struct ggml_tensor * hc_ffn_down    = nullptr;
     struct ggml_tensor * hc_ffn_up      = nullptr;
     struct ggml_tensor * hc_ffn_inject  = nullptr;
+    struct ggml_tensor * hc_attn_down_inject = nullptr; // [hc_attn_down | hc_attn_inject] in one weight
+    struct ggml_tensor * hc_ffn_down_inject  = nullptr; // [hc_ffn_down | hc_ffn_inject] in one weight
+    struct ggml_tensor * wqkv_z              = nullptr; // [wqkv | wqkv_gate] in one weight
+    struct ggml_tensor * ssm_beta_alpha_cat  = nullptr; // [ssm_beta | ssm_alpha] in one weight, unlike the grouped ssm_beta_alpha
 
     struct ggml_tensor * ple_key        = nullptr;
     struct ggml_tensor * ple_value      = nullptr;
@@ -730,6 +734,9 @@ struct llama_model {
     std::vector<std::string> classifier_labels;
     std::vector<std::string> decision_labels;
 
+    // deepseek41 Engram: token id -> id of its normalized form, hashed into the n-grams
+    std::vector<int32_t> engram_token_map;
+
     struct ggml_tensor * tok_embd   = nullptr;
     struct ggml_tensor * type_embd  = nullptr;
     struct ggml_tensor * pos_embd   = nullptr;

src/llama-quant.cpp

@@ -125,7 +125,8 @@ static tensor_category tensor_get_category(const std::string & tensor_name) {
     if (tensor_name_match_token_embd(tensor_name.c_str())) {
         return tensor_category::TOKEN_EMBD;
     }
-    if (tensor_name.find("attn_qkv.weight") != std::string::npos) {
+    // attn_qkv_gate is [attn_qkv | attn_gate] in one weight, the qkv rows decide its type
+    if (tensor_name.find("attn_qkv.weight") != std::string::npos || tensor_name.find("attn_qkv_gate.weight") != std::string::npos) {
         return tensor_category::ATTENTION_QKV;
     }
     if (tensor_name.find("attn_kv_b.weight") != std::string::npos) {
@@ -161,6 +162,12 @@ static tensor_category tensor_get_category(const std::string & tensor_name) {
     return tensor_category::OTHER;
 }
 
+// check if tensor is a MoE expert tensor (used by GGML_OP_MUL_MAT_ID)
+static bool tensor_is_expert(const std::string & tensor_name) {
+    return tensor_name.find("_exps.")   != std::string::npos ||
+           tensor_name.find("_chexps.") != std::string::npos;
+}
+
 // check if category is for attention-v-like tensors (more sensitive to quantization)
 static bool category_is_attn_v(tensor_category cat) {
     return cat == tensor_category::ATTENTION_V     ||
@@ -296,7 +303,7 @@ static bool tensor_allows_quantization(const llama_model_quantize_params * param
     // trivial checks first -- no string ops needed
     if (params->only_copy)       return false;
 
-    // quantize only 2D and 3D tensors (experts)
+    // quantize only 2D+ tensors (matrices, experts, conv kernels)
     if (ggml_n_dims(tensor) < 2) return false;
 
     const std::string name = ggml_get_name(tensor);
@@ -541,12 +548,7 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type
     } else if (ftype == LLAMA_FTYPE_MOSTLY_MXFP4_MOE) {
         // MoE   tensors -> MXFP4
         // other tensors -> Q8_0
-        // MLA projection tensors are also 3D, so match expert tensor roles explicitly.
-        const bool is_bailingmoe3_expert = arch == LLM_ARCH_BAILINGMOE3 &&
-            (category == tensor_category::FFN_UP ||
-             category == tensor_category::FFN_GATE ||
-             category == tensor_category::FFN_DOWN);
-        if (tensor->ne[2] > 1 && (arch != LLM_ARCH_BAILINGMOE3 || is_bailingmoe3_expert)) {
+        if (tensor_is_expert(name)) {
             new_type = GGML_TYPE_MXFP4;
         } else {
             new_type = GGML_TYPE_Q8_0;
@@ -1307,11 +1309,11 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
                     if (it == imatrix_data->end()) {
                         LLAMA_LOG_INFO("\n====== %s: did not find weights for %s\n", __func__, tensor->name);
                     } else {
-                        if (it->second.size() == (size_t)tensor->ne[0]*tensor->ne[2]) {
+                        if (it->second.size() == (size_t)tensor->ne[0]*tensor->ne[2]*tensor->ne[3]) {
                             imatrix = it->second.data();
                         } else {
                             LLAMA_LOG_INFO("\n====== %s: imatrix size %d is different from tensor size %d for %s\n", __func__,
-                                    int(it->second.size()), int(tensor->ne[0]*tensor->ne[2]), tensor->name);
+                                    int(it->second.size()), int(tensor->ne[0]*tensor->ne[2]*tensor->ne[3]), tensor->name);
 
                             // this can happen when quantizing an old mixtral model with split tensors with a new incompatible imatrix
                             // this is a significant error and it may be good idea to abort the process if this happens,
@@ -1319,7 +1321,7 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
                             // tok_embd should be ignored in this case, since it always causes this warning
                             if (!tensor_name_match_token_embd(tensor->name)) {
                                 throw std::runtime_error(format("imatrix size %d is different from tensor size %d for %s",
-                                        int(it->second.size()), int(tensor->ne[0]*tensor->ne[2]), tensor->name));
+                                        int(it->second.size()), int(tensor->ne[0]*tensor->ne[2]*tensor->ne[3]), tensor->name));
                             }
                         }
                     }
@@ -1341,7 +1343,7 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
 
                 const int64_t n_per_row = tensor->ne[0];
                 const int64_t nrows_per_expert = tensor->ne[1];
-                const int64_t nrows_total = tensor->ne[1] * tensor->ne[2];
+                const int64_t nrows_total = ggml_nrows(tensor);
 
                 const size_t row_size_src = ggml_row_size(tensor->type, n_per_row);
                 const size_t row_size_dst = ggml_row_size(new_type,     n_per_row);

src/llama-sampler.cpp

@@ -18,6 +18,7 @@
 #include <numeric>
 #include <random>
 #include <unordered_map>
+#include <unordered_set>
 #include <stdexcept>
 
 // the ring buffer works similarly to std::deque, but with a fixed capacity
@@ -3333,9 +3334,13 @@ struct llama_sampler_dry {
 };
 
 // Ported from Koboldcpp, original PR: https://github.com/LostRuins/koboldcpp/pull/982 (Original author: pi6am)
-static void get_overlapping_token_sequences(const llama_vocab & vocab, const std::string& str, std::unordered_multimap<llama_token, std::vector<llama_token>>& token_sequences, int max_tail_len = -1) {
-    for (llama_token token_id = 0; token_id < (llama_token) vocab.n_tokens(); token_id++) {
-        std::string word = vocab.detokenize({token_id}, true);
+static void get_overlapping_token_sequences(const llama_vocab & vocab, const std::vector<std::string> & vocab_words, const std::string& str, std::unordered_multimap<llama_token, std::vector<llama_token>>& token_sequences, int max_tail_len = -1) {
+    // the tail tokenization only depends on how many characters of str were matched, so compute each one at most once
+    std::vector<std::vector<llama_token>> tail_cache(str.size() + 1);
+    std::vector<bool> tail_cached(str.size() + 1, false);
+
+    for (llama_token token_id = 0; token_id < (llama_token) vocab_words.size(); token_id++) {
+        const std::string & word = vocab_words[token_id];
         if (word.find(str) != std::string::npos) {
             token_sequences.emplace(token_id, std::vector<llama_token>());
         } else {
@@ -3352,10 +3357,14 @@ static void get_overlapping_token_sequences(const llama_vocab & vocab, const std
                     }
                 }
                 if (match) {
-                    std::vector<llama_token> tokenization = vocab.tokenize(str.substr(i), false, false);
-                    if (max_tail_len >= 0 && tokenization.size() > (size_t)max_tail_len) {
-                        tokenization.resize(max_tail_len);
+                    if (!tail_cached[i]) {
+                        tail_cache[i] = vocab.tokenize(str.substr(i), false, false);
+                        if (max_tail_len >= 0 && tail_cache[i].size() > (size_t)max_tail_len) {
+                            tail_cache[i].resize(max_tail_len);
+                        }
+                        tail_cached[i] = true;
                     }
+                    const std::vector<llama_token> & tokenization = tail_cache[i];
 
                     // Ensure we don't already have a duplicate matching tokenization
                     auto its = token_sequences.equal_range(token_id);
@@ -3656,6 +3665,16 @@ struct llama_sampler * llama_sampler_init_dry(const struct llama_vocab * vocab,
 
     if (dry_enabled && seq_breakers != nullptr && num_breakers > 0) {
         // Process sequence breakers
+        // each distinct breaker costs a full vocabulary scan: drop duplicates and stop at the cap
+        std::unordered_set<std::string> seen_breakers;
+
+        // detokenize the vocabulary once instead of rescanning it for every breaker
+        std::vector<std::string> vocab_words;
+        vocab_words.reserve(vocab->n_tokens());
+        for (llama_token token_id = 0; token_id < (llama_token) vocab->n_tokens(); ++token_id) {
+            vocab_words.push_back(vocab->detokenize({token_id}, true));
+        }
+
         for (size_t i = 0; i < num_breakers; ++i) {
             if (seq_breakers[i] == nullptr || std::strlen(seq_breakers[i]) == 0) {
                 LLAMA_LOG_WARN("skipping null or empty DRY sequence breaker at index %zu\n", i);
@@ -3663,17 +3682,22 @@ struct llama_sampler * llama_sampler_init_dry(const struct llama_vocab * vocab,
             }
 
             std::string sequence_break(seq_breakers[i]);
-            if (sequence_break.empty()) {
-                LLAMA_LOG_WARN("skipping empty DRY sequence breaker\n");
-                continue;
-            }
-
             if (sequence_break.size() > MAX_CHAR_LEN) {
                 LLAMA_LOG_WARN("truncating DRY sequence breaker to %d characters\n", MAX_CHAR_LEN);
                 sequence_break.resize(MAX_CHAR_LEN);
             }
 
-            get_overlapping_token_sequences(*vocab, sequence_break, processed_breakers, MAX_SEQ_LEN);
+            if (!seen_breakers.insert(sequence_break).second) {
+                continue;
+            }
+
+            if (seen_breakers.size() > LLAMA_DRY_MAX_SEQ_BREAKERS) {
+                LLAMA_LOG_WARN("ignoring DRY sequence breaker at index %zu: more than %d distinct breakers\n",
+                        i, LLAMA_DRY_MAX_SEQ_BREAKERS);
+                break;
+            }
+
+            get_overlapping_token_sequences(*vocab, vocab_words, sequence_break, processed_breakers, MAX_SEQ_LEN);
         }
     }
 

src/llama-vocab.cpp

@@ -29,31 +29,25 @@
 struct naive_trie {
     naive_trie() : has_value(false), value(0) {
     }
+    // no recursion: a token from the GGUF can be long
     void insert(const char * key, size_t len, int32_t value = 0) {
-        if (len == 0) {
-            this->has_value = true;
-            this->value = value;
-            return;
-        }
-        char c = key[0];
-        auto res = children.find(c);
-        if (res != children.end()) {
-            res->second.insert(key + 1, len - 1, value);
-        } else {
-            auto res = children.insert(std::make_pair(c, naive_trie()));
-            res.first->second.insert(key + 1, len - 1, value);
+        naive_trie * node = this;
+        for (size_t i = 0; i < len; i++) {
+            node = &node->children[key[i]];
         }
+        node->has_value = true;
+        node->value = value;
     }
     std::pair<const char *, size_t> get_longest_prefix(const char * key, size_t len, size_t offset = 0) const {
-        if (len == 0 || offset == len) {
-            return std::make_pair(key, offset);
+        const naive_trie * node = this;
+        while (offset < len) {
+            auto res = node->children.find(key[offset]);
+            if (res == node->children.end()) {
+                break;
+            }
+            node = &res->second;
+            offset++;
         }
-        char c = key[offset];
-        auto res = children.find(c);
-        if (res != children.end()) {
-            return res->second.get_longest_prefix(key, len, offset + 1);
-        }
-
         return std::make_pair(key, offset);
     }
     const struct naive_trie * traverse(const char c) const {
@@ -174,30 +168,37 @@ struct llm_tokenizer_spm_session {
     }
 
 private:
+    // no recursion: the merges of a long token can be deep
     void resegment(llm_symbol & symbol, std::vector<llama_token> & output) {
-        auto text = std::string(symbol.text, symbol.n);
-        auto token = vocab.text_to_token(text);
+        std::vector<const llm_symbol *> todo = { &symbol };
+        while (!todo.empty()) {
+            const llm_symbol & cur = *todo.back();
+            todo.pop_back();
 
-        // Do we need to support is_unused?
-        if (token != LLAMA_TOKEN_NULL) {
-            output.push_back(token);
-            return;
-        }
+            auto text = std::string(cur.text, cur.n);
+            auto token = vocab.text_to_token(text);
 
-        const auto p = rev_merge.find(text);
-
-        if (p == rev_merge.end()) {
-            // output any symbols that did not form tokens as bytes.
-            output.reserve(output.size() + symbol.n);
-            for (int j = 0; j < (int)symbol.n; ++j) {
-                llama_token id = vocab.byte_to_token(symbol.text[j]);
-                output.push_back(id);
+            // Do we need to support is_unused?
+            if (token != LLAMA_TOKEN_NULL) {
+                output.push_back(token);
+                continue;
             }
-            return;
-        }
 
-        resegment(symbols[p->second.first], output);
-        resegment(symbols[p->second.second], output);
+            const auto p = rev_merge.find(text);
+
+            if (p == rev_merge.end()) {
+                // output any symbols that did not form tokens as bytes.
+                output.reserve(output.size() + cur.n);
+                for (int j = 0; j < (int)cur.n; ++j) {
+                    llama_token id = vocab.byte_to_token(cur.text[j]);
+                    output.push_back(id);
+                }
+                continue;
+            }
+
+            todo.push_back(&symbols[p->second.second]);
+            todo.push_back(&symbols[p->second.first]);
+        }
     }
 
     void try_add_bigram(int left, int right) {

src/models/deepseek4.cpp

@@ -350,7 +350,7 @@ ggml_tensor * llama_model_deepseek4::graph::build_hc_pre(
         ggml_tensor ** post,
         ggml_tensor ** comb,
         int il) const {
-    return build_hc_pre(x, build_hc_mixes(x, hc_fn, hc_scale, hc_base, post, comb, false, il), il);
+    return build_hc_pre(x, build_hc_mixes(x, hc_fn, hc_scale, hc_base, post, comb, il), il);
 }
 
 ggml_tensor * llama_model_deepseek4::graph::build_hc_mixes(
@@ -360,7 +360,6 @@ ggml_tensor * llama_model_deepseek4::graph::build_hc_mixes(
         ggml_tensor * hc_base,
         ggml_tensor ** post,
         ggml_tensor ** comb,
-        bool f32_src,
         int il) const {
     const int64_t hc         = hparams.dsv4_hc_mult;
     const int64_t hc_dim     = hc*n_embd;
@@ -373,10 +372,6 @@ ggml_tensor * llama_model_deepseek4::graph::build_hc_mixes(
     ggml_tensor * flat = ggml_reshape_2d(ctx0, x, hc_dim, nt);
     ggml_tensor * flat_norm = ggml_rms_norm(ctx0, flat, norm_rms_eps);
     ggml_tensor * mixes = ggml_mul_mat(ctx0, hc_fn, flat_norm);
-    if (f32_src && !ggml_is_quantized(hc_fn->type)) {
-        // F32 weights (model.py hc_pre runs in fp32): request F32 inputs, the mixes are near-tie sensitive
-        ggml_prec_set_src(mixes, GGML_PREC_F32, 1);
-    }
     cb(mixes, "hc_mixes", il);
 
     ggml_tensor * scale_pre  = dsv4_view_1d(ctx0, hc_scale, 1, 0);

src/models/deepseek41.cpp

@@ -11,9 +11,7 @@
 #include <string>
 
 // ref: deepseek-ai/DeepSeek-V4.1-Flash-1017 inference/model.py, engram.py
-// compressed layers attend over [window | all compressed rows] with top-k as a mask. with flash attention the default
-// is sparse: fused indexer, FA bounded to window + top-k, decode-sized ubatches gather only the selected rows (same
-// keys and mask, different key order). LLAMA_DSV41_SPARSE=0 restores the reference path
+// LLAMA_DSV41_SPARSE=0 disables the sparse attention path (see build_attention41)
 
 static constexpr int64_t DSV41_ENGRAM_DEAD = -1;
 
@@ -96,7 +94,7 @@ void llama_model_deepseek41::load_arch_hparams(llama_model_loader & ml) {
         ml.get_arr(LLM_KV_PLE_LAYER_MULTIPLIERS,      mult);
         ml.get_arr(LLM_KV_PLE_HEAD_OFFSETS,           offsets);
         ml.get_arr(LLM_KV_PLE_HEAD_VOCAB_SIZES,       primes);
-        ml.get_key(LLM_KV_ENGRAM_PAD_TOKEN_ID,        engram_pad_id);
+        ml.get_key(LLM_KV_ENGRAM_PAD_TOKEN_ID,        hparams.engram_pad_id);
         ml.get_arr(LLM_KV_ENGRAM_TOKEN_MAP,           engram_token_map);
 
         const size_t n_e    = e_layers.size();
@@ -105,7 +103,7 @@ void llama_model_deepseek41::load_arch_hparams(llama_model_loader & ml) {
         if (n_cols == 0 || hparams.ple_head_dim == 0 ||
                 mult.size() != n_e*n_gram || primes.size() != n_e*n_cols || offsets.size() != n_e*n_cols ||
                 mult.size() > LLAMA_MAX_PLE_NGRAM || primes.size() > LLAMA_MAX_PLE_HEADS ||
-                engram_pad_id >= engram_token_map.size()) {
+                hparams.engram_pad_id >= engram_token_map.size()) {
             throw std::runtime_error("DeepSeek-V4.1 engram metadata is inconsistent");
         }
         for (size_t e = 0; e < n_e; ++e) {
@@ -247,7 +245,7 @@ void llama_model_deepseek41::engram_hash(int e, const llama_ubatch & ubatch, con
 
     const uint64_t * mult   = hparams.ple_layer_multipliers.data() + (size_t) e*n_gram;
     const uint32_t * primes = hparams.ple_head_vocab_sizes.data()  + (size_t) e*n_cols;
-    const int64_t    pad    = engram_token_map[engram_pad_id];
+    const int64_t    pad    = engram_token_map[hparams.engram_pad_id];
     const uint32_t * offs   = hparams.ple_head_offsets.data()      + (size_t) e*n_cols;
 
     const auto cid = [&](llama_token t) -> int64_t {
@@ -279,8 +277,7 @@ void llama_model_deepseek41::engram_hash(int e, const llama_ubatch & ubatch, con
     }
 }
 
-// cand_pin: +1e30 (the reference uses +inf) on the block holding each query's newest compressed row, so a partly
-// filled block always survives; a finite pin keeps fully masked blocks at -inf instead of NaN
+// cand_pin: +1e30 on the block with the newest compressed row of each query, so a partly filled block is always kept; the reference uses +inf, but a finite pin keeps fully masked blocks at -inf instead of NaN
 class llm_graph_input_dsv41 : public llm_graph_input_i {
 public:
     llm_graph_input_dsv41(const llama_model_deepseek41 & model, const llama_kv_cache_dsv4_context * mctx) :
@@ -535,8 +532,7 @@ ggml_tensor * llama_model_deepseek41::graph::build_attention41(
         if (hparams.dsv41_topk_src[il] == il) {
             ggml_tensor * ik = ggml_view_4d(ctx0, comp_all, n_idx_embd, 1, n_comp, n_stream,
                     comp_all->nb[1], comp_all->nb[2], comp_all->nb[3], ggml_row_size(comp_all->type, n_embd_head));
-            // the fused indexer reads the strided view in place, the unfused path needs it contiguous; quantized
-            // caches are dequantized straight from the strided view (CUDA has no same-type copy of quantized views)
+            // the fused indexer reads the strided view in place, the unfused path needs it contiguous; CUDA has no same-type copy of quantized views, so dequantize them
             const bool fused_idx = sparse && cparams.fused_lid && ci.kq_mask->type == GGML_TYPE_F16;
             if (ggml_is_quantized(ik->type)) {
                 ik = ggml_cast(ctx0, ik, GGML_TYPE_F32);
@@ -551,8 +547,7 @@ ggml_tensor * llama_model_deepseek41::graph::build_attention41(
             ggml_tensor * iw = ggml_scale(ctx0, build_lora_mm(layer.indexer_proj, cur), 1.0f/sqrtf(float(n_idx_embd*n_idx_head)));
             iw = ggml_view_4d(ctx0, iw, iw->ne[0], nt/n_stream, 1, n_stream, iw->nb[1], iw->nb[1]*(nt/n_stream), iw->nb[1]*(nt/n_stream), 0);
 
-            // chunk queries so the [n_ih, tokens, n_comp] scores stay near 2^26 elements (at most 64 chunks); the fused
-            // indexer has no such intermediate and runs as one chunk
+            // chunk queries so the [n_ih, tokens, n_comp] scores stay near 2^26 elements (at most 64 chunks); the fused indexer runs as one chunk
             const int64_t n_tok = nt/n_stream;
             GGML_ASSERT(ci.kq_mask->ne[1] == n_tok);
             const int64_t chunk = fused_idx ? n_tok :
@@ -582,9 +577,6 @@ ggml_tensor * llama_model_deepseek41::graph::build_attention41(
                     ggml_tensor * iqc = ggml_view_4d(ctx0, iq, iq->ne[0], iq->ne[1], nc, n_stream, iq->nb[1], iq->nb[2], iq->nb[3], t0*iq->nb[2]);
 
                     ggml_tensor * kq = ggml_mul_mat(ctx0, ik_p, ggml_permute(ctx0, iqc, 0, 2, 1, 3));
-                    if (!ggml_is_quantized(ik->type)) {
-                        ggml_prec_set_src(kq, GGML_PREC_F32, 1);
-                    }
                     kq = ggml_cont(ctx0, ggml_permute(ctx0, kq, 2, 1, 0, 3));
                     score = ggml_sum_rows(ctx0, ggml_mul(ctx0, ggml_relu(ctx0, kq), rows(iw)));
                     ggml_tensor * mask  = rows(ci.kq_mask);
@@ -656,14 +648,12 @@ ggml_tensor * llama_model_deepseek41::graph::build_attention41(
         const int64_t k_sel    = top_k->ne[0];
         const int64_t n_raw    = raw_k->ne[2];
 
-        // few queries per stream: gather each query's top-k rows and attend over [window | selected], one FA stream
-        // per query, instead of a full n_comp mask per layer. same keys and mask as dense, different key order
+        // few queries per stream: attend over [window | top-k rows of the query], one FA stream per query; same keys and mask as dense, different key order
         const bool gather = sparse && cparams.flash_attn && raw_mask->type == GGML_TYPE_F16 && ci.kq_mask->type == GGML_TYPE_F16 &&
                 ((n_stream == 1 && nt <= 32) || nt == n_stream);
 
         if (gather) {
-            // ggml_top_k returns the rows in no particular order (CUB DeviceTopK can vary between runs) and FA sums the
-            // gathered keys in that order: sort the rows once per top-k source so the result is deterministic
+            // ggml_top_k returns the rows in no fixed order and FA sums the keys in that order: sort them so the result is deterministic
             ggml_tensor *& sel = gather_sel_of[topk_src];
             if (!sel) {
                 GGML_ASSERT(n_comp <= (INT64_C(1) << 24)); // row ids are exact in F32
@@ -714,8 +704,7 @@ ggml_tensor * llama_model_deepseek41::graph::build_attention41(
             }
             ggml_tensor * k_all = ggml_concat(ctx0, raw_k, comp_k, 2);
 
-            // mask: window row, then comp mask at the selected rows (-inf where top-k picked an invisible row);
-            // one per top-k source, shared by its layers
+            // window mask, then the comp mask at the selected rows (-inf where top-k picked an invisible row); one per top-k source
             ggml_tensor *& kq_mask = gather_mask_of[topk_src];
             if (!kq_mask) {
                 ggml_tensor * sel_mask = ggml_get_rows(ctx0, ggml_reshape_3d(ctx0, ci.kq_mask, 1, n_comp, nt), sel);
@@ -730,8 +719,7 @@ ggml_tensor * llama_model_deepseek41::graph::build_attention41(
             ggml_tensor * comp_k = ggml_view_4d(ctx0, comp_all, n_embd_head, 1, n_comp, n_stream,
                     comp_all->nb[1], comp_all->nb[2], comp_all->nb[3], 0);
             if (k_rot) {
-                // compressed rows are stored unrotated: rotate in F32, then cast to the cache type (dequantized from
-                // the strided view: CUDA has no same-type quantized copy)
+                // compressed rows are stored unrotated: rotate in F32, then cast to the cache type
                 comp_k = comp_k->type != GGML_TYPE_F32 ? ggml_cast(ctx0, comp_k, GGML_TYPE_F32) : ggml_cont(ctx0, comp_k);
                 comp_k = llama_mul_mat_hadamard(ctx0, comp_k, k_rot);
                 if (comp_k->type != raw_k->type) {
@@ -862,7 +850,7 @@ llama_model_deepseek41::graph::graph(const llama_model & model_base, const llm_g
         ggml_tensor * post = nullptr;
         ggml_tensor * comb = nullptr;
 
-        ggml_tensor * attn_pre = build_hc_mixes(inpL, layer.hc_attn_fn, layer.hc_attn_scale, layer.hc_attn_base, &post, &comb, true, il);
+        ggml_tensor * attn_pre = build_hc_mixes(inpL, layer.hc_attn_fn, layer.hc_attn_scale, layer.hc_attn_base, &post, &comb, il);
 
         // the first layer's attention collapses with the identity pre-mix: copy 0
         cur = pre_prev ? build_hc_pre(inpL, pre_prev, il) : ggml_cont(ctx0, ggml_view_2d(ctx0, inpL, n_embd, n_tokens, inpL->nb[2], 0));
@@ -877,7 +865,7 @@ llama_model_deepseek41::graph::graph(const llama_model & model_base, const llm_g
         cb(inpL, "hc_attn_post", il);
 
         residual = inpL;
-        pre_prev = build_hc_mixes(inpL, layer.hc_ffn_fn, layer.hc_ffn_scale, layer.hc_ffn_base, &post, &comb, true, il);
+        pre_prev = build_hc_mixes(inpL, layer.hc_ffn_fn, layer.hc_ffn_scale, layer.hc_ffn_base, &post, &comb, il);
 
         cur = build_hc_pre(inpL, attn_pre, il);
         cb(cur, "hc_ffn_pre", il);
@@ -889,12 +877,9 @@ llama_model_deepseek41::graph::graph(const llama_model & model_base, const llm_g
         cur = build_norm(cur, layer.ffn_norm, nullptr, LLM_NORM_RMS, il);
         cb(cur, "ffn_norm", il);
 
-        // F32 router inputs: expert selection is near-tie sensitive
+        // F32 router accumulation: expert selection is near-tie sensitive
         ggml_tensor * logits = build_lora_mm(layer.ffn_gate_inp, cur);
         ggml_prec_set_acc(logits, GGML_PREC_F32);
-        if (!ggml_is_quantized(layer.ffn_gate_inp->type)) {
-            ggml_prec_set_src(logits, GGML_PREC_F32, 1);
-        }
         cb(logits, "ffn_moe_logits", il);
 
         // image (embedding) tokens use the vision router bias, as DeepSeek-V4

src/models/dflash.cpp

@@ -1406,7 +1406,7 @@ llama_model_dflash::graph_dsv41::graph_dsv41(const llama_model & model, const ll
         ggml_tensor * post = nullptr;
         ggml_tensor * comb = nullptr;
 
-        ggml_tensor * attn_pre = build_hc_mixes(inpL, layer.hc_attn_fn, layer.hc_attn_scale, layer.hc_attn_base, &post, &comb, false, il);
+        ggml_tensor * attn_pre = build_hc_mixes(inpL, layer.hc_attn_fn, layer.hc_attn_scale, layer.hc_attn_base, &post, &comb, il);
 
         ggml_tensor * cur;
         if (pre_prev == nullptr) {
@@ -1465,7 +1465,7 @@ llama_model_dflash::graph_dsv41::graph_dsv41(const llama_model & model, const ll
         cb(inpL, "hc_attn_post", il);
 
         residual = inpL;
-        ggml_tensor * ffn_pre = build_hc_mixes(inpL, layer.hc_ffn_fn, layer.hc_ffn_scale, layer.hc_ffn_base, &post, &comb, false, il);
+        ggml_tensor * ffn_pre = build_hc_mixes(inpL, layer.hc_ffn_fn, layer.hc_ffn_scale, layer.hc_ffn_base, &post, &comb, il);
 
         cur = build_hc_pre(inpL, attn_pre, il);
         cb(cur, "hc_ffn_pre", il);

src/models/models.h

@@ -1376,7 +1376,6 @@ struct llama_model_deepseek4 : public llama_model_base {
                 ggml_tensor * hc_base,
                 ggml_tensor ** post,
                 ggml_tensor ** comb,
-                bool f32_src,
                 int il) const;
 
         ggml_tensor * build_hc_post(
@@ -1498,11 +1497,7 @@ struct llama_model_deepseek41 : public llama_model_deepseek4 {
     void load_arch_hparams(llama_model_loader & ml) override;
     void load_arch_tensors(llama_model_loader & ml) override;
 
-    // Engram: the hash constants are hparams.ple_* (one module per Engram layer), hashed on the host (ggml has no
-    // int64 mul/xor/mod); the token map and its pad id are V4.1 only
-    uint32_t engram_pad_id = 0;
-    std::vector<int32_t> engram_token_map;
-
+    // Engram: the hash constants are hparams.ple_* (one module per Engram layer), hashed on the host (ggml has no int64 mul/xor/mod)
     int engram_index(int il) const;
 
     void engram_hash(int e, const llama_ubatch & ubatch, const std::vector<llama_token> & prev, std::vector<int32_t> & out) const;
@@ -2741,6 +2736,7 @@ struct llama_model_qwen4exp : public llama_model_base {
                     ggml_tensor * w_down,
                     ggml_tensor * w_up,
                     ggml_tensor * w_inject,
+                    ggml_tensor * w_down_inject, // [w_down | w_inject] in one weight, can be null
                     ggml_tensor ** inject,
                             int   il);
 

src/models/qwen4exp.cpp

@@ -231,15 +231,32 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
         const int64_t value_dim  = head_v_dim * n_v_heads;
         const int64_t conv_dim   = key_dim * 2 + value_dim;
 
+        // matmuls that read the same input: try the merged [a | b] weight the converter writes, else the separate ones
+        auto create_merged = [&](llm_tensor t_merged, llm_tensor t_a, llm_tensor t_b, int64_t n_in, int64_t n_a, int64_t n_b,
+                ggml_tensor *& merged, ggml_tensor *& a, ggml_tensor *& b) {
+            if (flags & TENSOR_SKIP) {
+                const int skip = TENSOR_NOT_REQUIRED | TENSOR_SKIP;
+                create_tensor(tn(t_merged, "weight", il), { n_in, n_a + n_b }, skip | TENSOR_SKIP_IF_VIRTUAL);
+                create_tensor(tn(t_a,      "weight", il), { n_in, n_a },       skip);
+                create_tensor(tn(t_b,      "weight", il), { n_in, n_b },       skip);
+                return;
+            }
+            merged = create_tensor(tn(t_merged, "weight", il), { n_in, n_a + n_b }, flags | TENSOR_NOT_REQUIRED | TENSOR_SKIP_IF_VIRTUAL);
+            if (!merged) {
+                a = create_tensor(tn(t_a, "weight", il), { n_in, n_a }, flags);
+                b = create_tensor(tn(t_b, "weight", il), { n_in, n_b }, flags);
+            }
+        };
+
         // two HC modules per layer: before the token mixer, before the MoE
         layer.hc_attn_norm   = create_tensor(tn(LLM_TENSOR_HC_ATTN_NORM,   "weight", il), { n_embd, hc }, flags | TENSOR_ALLOW_RESHAPE);
-        layer.hc_attn_down   = create_tensor(tn(LLM_TENSOR_HC_ATTN_DOWN,   "weight", il), { hc_dim, hc_lr }, flags);
         layer.hc_attn_up     = create_tensor(tn(LLM_TENSOR_HC_ATTN_UP,     "weight", il), { hc_lr, hc_dim }, flags);
-        layer.hc_attn_inject = create_tensor(tn(LLM_TENSOR_HC_ATTN_INJECT, "weight", il), { hc_dim, hc }, flags);
+        create_merged(LLM_TENSOR_HC_ATTN_DOWN_INJECT, LLM_TENSOR_HC_ATTN_DOWN, LLM_TENSOR_HC_ATTN_INJECT, hc_dim, hc_lr, hc,
+                layer.hc_attn_down_inject, layer.hc_attn_down, layer.hc_attn_inject);
         layer.hc_ffn_norm    = create_tensor(tn(LLM_TENSOR_HC_FFN_NORM,    "weight", il), { n_embd, hc }, flags | TENSOR_ALLOW_RESHAPE);
-        layer.hc_ffn_down    = create_tensor(tn(LLM_TENSOR_HC_FFN_DOWN,    "weight", il), { hc_dim, hc_lr }, flags);
         layer.hc_ffn_up      = create_tensor(tn(LLM_TENSOR_HC_FFN_UP,      "weight", il), { hc_lr, hc_dim }, flags);
-        layer.hc_ffn_inject  = create_tensor(tn(LLM_TENSOR_HC_FFN_INJECT,  "weight", il), { hc_dim, hc }, flags);
+        create_merged(LLM_TENSOR_HC_FFN_DOWN_INJECT, LLM_TENSOR_HC_FFN_DOWN, LLM_TENSOR_HC_FFN_INJECT, hc_dim, hc_lr, hc,
+                layer.hc_ffn_down_inject, layer.hc_ffn_down, layer.hc_ffn_inject);
 
         if (!hparams.is_recr(il)) {
             // full attention: wq holds [q|gate] interleaved per head
@@ -255,13 +272,13 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
             layer.index_q_norm = create_tensor(tn(LLM_TENSOR_INDEXER_Q_NORM, "weight", il), { idx_dim }, flags);
             layer.index_k_norm = create_tensor(tn(LLM_TENSOR_INDEXER_K_NORM, "weight", il), { idx_dim }, flags);
         } else {
-            layer.wqkv       = create_tensor(tn(LLM_TENSOR_ATTN_QKV,   "weight", il), { n_embd, key_dim * 2 + value_dim }, flags);
-            layer.wqkv_gate  = create_tensor(tn(LLM_TENSOR_ATTN_GATE,  "weight", il), { n_embd, value_dim }, flags);
+            create_merged(LLM_TENSOR_ATTN_QKV_GATE, LLM_TENSOR_ATTN_QKV, LLM_TENSOR_ATTN_GATE, n_embd, key_dim * 2 + value_dim, value_dim,
+                    layer.wqkv_z, layer.wqkv, layer.wqkv_gate);
             layer.ssm_conv1d = create_tensor(tn(LLM_TENSOR_SSM_CONV1D, "weight", il), { hparams.ssm_d_conv, conv_dim }, flags);
             layer.ssm_dt     = create_tensor(tn(LLM_TENSOR_SSM_DT,     "bias",   il), { hparams.ssm_dt_rank }, flags);
             layer.ssm_a      = create_tensor(tn(LLM_TENSOR_SSM_A_NOSCAN,         il), { hparams.ssm_dt_rank }, flags);
-            layer.ssm_beta   = create_tensor(tn(LLM_TENSOR_SSM_BETA,   "weight", il), { n_embd, n_v_heads }, flags);
-            layer.ssm_alpha  = create_tensor(tn(LLM_TENSOR_SSM_ALPHA,  "weight", il), { n_embd, n_v_heads }, flags);
+            create_merged(LLM_TENSOR_SSM_BETA_ALPHA_CAT, LLM_TENSOR_SSM_BETA, LLM_TENSOR_SSM_ALPHA, n_embd, n_v_heads, n_v_heads,
+                    layer.ssm_beta_alpha_cat, layer.ssm_beta, layer.ssm_alpha);
             layer.ssm_norm   = create_tensor(tn(LLM_TENSOR_SSM_NORM,   "weight", il), { head_v_dim }, flags);
             layer.ssm_out    = create_tensor(tn(LLM_TENSOR_SSM_OUT,    "weight", il), { value_dim, n_embd }, flags);
         }
@@ -333,6 +350,7 @@ ggml_tensor * llama_model_qwen4exp::graph::build_hc_mix(
         ggml_tensor *  w_down,
         ggml_tensor *  w_up,
         ggml_tensor *  w_inject,
+        ggml_tensor *  w_down_inject,
         ggml_tensor ** inject,
         int            il) {
     const int64_t hc     = hparams.dsv4_hc_mult;
@@ -345,8 +363,20 @@ ggml_tensor * llama_model_qwen4exp::graph::build_hc_mix(
     xn = ggml_reshape_2d(ctx0, xn, hc_dim, nt);
     cb(xn, "hc_norm", il);
 
-    ggml_tensor * lo = build_lora_mm(w_down, xn);
-    lo = ggml_silu(ctx0, ggml_scale(ctx0, lo, 1.0f / (float) hc));
+    ggml_tensor * lo = nullptr;
+    if (inject && w_down_inject) {
+        // one matmul for both: the hc_down rows, then the hc_inject rows
+        const int64_t n_lo = w_down_inject->ne[1] - hc;
+        ggml_tensor * lo_inject = build_lora_mm(w_down_inject, xn);
+        // scale the combined rows before the split; both lo and inject use 1/hc
+        lo_inject = ggml_scale(ctx0, lo_inject, 1.0f / (float) hc);
+        lo      = ggml_view_2d(ctx0, lo_inject, n_lo, nt, lo_inject->nb[1], 0);
+        *inject = ggml_view_2d(ctx0, lo_inject, hc, nt, lo_inject->nb[1], ggml_row_size(lo_inject->type, n_lo));
+        cb(*inject, "hc_inject", il);
+    } else {
+        lo = ggml_scale(ctx0, build_lora_mm(w_down, xn), 1.0f / (float) hc);
+    }
+    lo = ggml_silu(ctx0, lo);
     ggml_tensor * gate = build_lora_mm(w_up, lo);
     cb(gate, "hc_gate", il);
 
@@ -375,8 +405,8 @@ ggml_tensor * llama_model_qwen4exp::graph::build_hc_mix(
     }
     cb(mixed, "hc_mixed", il);
 
-    if (inject) {
-        *inject = build_lora_mm(w_inject, xn);
+    if (inject && !w_down_inject) {
+        *inject = ggml_scale(ctx0, build_lora_mm(w_inject, xn), 1.0f / (float) hc);
         cb(*inject, "hc_inject", il);
     }
 
@@ -392,7 +422,7 @@ ggml_tensor * llama_model_qwen4exp::graph::build_hc_combine(
     const int64_t nt = residual->ne[2];
 
     // 2*sigmoid centres the scatter weights on 1, so a zero injection is a plain residual add
-    ggml_tensor * w = ggml_sigmoid(ctx0, ggml_scale(ctx0, inject, 1.0f / (float) hc));
+    ggml_tensor * w = ggml_sigmoid(ctx0, inject);
     w = ggml_scale(ctx0, w, 2.0f);
 
     ggml_tensor * cur = nullptr;
@@ -480,6 +510,7 @@ llama_model_qwen4exp::graph::graph(const llama_model & model, const llm_graph_pa
                 model.layers[il].hc_attn_down,
                 model.layers[il].hc_attn_up,
                 model.layers[il].hc_attn_inject,
+                model.layers[il].hc_attn_down_inject,
                 &inject, il);
 
         ggml_build_forward_expand(gf, cur);
@@ -507,6 +538,7 @@ llama_model_qwen4exp::graph::graph(const llama_model & model, const llm_graph_pa
                 model.layers[il].hc_ffn_down,
                 model.layers[il].hc_ffn_up,
                 model.layers[il].hc_ffn_inject,
+                model.layers[il].hc_ffn_down_inject,
                 &inject, il);
 
         cur = build_layer_ffn(cur, il);
@@ -532,7 +564,7 @@ llama_model_qwen4exp::graph::graph(const llama_model & model, const llm_graph_pa
     // the final mixer is the output norm: there is no separate one
     ggml_tensor * cur = build_hc_mix(res_hc,
             model.hc_head_norm, model.hc_head_down, model.hc_head_up,
-            nullptr, nullptr, -1);
+            nullptr, nullptr, nullptr, -1);
 
     cb(cur, "result_norm", -1);
     res->t_embd = cur;
@@ -610,11 +642,11 @@ llama_model_qwen4exp::graph_mtp::graph_mtp(const llama_model & model, const llm_
     cb(res_hc, "mtp_eh_proj", il);
 
     ggml_tensor * inject = nullptr;
-    ggml_tensor * cur = build_hc_mix(res_hc, layer.hc_attn_norm, layer.hc_attn_down, layer.hc_attn_up, layer.hc_attn_inject, &inject, il);
+    ggml_tensor * cur = build_hc_mix(res_hc, layer.hc_attn_norm, layer.hc_attn_down, layer.hc_attn_up, layer.hc_attn_inject, layer.hc_attn_down_inject, &inject, il);
     cur    = build_layer_attn(inp_hyb->get_attn(), mctx_hyb, inp_kpool, cur, inp_pos, sections, il);
     res_hc = build_hc_combine(res_hc, cur, inject, il);
 
-    cur    = build_hc_mix(res_hc, layer.hc_ffn_norm, layer.hc_ffn_down, layer.hc_ffn_up, layer.hc_ffn_inject, &inject, il);
+    cur    = build_hc_mix(res_hc, layer.hc_ffn_norm, layer.hc_ffn_down, layer.hc_ffn_up, layer.hc_ffn_inject, layer.hc_ffn_down_inject, &inject, il);
     cur    = build_layer_ffn(cur, il);
     res_hc = build_hc_combine(res_hc, cur, inject, il);
 
@@ -627,7 +659,7 @@ llama_model_qwen4exp::graph_mtp::graph_mtp(const llama_model & model, const llm_
 
     cur = build_hc_mix(ggml_reshape_3d(ctx0, flat_out, n_embd, hc, flat_out->ne[1]),
             layer.nextn.hc_head_norm, layer.nextn.hc_head_down, layer.nextn.hc_head_up,
-            nullptr, nullptr, il);
+            nullptr, nullptr, nullptr, il);
     cb(cur, "result_norm", -1);
     res->t_embd = cur;
 
@@ -642,6 +674,23 @@ std::pair<ggml_tensor *, ggml_tensor *> llama_model_qwen4exp::graph::build_qkvz(
     const int64_t n_seqs       = ubatch.n_seqs;
     const int64_t n_seq_tokens = ubatch.n_seq_tokens;
 
+    if (model.layers[il].wqkv_z) {
+        // qkv and z read the same input: one matmul, then the qkv rows and the z rows
+        ggml_tensor * qkvz = build_lora_mm(model.layers[il].wqkv_z, input);
+        cb(qkvz, "linear_attn_qkvz", il);
+
+        const int64_t n_z   = hparams.ssm_d_state * hparams.ssm_dt_rank;
+        const int64_t n_qkv = qkvz->ne[0] - n_z;
+
+        ggml_tensor * qkv_mixed = ggml_view_3d(ctx0, qkvz, n_qkv, n_seq_tokens, n_seqs, qkvz->nb[1], qkvz->nb[1] * n_seq_tokens, 0);
+        cb(qkv_mixed, "linear_attn_qkv_mixed", il);
+
+        ggml_tensor * z = ggml_view_2d(ctx0, qkvz, n_z, n_seq_tokens * n_seqs, qkvz->nb[1], ggml_row_size(qkvz->type, n_qkv));
+        cb(z, "z", il);
+
+        return { qkv_mixed, z };
+    }
+
     ggml_tensor * qkv_mixed = build_lora_mm(model.layers[il].wqkv, input, model.layers[il].wqkv_s);
     qkv_mixed = ggml_reshape_3d(ctx0, qkv_mixed, qkv_mixed->ne[0], n_seq_tokens, n_seqs);
     cb(qkv_mixed, "linear_attn_qkv_mixed", il);
@@ -1072,17 +1121,30 @@ ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn_linear(
     ggml_tensor * qkv_mixed = qkvz.first;
     ggml_tensor * z         = qkvz.second;
 
-    ggml_tensor * beta = build_lora_mm(model.layers[il].ssm_beta, cur, model.layers[il].ssm_beta_s);
-    beta = ggml_reshape_4d(ctx0, beta, 1, num_v_heads, n_seq_tokens, n_seqs);
+    ggml_tensor * beta  = nullptr;
+    ggml_tensor * alpha = nullptr;
+    if (model.layers[il].ssm_beta_alpha_cat) {
+        // beta and alpha read the same input: one matmul, then the beta rows and the alpha rows
+        ggml_tensor * ba = build_lora_mm(model.layers[il].ssm_beta_alpha_cat, cur);
+        cb(ba, "beta_alpha", il);
+
+        beta  = ggml_view_4d(ctx0, ba, 1, num_v_heads, n_seq_tokens, n_seqs,
+                ggml_element_size(ba), ba->nb[1], ba->nb[1] * n_seq_tokens, 0);
+        alpha = ggml_view_3d(ctx0, ba, num_v_heads, n_seq_tokens, n_seqs,
+                ba->nb[1], ba->nb[1] * n_seq_tokens, ggml_row_size(ba->type, num_v_heads));
+    } else {
+        beta = build_lora_mm(model.layers[il].ssm_beta, cur, model.layers[il].ssm_beta_s);
+        beta = ggml_reshape_4d(ctx0, beta, 1, num_v_heads, n_seq_tokens, n_seqs);
+
+        alpha = build_lora_mm(model.layers[il].ssm_alpha, cur, model.layers[il].ssm_alpha_s);
+        alpha = ggml_reshape_3d(ctx0, alpha, num_v_heads, n_seq_tokens, n_seqs);
+    }
     cb(beta, "beta", il);
+    cb(alpha, "alpha", il);
 
     beta = ggml_sigmoid(ctx0, beta);
     cb(beta, "beta_sigmoid", il);
 
-    ggml_tensor * alpha = build_lora_mm(model.layers[il].ssm_alpha, cur, model.layers[il].ssm_alpha_s);
-    alpha = ggml_reshape_3d(ctx0, alpha, num_v_heads, n_seq_tokens, n_seqs);
-    cb(alpha, "alpha", il);
-
     ggml_tensor * alpha_biased   = ggml_add(ctx0, alpha, model.layers[il].ssm_dt);
     ggml_tensor * alpha_softplus = ggml_softplus(ctx0, alpha_biased);
     cb(alpha_softplus, "a_softplus", il);
@@ -1160,7 +1222,8 @@ ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn_linear(
 
     ggml_tensor * output = build_recurrent_attn(inp, ssm_states_all, q_conv, k_conv, v_conv, gate, beta, state, il);
 
-    ggml_tensor * z_2d = ggml_reshape_4d(ctx0, z, head_v_dim, num_v_heads, n_seq_tokens, n_seqs);
+    ggml_tensor * z_2d = ggml_view_4d(ctx0, z, head_v_dim, num_v_heads, n_seq_tokens, n_seqs,
+            ggml_row_size(z->type, head_v_dim), z->nb[1], z->nb[1] * n_seq_tokens, 0);
 
     // gated normalization, as self.norm(core_attn_out, z) in the reference
     ggml_tensor * attn_out_norm = build_norm_gated(output, model.layers[il].ssm_norm, z_2d, il);

tests/CMakeLists.txt

@@ -223,58 +223,57 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
 
     # llama_build_and_test(test-double-float.cpp) # SLOW
 
-    # TODO: after `test-llama-archs` is refactored to not link internal symbols, we can move all related tests out of
-    #       this branch and run them on all platforms
-    llama_build(test-llama-archs.cpp)
-
-    set(MODEL_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-models/")
-    file(MAKE_DIRECTORY "${MODEL_DIR}")
-
-    llama_test(
-        test-llama-archs
-        NAME test-generate-models
-        LABEL "main behavior backend"
-        ARGS -o "${MODEL_DIR}"
-    )
-    set_tests_properties(test-generate-models PROPERTIES
-        FIXTURES_SETUP generate-models
-    )
-
-    # Target model for test-speculative-rejection-sampling
-    set(SPEC_TARGET_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-speculative-models/target")
-    file(MAKE_DIRECTORY "${SPEC_TARGET_DIR}")
-    llama_test(
-        test-llama-archs
-        NAME test-generate-speculative-target
-        LABEL main
-        ARGS -a llama -s 999 -o "${SPEC_TARGET_DIR}"
-    )
-    set_tests_properties(
-        test-generate-speculative-target
-        PROPERTIES FIXTURES_SETUP speculative-model
-    )
-
-    # To avoid recreating the model use the fixture exclude any option:
-    # -FA "^speculative-model$"
-    llama_build_and_test(test-speculative-rejection-sampling.cpp ARGS
-        -m "${SPEC_TARGET_DIR}/llama-dense.gguf")
-    set_tests_properties(test-speculative-rejection-sampling PROPERTIES FIXTURES_REQUIRED speculative-model)
-
-    # Test context state handling across all architectures, using the generated dummy models
-    llama_test(
-        test-llama-context
-        LABEL "main behavior backend"
-        ARGS --models "${MODEL_DIR}"
-    )
-    # abort on graph reallocations that the initial reserve should have covered
-    set_tests_properties(test-llama-context PROPERTIES
-        FIXTURES_REQUIRED generate-models
-        ENVIRONMENT "GGML_SCHED_DEBUG_REALLOC=1"
-    )
-
-    llama_build(test-fusion.cpp)
 endif()
 
+llama_build(test-llama-archs.cpp)
+
+set(MODEL_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-models/")
+file(MAKE_DIRECTORY "${MODEL_DIR}")
+
+llama_test(
+    test-llama-archs
+    NAME test-generate-models
+    LABEL "main behavior backend"
+    ARGS -o "${MODEL_DIR}"
+)
+set_tests_properties(test-generate-models PROPERTIES
+    FIXTURES_SETUP generate-models
+)
+
+# Target model for test-speculative-rejection-sampling
+set(SPEC_TARGET_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-speculative-models/target")
+file(MAKE_DIRECTORY "${SPEC_TARGET_DIR}")
+llama_test(
+    test-llama-archs
+    NAME test-generate-speculative-target
+    LABEL main
+    ARGS -a llama -s 999 -o "${SPEC_TARGET_DIR}"
+)
+set_tests_properties(
+    test-generate-speculative-target
+    PROPERTIES FIXTURES_SETUP speculative-model
+)
+
+# To avoid recreating the model use the fixture exclude any option:
+# -FA "^speculative-model$"
+llama_build_and_test(test-speculative-rejection-sampling.cpp ARGS
+    -m "${SPEC_TARGET_DIR}/llama-dense.gguf")
+set_tests_properties(test-speculative-rejection-sampling PROPERTIES FIXTURES_REQUIRED speculative-model)
+
+# Test context state handling across all architectures, using the generated dummy models
+llama_test(
+    test-llama-context
+    LABEL "main behavior backend"
+    ARGS --models "${MODEL_DIR}"
+)
+# abort on graph reallocations that the initial reserve should have covered
+set_tests_properties(test-llama-context PROPERTIES
+    FIXTURES_REQUIRED generate-models
+    ENVIRONMENT "GGML_SCHED_DEBUG_REALLOC=1"
+)
+
+llama_build(test-fusion.cpp)
+
 llama_build_and_test(test-chat-peg-parser.cpp peg-parser/simple-tokenize.cpp  LABEL behavior)
 llama_build_and_test(test-jinja.cpp  LABEL behavior)
 llama_test(test-jinja  NAME test-jinja-py  ARGS -py  LABEL "python behavior")

tests/peg-parser/test-json-parser.cpp

@@ -130,4 +130,23 @@ void test_json_parser(testing &t) {
             t.assert_true("fail", result.fail());
         });
     });
+
+    // json-value is entered once per level
+    t.test("deep JSON", [](testing &t) {
+        auto json = build_peg_parser([](common_peg_parser_builder & p) { return p.json(); });
+
+        const auto parse = [&](size_t depth) {
+            std::string input = std::string(depth, '[') + std::string(depth, ']');
+            common_peg_parse_context ctx(input);
+            try {
+                return json.parse(ctx).success() ? std::string() : std::string("fail");
+            } catch (const common_peg_recursion_error & e) {
+                return std::string(e.what());
+            }
+        };
+
+        t.assert_equal("64 levels", std::string(), parse(64));
+        t.assert_true("65 levels", parse(65).find("nested more than 64 levels deep") != std::string::npos);
+        t.assert_true("100000 levels", parse(100000).find("nested more than 64 levels deep") != std::string::npos);
+    });
 }

tests/snapshots/deepseek-v3.1.schema

@@ -2927,430 +2927,552 @@ blk.60.attn_k_b.weight q4_0
 [MXFP4_MOE] mxfp4
 output.weight q8_0
 token_embd.weight q8_0
+blk.0.attn_k_b.weight q8_0
 blk.0.attn_kv_a_mqa.weight q8_0
 blk.0.attn_output.weight q8_0
 blk.0.attn_q_a.weight q8_0
 blk.0.attn_q_b.weight q8_0
+blk.0.attn_v_b.weight q8_0
 blk.0.ffn_down.weight q8_0
 blk.0.ffn_gate.weight q8_0
 blk.0.ffn_up.weight q8_0
+blk.1.attn_k_b.weight q8_0
 blk.1.attn_kv_a_mqa.weight q8_0
 blk.1.attn_output.weight q8_0
 blk.1.attn_q_a.weight q8_0
 blk.1.attn_q_b.weight q8_0
+blk.1.attn_v_b.weight q8_0
 blk.1.ffn_down.weight q8_0
 blk.1.ffn_gate.weight q8_0
 blk.1.ffn_up.weight q8_0
+blk.2.attn_k_b.weight q8_0
 blk.2.attn_kv_a_mqa.weight q8_0
 blk.2.attn_output.weight q8_0
 blk.2.attn_q_a.weight q8_0
 blk.2.attn_q_b.weight q8_0
+blk.2.attn_v_b.weight q8_0
 blk.2.ffn_down.weight q8_0
 blk.2.ffn_gate.weight q8_0
 blk.2.ffn_up.weight q8_0
+blk.3.attn_k_b.weight q8_0
 blk.3.attn_kv_a_mqa.weight q8_0
 blk.3.attn_output.weight q8_0
 blk.3.attn_q_a.weight q8_0
 blk.3.attn_q_b.weight q8_0
+blk.3.attn_v_b.weight q8_0
 blk.3.ffn_down_shexp.weight q8_0
 blk.3.ffn_gate_shexp.weight q8_0
 blk.3.ffn_up_shexp.weight q8_0
+blk.4.attn_k_b.weight q8_0
 blk.4.attn_kv_a_mqa.weight q8_0
 blk.4.attn_output.weight q8_0
 blk.4.attn_q_a.weight q8_0
 blk.4.attn_q_b.weight q8_0
+blk.4.attn_v_b.weight q8_0
 blk.4.ffn_down_shexp.weight q8_0
 blk.4.ffn_gate_shexp.weight q8_0
 blk.4.ffn_up_shexp.weight q8_0
+blk.5.attn_k_b.weight q8_0
 blk.5.attn_kv_a_mqa.weight q8_0
 blk.5.attn_output.weight q8_0
 blk.5.attn_q_a.weight q8_0
 blk.5.attn_q_b.weight q8_0
+blk.5.attn_v_b.weight q8_0
 blk.5.ffn_down_shexp.weight q8_0
 blk.5.ffn_gate_shexp.weight q8_0
 blk.5.ffn_up_shexp.weight q8_0
+blk.6.attn_k_b.weight q8_0
 blk.6.attn_kv_a_mqa.weight q8_0
 blk.6.attn_output.weight q8_0
 blk.6.attn_q_a.weight q8_0
 blk.6.attn_q_b.weight q8_0
+blk.6.attn_v_b.weight q8_0
 blk.6.ffn_down_shexp.weight q8_0
 blk.6.ffn_gate_shexp.weight q8_0
 blk.6.ffn_up_shexp.weight q8_0
+blk.7.attn_k_b.weight q8_0
 blk.7.attn_kv_a_mqa.weight q8_0
 blk.7.attn_output.weight q8_0
 blk.7.attn_q_a.weight q8_0
 blk.7.attn_q_b.weight q8_0
+blk.7.attn_v_b.weight q8_0
 blk.7.ffn_down_shexp.weight q8_0
 blk.7.ffn_gate_shexp.weight q8_0
 blk.7.ffn_up_shexp.weight q8_0
+blk.8.attn_k_b.weight q8_0
 blk.8.attn_kv_a_mqa.weight q8_0
 blk.8.attn_output.weight q8_0
 blk.8.attn_q_a.weight q8_0
 blk.8.attn_q_b.weight q8_0
+blk.8.attn_v_b.weight q8_0
 blk.8.ffn_down_shexp.weight q8_0
 blk.8.ffn_gate_shexp.weight q8_0
 blk.8.ffn_up_shexp.weight q8_0
+blk.9.attn_k_b.weight q8_0
 blk.9.attn_kv_a_mqa.weight q8_0
 blk.9.attn_output.weight q8_0
 blk.9.attn_q_a.weight q8_0
 blk.9.attn_q_b.weight q8_0
+blk.9.attn_v_b.weight q8_0
 blk.9.ffn_down_shexp.weight q8_0
 blk.9.ffn_gate_shexp.weight q8_0
 blk.9.ffn_up_shexp.weight q8_0
+blk.10.attn_k_b.weight q8_0
 blk.10.attn_kv_a_mqa.weight q8_0
 blk.10.attn_output.weight q8_0
 blk.10.attn_q_a.weight q8_0
 blk.10.attn_q_b.weight q8_0
+blk.10.attn_v_b.weight q8_0
 blk.10.ffn_down_shexp.weight q8_0
 blk.10.ffn_gate_shexp.weight q8_0
 blk.10.ffn_up_shexp.weight q8_0
+blk.11.attn_k_b.weight q8_0
 blk.11.attn_kv_a_mqa.weight q8_0
 blk.11.attn_output.weight q8_0
 blk.11.attn_q_a.weight q8_0
 blk.11.attn_q_b.weight q8_0
+blk.11.attn_v_b.weight q8_0
 blk.11.ffn_down_shexp.weight q8_0
 blk.11.ffn_gate_shexp.weight q8_0
 blk.11.ffn_up_shexp.weight q8_0
+blk.12.attn_k_b.weight q8_0
 blk.12.attn_kv_a_mqa.weight q8_0
 blk.12.attn_output.weight q8_0
 blk.12.attn_q_a.weight q8_0
 blk.12.attn_q_b.weight q8_0
+blk.12.attn_v_b.weight q8_0
 blk.12.ffn_down_shexp.weight q8_0
 blk.12.ffn_gate_shexp.weight q8_0
 blk.12.ffn_up_shexp.weight q8_0
+blk.13.attn_k_b.weight q8_0
 blk.13.attn_kv_a_mqa.weight q8_0
 blk.13.attn_output.weight q8_0
 blk.13.attn_q_a.weight q8_0
 blk.13.attn_q_b.weight q8_0
+blk.13.attn_v_b.weight q8_0
 blk.13.ffn_down_shexp.weight q8_0
 blk.13.ffn_gate_shexp.weight q8_0
 blk.13.ffn_up_shexp.weight q8_0
+blk.14.attn_k_b.weight q8_0
 blk.14.attn_kv_a_mqa.weight q8_0
 blk.14.attn_output.weight q8_0
 blk.14.attn_q_a.weight q8_0
 blk.14.attn_q_b.weight q8_0
+blk.14.attn_v_b.weight q8_0
 blk.14.ffn_down_shexp.weight q8_0
 blk.14.ffn_gate_shexp.weight q8_0
 blk.14.ffn_up_shexp.weight q8_0
+blk.15.attn_k_b.weight q8_0
 blk.15.attn_kv_a_mqa.weight q8_0
 blk.15.attn_output.weight q8_0
 blk.15.attn_q_a.weight q8_0
 blk.15.attn_q_b.weight q8_0
+blk.15.attn_v_b.weight q8_0
 blk.15.ffn_down_shexp.weight q8_0
 blk.15.ffn_gate_shexp.weight q8_0
 blk.15.ffn_up_shexp.weight q8_0
+blk.16.attn_k_b.weight q8_0
 blk.16.attn_kv_a_mqa.weight q8_0
 blk.16.attn_output.weight q8_0
 blk.16.attn_q_a.weight q8_0
 blk.16.attn_q_b.weight q8_0
+blk.16.attn_v_b.weight q8_0
 blk.16.ffn_down_shexp.weight q8_0
 blk.16.ffn_gate_shexp.weight q8_0
 blk.16.ffn_up_shexp.weight q8_0
+blk.17.attn_k_b.weight q8_0
 blk.17.attn_kv_a_mqa.weight q8_0
 blk.17.attn_output.weight q8_0
 blk.17.attn_q_a.weight q8_0
 blk.17.attn_q_b.weight q8_0
+blk.17.attn_v_b.weight q8_0
 blk.17.ffn_down_shexp.weight q8_0
 blk.17.ffn_gate_shexp.weight q8_0
 blk.17.ffn_up_shexp.weight q8_0
+blk.18.attn_k_b.weight q8_0
 blk.18.attn_kv_a_mqa.weight q8_0
 blk.18.attn_output.weight q8_0
 blk.18.attn_q_a.weight q8_0
 blk.18.attn_q_b.weight q8_0
+blk.18.attn_v_b.weight q8_0
 blk.18.ffn_down_shexp.weight q8_0
 blk.18.ffn_gate_shexp.weight q8_0
 blk.18.ffn_up_shexp.weight q8_0
+blk.19.attn_k_b.weight q8_0
 blk.19.attn_kv_a_mqa.weight q8_0
 blk.19.attn_output.weight q8_0
 blk.19.attn_q_a.weight q8_0
 blk.19.attn_q_b.weight q8_0
+blk.19.attn_v_b.weight q8_0
 blk.19.ffn_down_shexp.weight q8_0
 blk.19.ffn_gate_shexp.weight q8_0
 blk.19.ffn_up_shexp.weight q8_0
+blk.20.attn_k_b.weight q8_0
 blk.20.attn_kv_a_mqa.weight q8_0
 blk.20.attn_output.weight q8_0
 blk.20.attn_q_a.weight q8_0
 blk.20.attn_q_b.weight q8_0
+blk.20.attn_v_b.weight q8_0
 blk.20.ffn_down_shexp.weight q8_0
 blk.20.ffn_gate_shexp.weight q8_0
 blk.20.ffn_up_shexp.weight q8_0
+blk.21.attn_k_b.weight q8_0
 blk.21.attn_kv_a_mqa.weight q8_0
 blk.21.attn_output.weight q8_0
 blk.21.attn_q_a.weight q8_0
 blk.21.attn_q_b.weight q8_0
+blk.21.attn_v_b.weight q8_0
 blk.21.ffn_down_shexp.weight q8_0
 blk.21.ffn_gate_shexp.weight q8_0
 blk.21.ffn_up_shexp.weight q8_0
+blk.22.attn_k_b.weight q8_0
 blk.22.attn_kv_a_mqa.weight q8_0
 blk.22.attn_output.weight q8_0
 blk.22.attn_q_a.weight q8_0
 blk.22.attn_q_b.weight q8_0
+blk.22.attn_v_b.weight q8_0
 blk.22.ffn_down_shexp.weight q8_0
 blk.22.ffn_gate_shexp.weight q8_0
 blk.22.ffn_up_shexp.weight q8_0
+blk.23.attn_k_b.weight q8_0
 blk.23.attn_kv_a_mqa.weight q8_0
 blk.23.attn_output.weight q8_0
 blk.23.attn_q_a.weight q8_0
 blk.23.attn_q_b.weight q8_0
+blk.23.attn_v_b.weight q8_0
 blk.23.ffn_down_shexp.weight q8_0
 blk.23.ffn_gate_shexp.weight q8_0
 blk.23.ffn_up_shexp.weight q8_0
+blk.24.attn_k_b.weight q8_0
 blk.24.attn_kv_a_mqa.weight q8_0
 blk.24.attn_output.weight q8_0
 blk.24.attn_q_a.weight q8_0
 blk.24.attn_q_b.weight q8_0
+blk.24.attn_v_b.weight q8_0
 blk.24.ffn_down_shexp.weight q8_0
 blk.24.ffn_gate_shexp.weight q8_0
 blk.24.ffn_up_shexp.weight q8_0
+blk.25.attn_k_b.weight q8_0
 blk.25.attn_kv_a_mqa.weight q8_0
 blk.25.attn_output.weight q8_0
 blk.25.attn_q_a.weight q8_0
 blk.25.attn_q_b.weight q8_0
+blk.25.attn_v_b.weight q8_0
 blk.25.ffn_down_shexp.weight q8_0
 blk.25.ffn_gate_shexp.weight q8_0
 blk.25.ffn_up_shexp.weight q8_0
+blk.26.attn_k_b.weight q8_0
 blk.26.attn_kv_a_mqa.weight q8_0
 blk.26.attn_output.weight q8_0
 blk.26.attn_q_a.weight q8_0
 blk.26.attn_q_b.weight q8_0
+blk.26.attn_v_b.weight q8_0
 blk.26.ffn_down_shexp.weight q8_0
 blk.26.ffn_gate_shexp.weight q8_0
 blk.26.ffn_up_shexp.weight q8_0
+blk.27.attn_k_b.weight q8_0
 blk.27.attn_kv_a_mqa.weight q8_0
 blk.27.attn_output.weight q8_0
 blk.27.attn_q_a.weight q8_0
 blk.27.attn_q_b.weight q8_0
+blk.27.attn_v_b.weight q8_0
 blk.27.ffn_down_shexp.weight q8_0
 blk.27.ffn_gate_shexp.weight q8_0
 blk.27.ffn_up_shexp.weight q8_0
+blk.28.attn_k_b.weight q8_0
 blk.28.attn_kv_a_mqa.weight q8_0
 blk.28.attn_output.weight q8_0
 blk.28.attn_q_a.weight q8_0
 blk.28.attn_q_b.weight q8_0
+blk.28.attn_v_b.weight q8_0
 blk.28.ffn_down_shexp.weight q8_0
 blk.28.ffn_gate_shexp.weight q8_0
 blk.28.ffn_up_shexp.weight q8_0
+blk.29.attn_k_b.weight q8_0
 blk.29.attn_kv_a_mqa.weight q8_0
 blk.29.attn_output.weight q8_0
 blk.29.attn_q_a.weight q8_0
 blk.29.attn_q_b.weight q8_0
+blk.29.attn_v_b.weight q8_0
 blk.29.ffn_down_shexp.weight q8_0
 blk.29.ffn_gate_shexp.weight q8_0
 blk.29.ffn_up_shexp.weight q8_0
+blk.30.attn_k_b.weight q8_0
 blk.30.attn_kv_a_mqa.weight q8_0
 blk.30.attn_output.weight q8_0
 blk.30.attn_q_a.weight q8_0
 blk.30.attn_q_b.weight q8_0
+blk.30.attn_v_b.weight q8_0
 blk.30.ffn_down_shexp.weight q8_0
 blk.30.ffn_gate_shexp.weight q8_0
 blk.30.ffn_up_shexp.weight q8_0
+blk.31.attn_k_b.weight q8_0
 blk.31.attn_kv_a_mqa.weight q8_0
 blk.31.attn_output.weight q8_0
 blk.31.attn_q_a.weight q8_0
 blk.31.attn_q_b.weight q8_0
+blk.31.attn_v_b.weight q8_0
 blk.31.ffn_down_shexp.weight q8_0
 blk.31.ffn_gate_shexp.weight q8_0
 blk.31.ffn_up_shexp.weight q8_0
+blk.32.attn_k_b.weight q8_0
 blk.32.attn_kv_a_mqa.weight q8_0
 blk.32.attn_output.weight q8_0
 blk.32.attn_q_a.weight q8_0
 blk.32.attn_q_b.weight q8_0
+blk.32.attn_v_b.weight q8_0
 blk.32.ffn_down_shexp.weight q8_0
 blk.32.ffn_gate_shexp.weight q8_0
 blk.32.ffn_up_shexp.weight q8_0
+blk.33.attn_k_b.weight q8_0
 blk.33.attn_kv_a_mqa.weight q8_0
 blk.33.attn_output.weight q8_0
 blk.33.attn_q_a.weight q8_0
 blk.33.attn_q_b.weight q8_0
+blk.33.attn_v_b.weight q8_0
 blk.33.ffn_down_shexp.weight q8_0
 blk.33.ffn_gate_shexp.weight q8_0
 blk.33.ffn_up_shexp.weight q8_0
+blk.34.attn_k_b.weight q8_0
 blk.34.attn_kv_a_mqa.weight q8_0
 blk.34.attn_output.weight q8_0
 blk.34.attn_q_a.weight q8_0
 blk.34.attn_q_b.weight q8_0
+blk.34.attn_v_b.weight q8_0
 blk.34.ffn_down_shexp.weight q8_0
 blk.34.ffn_gate_shexp.weight q8_0
 blk.34.ffn_up_shexp.weight q8_0
+blk.35.attn_k_b.weight q8_0
 blk.35.attn_kv_a_mqa.weight q8_0
 blk.35.attn_output.weight q8_0
 blk.35.attn_q_a.weight q8_0
 blk.35.attn_q_b.weight q8_0
+blk.35.attn_v_b.weight q8_0
 blk.35.ffn_down_shexp.weight q8_0
 blk.35.ffn_gate_shexp.weight q8_0
 blk.35.ffn_up_shexp.weight q8_0
+blk.36.attn_k_b.weight q8_0
 blk.36.attn_kv_a_mqa.weight q8_0
 blk.36.attn_output.weight q8_0
 blk.36.attn_q_a.weight q8_0
 blk.36.attn_q_b.weight q8_0
+blk.36.attn_v_b.weight q8_0
 blk.36.ffn_down_shexp.weight q8_0
 blk.36.ffn_gate_shexp.weight q8_0
 blk.36.ffn_up_shexp.weight q8_0
+blk.37.attn_k_b.weight q8_0
 blk.37.attn_kv_a_mqa.weight q8_0
 blk.37.attn_output.weight q8_0
 blk.37.attn_q_a.weight q8_0
 blk.37.attn_q_b.weight q8_0
+blk.37.attn_v_b.weight q8_0
 blk.37.ffn_down_shexp.weight q8_0
 blk.37.ffn_gate_shexp.weight q8_0
 blk.37.ffn_up_shexp.weight q8_0
+blk.38.attn_k_b.weight q8_0
 blk.38.attn_kv_a_mqa.weight q8_0
 blk.38.attn_output.weight q8_0
 blk.38.attn_q_a.weight q8_0
 blk.38.attn_q_b.weight q8_0
+blk.38.attn_v_b.weight q8_0
 blk.38.ffn_down_shexp.weight q8_0
 blk.38.ffn_gate_shexp.weight q8_0
 blk.38.ffn_up_shexp.weight q8_0
+blk.39.attn_k_b.weight q8_0
 blk.39.attn_kv_a_mqa.weight q8_0
 blk.39.attn_output.weight q8_0
 blk.39.attn_q_a.weight q8_0
 blk.39.attn_q_b.weight q8_0
+blk.39.attn_v_b.weight q8_0
 blk.39.ffn_down_shexp.weight q8_0
 blk.39.ffn_gate_shexp.weight q8_0
 blk.39.ffn_up_shexp.weight q8_0
+blk.40.attn_k_b.weight q8_0
 blk.40.attn_kv_a_mqa.weight q8_0
 blk.40.attn_output.weight q8_0
 blk.40.attn_q_a.weight q8_0
 blk.40.attn_q_b.weight q8_0
+blk.40.attn_v_b.weight q8_0
 blk.40.ffn_down_shexp.weight q8_0
 blk.40.ffn_gate_shexp.weight q8_0
 blk.40.ffn_up_shexp.weight q8_0
+blk.41.attn_k_b.weight q8_0
 blk.41.attn_kv_a_mqa.weight q8_0
 blk.41.attn_output.weight q8_0
 blk.41.attn_q_a.weight q8_0
 blk.41.attn_q_b.weight q8_0
+blk.41.attn_v_b.weight q8_0
 blk.41.ffn_down_shexp.weight q8_0
 blk.41.ffn_gate_shexp.weight q8_0
 blk.41.ffn_up_shexp.weight q8_0
+blk.42.attn_k_b.weight q8_0
 blk.42.attn_kv_a_mqa.weight q8_0
 blk.42.attn_output.weight q8_0
 blk.42.attn_q_a.weight q8_0
 blk.42.attn_q_b.weight q8_0
+blk.42.attn_v_b.weight q8_0
 blk.42.ffn_down_shexp.weight q8_0
 blk.42.ffn_gate_shexp.weight q8_0
 blk.42.ffn_up_shexp.weight q8_0
+blk.43.attn_k_b.weight q8_0
 blk.43.attn_kv_a_mqa.weight q8_0
 blk.43.attn_output.weight q8_0
 blk.43.attn_q_a.weight q8_0
 blk.43.attn_q_b.weight q8_0
+blk.43.attn_v_b.weight q8_0
 blk.43.ffn_down_shexp.weight q8_0
 blk.43.ffn_gate_shexp.weight q8_0
 blk.43.ffn_up_shexp.weight q8_0
+blk.44.attn_k_b.weight q8_0
 blk.44.attn_kv_a_mqa.weight q8_0
 blk.44.attn_output.weight q8_0
 blk.44.attn_q_a.weight q8_0
 blk.44.attn_q_b.weight q8_0
+blk.44.attn_v_b.weight q8_0
 blk.44.ffn_down_shexp.weight q8_0
 blk.44.ffn_gate_shexp.weight q8_0
 blk.44.ffn_up_shexp.weight q8_0
+blk.45.attn_k_b.weight q8_0
 blk.45.attn_kv_a_mqa.weight q8_0
 blk.45.attn_output.weight q8_0
 blk.45.attn_q_a.weight q8_0
 blk.45.attn_q_b.weight q8_0
+blk.45.attn_v_b.weight q8_0
 blk.45.ffn_down_shexp.weight q8_0
 blk.45.ffn_gate_shexp.weight q8_0
 blk.45.ffn_up_shexp.weight q8_0
+blk.46.attn_k_b.weight q8_0
 blk.46.attn_kv_a_mqa.weight q8_0
 blk.46.attn_output.weight q8_0
 blk.46.attn_q_a.weight q8_0
 blk.46.attn_q_b.weight q8_0
+blk.46.attn_v_b.weight q8_0
 blk.46.ffn_down_shexp.weight q8_0
 blk.46.ffn_gate_shexp.weight q8_0
 blk.46.ffn_up_shexp.weight q8_0
+blk.47.attn_k_b.weight q8_0
 blk.47.attn_kv_a_mqa.weight q8_0
 blk.47.attn_output.weight q8_0
 blk.47.attn_q_a.weight q8_0
 blk.47.attn_q_b.weight q8_0
+blk.47.attn_v_b.weight q8_0
 blk.47.ffn_down_shexp.weight q8_0
 blk.47.ffn_gate_shexp.weight q8_0
 blk.47.ffn_up_shexp.weight q8_0
+blk.48.attn_k_b.weight q8_0
 blk.48.attn_kv_a_mqa.weight q8_0
 blk.48.attn_output.weight q8_0
 blk.48.attn_q_a.weight q8_0
 blk.48.attn_q_b.weight q8_0
+blk.48.attn_v_b.weight q8_0
 blk.48.ffn_down_shexp.weight q8_0
 blk.48.ffn_gate_shexp.weight q8_0
 blk.48.ffn_up_shexp.weight q8_0
+blk.49.attn_k_b.weight q8_0
 blk.49.attn_kv_a_mqa.weight q8_0
 blk.49.attn_output.weight q8_0
 blk.49.attn_q_a.weight q8_0
 blk.49.attn_q_b.weight q8_0
+blk.49.attn_v_b.weight q8_0
 blk.49.ffn_down_shexp.weight q8_0
 blk.49.ffn_gate_shexp.weight q8_0
 blk.49.ffn_up_shexp.weight q8_0
+blk.50.attn_k_b.weight q8_0
 blk.50.attn_kv_a_mqa.weight q8_0
 blk.50.attn_output.weight q8_0
 blk.50.attn_q_a.weight q8_0
 blk.50.attn_q_b.weight q8_0
+blk.50.attn_v_b.weight q8_0
 blk.50.ffn_down_shexp.weight q8_0
 blk.50.ffn_gate_shexp.weight q8_0
 blk.50.ffn_up_shexp.weight q8_0
+blk.51.attn_k_b.weight q8_0
 blk.51.attn_kv_a_mqa.weight q8_0
 blk.51.attn_output.weight q8_0
 blk.51.attn_q_a.weight q8_0
 blk.51.attn_q_b.weight q8_0
+blk.51.attn_v_b.weight q8_0
 blk.51.ffn_down_shexp.weight q8_0
 blk.51.ffn_gate_shexp.weight q8_0
 blk.51.ffn_up_shexp.weight q8_0
+blk.52.attn_k_b.weight q8_0
 blk.52.attn_kv_a_mqa.weight q8_0
 blk.52.attn_output.weight q8_0
 blk.52.attn_q_a.weight q8_0
 blk.52.attn_q_b.weight q8_0
+blk.52.attn_v_b.weight q8_0
 blk.52.ffn_down_shexp.weight q8_0
 blk.52.ffn_gate_shexp.weight q8_0
 blk.52.ffn_up_shexp.weight q8_0
+blk.53.attn_k_b.weight q8_0
 blk.53.attn_kv_a_mqa.weight q8_0
 blk.53.attn_output.weight q8_0
 blk.53.attn_q_a.weight q8_0
 blk.53.attn_q_b.weight q8_0
+blk.53.attn_v_b.weight q8_0
 blk.53.ffn_down_shexp.weight q8_0
 blk.53.ffn_gate_shexp.weight q8_0
 blk.53.ffn_up_shexp.weight q8_0
+blk.54.attn_k_b.weight q8_0
 blk.54.attn_kv_a_mqa.weight q8_0
 blk.54.attn_output.weight q8_0
 blk.54.attn_q_a.weight q8_0
 blk.54.attn_q_b.weight q8_0
+blk.54.attn_v_b.weight q8_0
 blk.54.ffn_down_shexp.weight q8_0
 blk.54.ffn_gate_shexp.weight q8_0
 blk.54.ffn_up_shexp.weight q8_0
+blk.55.attn_k_b.weight q8_0
 blk.55.attn_kv_a_mqa.weight q8_0
 blk.55.attn_output.weight q8_0
 blk.55.attn_q_a.weight q8_0
 blk.55.attn_q_b.weight q8_0
+blk.55.attn_v_b.weight q8_0
 blk.55.ffn_down_shexp.weight q8_0
 blk.55.ffn_gate_shexp.weight q8_0
 blk.55.ffn_up_shexp.weight q8_0
+blk.56.attn_k_b.weight q8_0
 blk.56.attn_kv_a_mqa.weight q8_0
 blk.56.attn_output.weight q8_0
 blk.56.attn_q_a.weight q8_0
 blk.56.attn_q_b.weight q8_0
+blk.56.attn_v_b.weight q8_0
 blk.56.ffn_down_shexp.weight q8_0
 blk.56.ffn_gate_shexp.weight q8_0
 blk.56.ffn_up_shexp.weight q8_0
+blk.57.attn_k_b.weight q8_0
 blk.57.attn_kv_a_mqa.weight q8_0
 blk.57.attn_output.weight q8_0
 blk.57.attn_q_a.weight q8_0
 blk.57.attn_q_b.weight q8_0
+blk.57.attn_v_b.weight q8_0
 blk.57.ffn_down_shexp.weight q8_0
 blk.57.ffn_gate_shexp.weight q8_0
 blk.57.ffn_up_shexp.weight q8_0
+blk.58.attn_k_b.weight q8_0
 blk.58.attn_kv_a_mqa.weight q8_0
 blk.58.attn_output.weight q8_0
 blk.58.attn_q_a.weight q8_0
 blk.58.attn_q_b.weight q8_0
+blk.58.attn_v_b.weight q8_0
 blk.58.ffn_down_shexp.weight q8_0
 blk.58.ffn_gate_shexp.weight q8_0
 blk.58.ffn_up_shexp.weight q8_0
+blk.59.attn_k_b.weight q8_0
 blk.59.attn_kv_a_mqa.weight q8_0
 blk.59.attn_output.weight q8_0
 blk.59.attn_q_a.weight q8_0
 blk.59.attn_q_b.weight q8_0
+blk.59.attn_v_b.weight q8_0
 blk.59.ffn_down_shexp.weight q8_0
 blk.59.ffn_gate_shexp.weight q8_0
 blk.59.ffn_up_shexp.weight q8_0
+blk.60.attn_k_b.weight q8_0
 blk.60.attn_kv_a_mqa.weight q8_0
 blk.60.attn_output.weight q8_0
 blk.60.attn_q_a.weight q8_0
 blk.60.attn_q_b.weight q8_0
+blk.60.attn_v_b.weight q8_0
 blk.60.ffn_down_shexp.weight q8_0
 blk.60.ffn_gate_shexp.weight q8_0
 blk.60.ffn_up_shexp.weight q8_0

tests/test-backend-ops.cpp

@@ -9996,6 +9996,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     std::default_random_engine rng(0);
 
     // unary ops
+    // TODO: add GGML_UNARY_OP_SOFTPLUS test that is sensitive to using `log1p(exp(x))` vs `log(exp(x) + 1)`
     for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32, GGML_TYPE_BF16}) {
         for (int v : {0, 1}) {
             for (int op = 0; op < GGML_UNARY_OP_COUNT; op++) {
@@ -11079,6 +11080,17 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_mul_mat_id_w4a8(GGML_TYPE_MXFP4, GGML_TYPE_F32, 8, 2, false, 32, 32, 256));
     test_cases.emplace_back(new test_mul_mat_id_w4a4(GGML_TYPE_MXFP4, GGML_TYPE_F32, 8, 2, false, 32, 32, 256));
 
+#if 0
+    // > 2^31-element dst (8.9 GB), about 45 GB peak memory. Too large to be enabled by default.
+    // kernel_mul_mm dst offset of the last batch is past INT32_MAX: rows 0-31 use full tiles, rows 32-39 use edge tiles.
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 1048576, 40, 64, {1, 1}, {53, 1}));
+
+    // kernel_mul_mm_id: the index of tokens 32768-33791 does not fit in 16 bits (about 0.3 GB peak memory).
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F32, GGML_TYPE_F32, 4, 2, false, 64, 33792, 64));
+    // kernel_mul_mm_id: > 2^31-element dst (8.7 GB), about 45 GB peak memory. Tokens 32514+ have a dst offset past INT32_MAX, tokens 32768+ also need a 32-bit token index.
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F32, GGML_TYPE_F32, 16, 8, false, 8256, 32832, 64));
+#endif
+
 #if 0
     // > 4GB A matrix. Too slow to be enabled by default.
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16,  900000,  3, 2592, {1, 1}, {1, 1}));

tests/test-grammar-integration.cpp

@@ -879,6 +879,18 @@ static void test_failure_left_recursion() {
         empty ::= "blah" | )""";
     assert(test_build_grammar_fails(hardest_str));
 
+    const auto chain = [](int n) {
+        std::string s = "root ::= r0\n";
+        for (int i = 0; i < n; i++) {
+            s += "r" + std::to_string(i) + " ::= r" + std::to_string(i + 1) + "\n";
+        }
+        return s + "r" + std::to_string(n) + " ::= \"a\"\n";
+    };
+    llama_grammar * grammar = build_grammar(chain(100));
+    assert(grammar != nullptr);
+    llama_grammar_free_impl(grammar);
+    assert(test_build_grammar_fails(chain(1000)));
+
     fprintf(stderr, "  ✅︎ Passed\n");
 }
 
@@ -1476,6 +1488,25 @@ static void test_json_schema() {
             // R"""({"productId": 1, "productName": "A green door", "price": 12.50, "tags": ["home", "green", "home"]})""",
         }
     );
+
+    // the deepest schema of json_schema_to_grammar must give a valid grammar
+    std::string deep_schema;
+    for (int i = 0; i < 127; i++) {
+        deep_schema += R"({"type": "array", "items": )";
+    }
+    deep_schema += R"({"type": "string"})" + std::string(127, '}');
+    test_schema(
+        "127 nested arrays",
+        deep_schema,
+        // Passing strings
+        {
+            std::string(127, '[') + R"("a")" + std::string(127, ']'),
+        },
+        // Failing strings
+        {
+            std::string(126, '[') + R"("a")" + std::string(126, ']'),
+        }
+    );
 }
 
 int main() {

tests/test-grammar-parser.cpp

@@ -165,6 +165,20 @@ int main()
         root ::= "a"{5000,6000}
     )""");
 
+    // Recursion depth limit exceeded (Issue #29495)
+    {
+        std::string deep_grammar = "root ::= ";
+        int n_depth = 300;
+        for (int i = 0; i < n_depth; ++i) {
+            deep_grammar += "(";
+        }
+        deep_grammar += "\"x\"";
+        for (int i = 0; i < n_depth; ++i) {
+            deep_grammar += ")";
+        }
+        verify_failure(deep_grammar.c_str());
+    }
+
     verify_parsing(R"""(
         root ::= "a"{0,5000}
     )""", {

tests/test-jinja.cpp

@@ -17,6 +17,7 @@
 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);
+static void test_template_error(testing & t, const std::string & name, const std::string & tmpl, const std::string & expect);
 
 static void test_whitespace_control(testing & t);
 static void test_conditionals(testing & t);
@@ -606,6 +607,17 @@ static void test_expressions(testing & t) {
         {{"cond", true}},
         "yes"
     );
+
+    test_template(t, "nested parentheses",
+        "{{ " + std::string(30, '(') + "1" + std::string(30, ')') + " }}",
+        json::object(),
+        "1"
+    );
+
+    test_template_error(t, "too deeply nested parentheses",
+        "{{ " + std::string(1000, '(') + "1" + std::string(1000, ')') + " }}",
+        "template is nested more than 128 levels deep"
+    );
 }
 
 static void test_set_statement(testing & t) {
@@ -1029,6 +1041,11 @@ static void test_filters(testing & t) {
         {{"x", nullptr}},
         "None"
     );
+
+    test_template_error(t, "tojson of a list that contains itself",
+        "{% set arr = [] %}{% set _ = arr.append(arr) %}{{ arr|tojson }}",
+        "value is nested more than 128 levels deep"
+    );
 }
 
 static void test_literals(testing & t) {
@@ -1291,6 +1308,17 @@ static void test_macros(testing & t) {
         json::object(),
         "Dict {\n    root1: 1\n    root2: {\n        nest1: 1\n        nest2: {\n            nest3: 2\n        }\n    }\n}"
     );
+
+    test_template(t, "macro recursion",
+        "{% macro f(n) %}{% if n < 100 %}{{ f(n + 1) }}{% else %}{{ n }}{% endif %}{% endmacro %}{{ f(0) }}",
+        json::object(),
+        "100"
+    );
+
+    test_template_error(t, "endless macro recursion",
+        "{% macro f(n) %}{{ f(n + 1) }}{% endmacro %}{{ f(0) }}",
+        "macro calls are nested more than 128 levels deep"
+    );
 }
 
 static void test_namespace(testing & t) {
@@ -2554,6 +2582,22 @@ static void test_string_parts(testing & t) {
 
 }
 
+static void test_template_error(testing & t, const std::string & name, const std::string & tmpl, const std::string & expect) {
+    t.test(name, [&tmpl, &expect](testing & t) {
+        std::string error;
+        try {
+            jinja::lexer lexer;
+            jinja::program ast = jinja::parse_from_tokens(lexer.tokenize(tmpl));
+            jinja::context ctx(tmpl);
+            jinja::runtime runtime(ctx);
+            runtime.gather_string_parts(runtime.execute(ast));
+        } catch (const std::exception & e) {
+            error = e.what();
+        }
+        t.assert_true("expected error: " + expect + ", got: " + error, error.find(expect) != std::string::npos);
+    });
+}
+
 static void test_template_cpp(testing & t, const std::string & name, const std::string & tmpl, const json & vars, const std::string & expect) {
     t.test(name, [&tmpl, &vars, &expect](testing & t) {
         jinja::lexer lexer;

tests/test-json-schema-to-grammar.cpp

@@ -1513,6 +1513,23 @@ static void test_all(const std::string & title, std::function<void(const TestCas
         ""
     });
 
+    test({
+        FAILURE,
+        "allOf with a $ref cycle",
+        R"""({
+            "allOf": [{"$ref": "#/$defs/a"}],
+            "$defs": {"a": {"$ref": "#/$defs/a"}}
+        })""",
+        ""
+    });
+
+    test({
+        FAILURE,
+        "long property name with additionalProperties",
+        R"""({"type": "object", "properties": {")""" + std::string(1000, 'a') + R"""(": {"type": "string"}}, "additionalProperties": true})""",
+        ""
+    });
+
     // only the property with the bad pattern degrades
     test({
         SUCCESS,

tests/test-llama-archs.cpp

@@ -6,10 +6,7 @@
 #include "ggml-cpp.h"
 #include "llama.h"
 #include "llama-cpp.h"
-
-// TODO: replace with #include "llama-ext.h" in the future
-#include "../src/llama-arch.h"
-#include "../src/llama-model-saver.h"
+#include "../src/llama-ext.h"
 
 #include <cinttypes>
 #include <cmath>
@@ -17,6 +14,7 @@
 #include <cstdio>
 #include <cstring>
 #include <cstdint>
+#include <filesystem>
 #include <random>
 #include <regex>
 #include <stdexcept>
@@ -28,7 +26,7 @@ static bool arch_matches(const std::string & filter, llm_arch arch) {
     if (filter.empty()) {
         return true;
     }
-    return std::regex_search(llm_arch_name(arch), std::regex(filter));
+    return std::regex_search(llama_arch_name(arch), std::regex(filter));
 }
 
 // normalized mean squared error = mse(a, b) / mse(a, 0)
@@ -164,9 +162,50 @@ static bool host_experts_test(const llm_arch arch) {
     }
 }
 
+// writes arch-resolved GGUF keys into the fixture context, the test-side equivalent of llama_model_saver
+struct kv_writer {
+    struct gguf_context * gguf_ctx;
+    const llm_arch        arch;
+
+    kv_writer(enum llm_arch arch, struct gguf_context * gguf_ctx) : gguf_ctx(gguf_ctx), arch(arch) {}
+
+    void add_kv(enum llm_kv key, uint32_t     value) { gguf_set_val_u32 (gguf_ctx, llama_arch_kv_name(arch, key).c_str(), value); }
+    void add_kv(enum llm_kv key, int32_t      value) { gguf_set_val_i32 (gguf_ctx, llama_arch_kv_name(arch, key).c_str(), value); }
+    void add_kv(enum llm_kv key, uint64_t     value) { gguf_set_val_u64 (gguf_ctx, llama_arch_kv_name(arch, key).c_str(), value); }
+    void add_kv(enum llm_kv key, float        value) { gguf_set_val_f32 (gguf_ctx, llama_arch_kv_name(arch, key).c_str(), value); }
+    void add_kv(enum llm_kv key, bool         value) { gguf_set_val_bool(gguf_ctx, llama_arch_kv_name(arch, key).c_str(), value); }
+    void add_kv(enum llm_kv key, const char * value) { gguf_set_val_str (gguf_ctx, llama_arch_kv_name(arch, key).c_str(), value); }
+
+    template <typename T>
+    void add_kv(enum llm_kv key, const std::vector<T> & value) {
+        if (value.empty()) {
+            return;
+        }
+        if (std::is_same<T, uint32_t>::value) {
+            gguf_set_arr_data(gguf_ctx, llama_arch_kv_name(arch, key).c_str(), GGUF_TYPE_UINT32, value.data(), value.size());
+        } else if (std::is_same<T, int32_t>::value) {
+            gguf_set_arr_data(gguf_ctx, llama_arch_kv_name(arch, key).c_str(), GGUF_TYPE_INT32, value.data(), value.size());
+        } else if (std::is_same<T, uint64_t>::value) {
+            gguf_set_arr_data(gguf_ctx, llama_arch_kv_name(arch, key).c_str(), GGUF_TYPE_UINT64, value.data(), value.size());
+        } else if (std::is_same<T, float>::value) {
+            gguf_set_arr_data(gguf_ctx, llama_arch_kv_name(arch, key).c_str(), GGUF_TYPE_FLOAT32, value.data(), value.size());
+        } else {
+            GGML_ABORT("fatal error");
+        }
+    }
+
+    void add_kv(enum llm_kv key, const std::vector<std::string> & value) {
+        std::vector<const char *> tmp(value.size());
+        for (size_t i = 0; i < value.size(); ++i) {
+            tmp[i] = value[i].c_str();
+        }
+        gguf_set_arr_str(gguf_ctx, llama_arch_kv_name(arch, key).c_str(), tmp.data(), tmp.size());
+    }
+};
+
 static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
     gguf_context_ptr ret(gguf_init_empty());
-    llama_model_saver ms(arch, ret.get());
+    kv_writer ms(arch, ret.get());
     const uint32_t n_ctx = 256;
 
     uint32_t n_vocab = 128;
@@ -243,7 +282,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
     }
     const uint32_t n_embd_head = n_embd / n_head;
 
-    ms.add_kv(LLM_KV_GENERAL_ARCHITECTURE,      llm_arch_name(arch));
+    ms.add_kv(LLM_KV_GENERAL_ARCHITECTURE,      llama_arch_name(arch));
     ms.add_kv(LLM_KV_VOCAB_SIZE,                n_vocab);
     ms.add_kv(LLM_KV_CONTEXT_LENGTH,            n_ctx);
     ms.add_kv(LLM_KV_EMBEDDING_LENGTH,          n_embd);
@@ -638,6 +677,11 @@ static bool silent_model_load_progress(float /*progress*/, void * /*user_data*/)
     return true;
 }
 
+static std::string get_tmp_gguf_path() {
+    static size_t id = 0;
+    return (std::filesystem::temp_directory_path() / ("test-llama-archs-" + std::to_string(++id) + ".gguf")).string();
+}
+
 static std::pair<llama_model_ptr, llama_context_ptr> get_model_and_ctx(
         struct gguf_context * gguf_ctx, FILE * file, const size_t seed, const float stdev,
         const std::vector<ggml_backend_dev_t> & devs,
@@ -958,7 +1002,7 @@ static int save_models(const std::string & arch_filter, const size_t seed, const
         }
     }, &ud);
 
-    for (const llm_arch & arch : llm_arch_all()) {
+    for (const llm_arch & arch : llama_arch_all()) {
         if (arch == LLM_ARCH_UNKNOWN) {
             continue;
         }
@@ -978,14 +1022,14 @@ static int save_models(const std::string & arch_filter, const size_t seed, const
             if (!moe && moe_mandatory(arch)) {
                 continue;
             }
-            if (!llama_model_saver_supports_arch(arch) || !arch_supported(arch)) {
-                LOG_INF("%s: %s model (%s) is unsupported, skipping\n", __func__, llm_arch_name(arch), moe ? "MoE" : "dense");
+            if (!llama_arch_supports_save(arch) || !arch_supported(arch)) {
+                LOG_INF("%s: %s model (%s) is unsupported, skipping\n", __func__, llama_arch_name(arch), moe ? "MoE" : "dense");
                 continue;
             }
             gguf_context_ptr gguf_ctx = get_gguf_ctx(arch, moe);
             auto model_and_ctx = get_model_and_ctx(gguf_ctx.get(), nullptr, seed, stdev, {});
-            const std::string path = dir + "/" + llm_arch_name(arch) + (moe ? "-moe.gguf" : "-dense.gguf");
-            LOG_INF("%s: Saving %s model (%s) to %s...\n", __func__, llm_arch_name(arch), moe ? "MoE" : "dense", path.c_str());
+            const std::string path = dir + "/" + llama_arch_name(arch) + (moe ? "-moe.gguf" : "-dense.gguf");
+            LOG_INF("%s: Saving %s model (%s) to %s...\n", __func__, llama_arch_name(arch), moe ? "MoE" : "dense", path.c_str());
             llama_model_save_to_file(model_and_ctx.first.get(), path.c_str());
         }
     }
@@ -1094,8 +1138,8 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con
     }
 
     size_t max_arch_name_length = 0;
-    for (const llm_arch & arch : llm_arch_all()) {
-        max_arch_name_length = std::max(max_arch_name_length, strlen(llm_arch_name(arch)));
+    for (const llm_arch & arch : llama_arch_all()) {
+        max_arch_name_length = std::max(max_arch_name_length, strlen(llama_arch_name(arch)));
     }
 
     const std::string template_header  = std::string("|%" + std::to_string(max_arch_name_length) + "s|%") + std::to_string(max_device_label_length) + "s|%6s|%15s|%9s|%15s|\n";
@@ -1116,7 +1160,7 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con
         LOG("-");
     }
     LOG("|------|---------------|---------|---------------|\n");
-    for (const llm_arch & arch : llm_arch_all()) {
+    for (const llm_arch & arch : llama_arch_all()) {
         if (arch == LLM_ARCH_UNKNOWN) {
             continue;
         }
@@ -1152,7 +1196,7 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con
                 const llama_model_tensor_buft_override * overrides = dc.host_experts ? host_experts_overrides : nullptr;
 
                 // print test config first; should anything fail during model loading or inference, at least we know which test case caused it
-                LOG(template_row_cfg.c_str(), llm_arch_name(arch), dc.label.c_str(), config_name.c_str());
+                LOG(template_row_cfg.c_str(), llama_arch_name(arch), dc.label.c_str(), config_name.c_str());
                 fflush(stdout);
 
                 std::pair<llama_model_ptr, llama_context_ptr> model_and_ctx_dev;
@@ -1171,7 +1215,7 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con
                         model_and_ctx_cpu = get_model_and_ctx(gguf_ctx.get(), nullptr, seed, stdev, {}, LLAMA_SPLIT_MODE_LAYER, encode);
                         logits_cpu = get_logits(model_and_ctx_cpu.first.get(), model_and_ctx_cpu.second.get(), tokens, encode);
                     }
-                    if (dc.split_mode != LLAMA_SPLIT_MODE_TENSOR || llm_arch_supports_sm_tensor(arch)) {
+                    if (dc.split_mode != LLAMA_SPLIT_MODE_TENSOR || llama_arch_supports_sm_tensor(arch)) {
                         test_executed = true;
                         model_and_ctx_dev = get_model_and_ctx(gguf_ctx.get(), nullptr, seed, stdev, dc.devs, dc.split_mode, encode, overrides, dc.moe_cache_size);
                         logits_dev = get_logits(model_and_ctx_dev.first.get(), model_and_ctx_dev.second.get(), tokens, encode);
@@ -1194,7 +1238,7 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con
                             }
                             std::vector<float> logits_mixed;
                             std::vector<float> logits_chunks;
-                            if (llm_arch_supports_mixed_batch(arch)) {
+                            if (llama_arch_supports_mixed_batch(arch)) {
                                 if (get_logits_mixed(model_and_ctx_dev.first.get(), lctx_dev, tokens, embd_mixed, false, logits_chunks) != 0 ||
                                     get_logits_mixed(model_and_ctx_dev.first.get(), lctx_dev, tokens, embd_mixed, true,  logits_mixed)  != 0) {
                                     throw std::runtime_error("failed to decode mixed batch");
@@ -1227,30 +1271,33 @@ static int test_backends(const std::string & arch_filter, const size_t seed, con
                         }
                     }
 
-                    FILE * file = tmpfile(); // Can be null on Windows without administrator privileges.
-                    // FIXME: when adding a tensor to a gguf_context a copy is made, this changes the pointer which the meta backend
-                    //     in turn uses to map the tensors to their simple equivalents - this is fundamentally incompatible
-                    if (file != nullptr && llama_model_saver_supports_arch(arch) && dc.split_mode != LLAMA_SPLIT_MODE_TENSOR) {
+                    if (llama_arch_supports_save(arch) && dc.split_mode != LLAMA_SPLIT_MODE_TENSOR) {
                         test_executed = true;
                         GGML_ASSERT(model_and_ctx_dev.first && model_and_ctx_dev.second);
-                        llama_model_saver ms = llama_model_saver(model_and_ctx_dev.first.get());
-                        ms.add_kv_from_model();
-                        ms.add_tensors_from_model();
-                        ms.save(file);
-                        rewind(file);
+                        // FIXME: when adding a tensor to a gguf_context a copy is made, this changes the pointer which the meta backend
+                        //     in turn uses to map the tensors to their simple equivalents - this is fundamentally incompatible
+                        const std::string path = get_tmp_gguf_path();
+                        llama_model_save_to_file(model_and_ctx_dev.first.get(), path.c_str());
+                        FILE * file = fopen(path.c_str(), "rb");
+                        GGML_ASSERT(file != nullptr);
 
-                        auto model_and_ctx_roundtrip = get_model_and_ctx(nullptr, file, seed, stdev, dc.devs, dc.split_mode, encode, overrides, dc.moe_cache_size);
-                        const std::vector<float> logits_roundtrip = get_logits(
-                            model_and_ctx_roundtrip.first.get(), model_and_ctx_roundtrip.second.get(), tokens, encode);
-                        status_roundtrip = "\033[1;32mOK\033[0m";
-                        GGML_ASSERT(logits_roundtrip.size() == logits_dev.size());
-                        for (size_t i = 0; i < logits_roundtrip.size(); i++) {
-                            if (logits_roundtrip[i] != logits_dev[i]) {
-                                test_ok = false;
-                                status_roundtrip = "\033[1;31mFAIL\033[0m";
-                                break;
+                        {
+                            auto model_and_ctx_roundtrip = get_model_and_ctx(nullptr, file, seed, stdev, dc.devs, dc.split_mode, encode, overrides, dc.moe_cache_size);
+                            const std::vector<float> logits_roundtrip = get_logits(
+                                model_and_ctx_roundtrip.first.get(), model_and_ctx_roundtrip.second.get(), tokens, encode);
+                            status_roundtrip = "\033[1;32mOK\033[0m";
+                            GGML_ASSERT(logits_roundtrip.size() == logits_dev.size());
+                            for (size_t i = 0; i < logits_roundtrip.size(); i++) {
+                                if (logits_roundtrip[i] != logits_dev[i]) {
+                                    test_ok = false;
+                                    status_roundtrip = "\033[1;31mFAIL\033[0m";
+                                    break;
+                                }
                             }
                         }
+
+                        fclose(file);
+                        std::filesystem::remove(path);
                     }
                 }
 
@@ -1303,7 +1350,7 @@ int main(int argc, char ** argv) {
         } else if (strcmp(argv[i], "-a") == 0 || strcmp(argv[i], "--arch") == 0) {
             if (i + 1 < argc) {
                 const std::string arch_name = argv[++i];
-                if (llm_arch_from_string(arch_name) != LLM_ARCH_UNKNOWN) {
+                if (llama_arch_from_string(arch_name) != LLM_ARCH_UNKNOWN) {
                     // exact architecture name
                     arch_filter = "^" + arch_name + "$";
                 } else {

tests/test-llama-context.cpp

@@ -220,7 +220,7 @@ int main(int argc, char ** argv) {
     bool list_cases        = false;
     bool dump_logs_on_fail = true;
 
-    int parallel_workers = 1; // 0 = auto, same flag as in test-backend-ops
+    int parallel_workers = 0; // 0 = auto, same flag as in test-backend-ops
 
     auto parse_parallel_workers = [&](const char * flag, const char * value) {
         try {
@@ -361,12 +361,9 @@ int main(int argc, char ** argv) {
 
     llama_backend_init();
 
-    // the threads a context runs on, resolved by the common parser - in table mode every worker
-    //   builds a context of its own, so the run uses n_pool times this many
-    const int32_t n_threads = params.cpuparams.n_threads;
-
     // run the full suite for a single model
     if (models_dir.empty()) {
+        const int32_t n_threads = params.cpuparams.n_threads;
         LOG("threads: %d\n", n_threads);
 
         const model_result res = run_model(params.model.path, params, cases, logs_dir,
@@ -404,14 +401,14 @@ int main(int argc, char ** argv) {
         return 1;
     }
 
-    // the dummy models are tiny, but the same flag is meant to run sweeps of real checkpoints,
-    //   where every worker holds the weights it loaded - so do not scale with the cores
-    const unsigned n_hw   = std::thread::hardware_concurrency();
-    const size_t   n_auto = n_hw > 0 ? std::min<size_t>(n_hw, 4) : 1;
+    const unsigned n_hw_2 = std::thread::hardware_concurrency() / 2;
+    const size_t   n_auto = std::max<size_t>(1, std::min<size_t>(n_hw_2, 4));
     const size_t   n_pool = parallel_workers > 0
         ? std::min<size_t>((size_t) parallel_workers, models.size())
         : n_auto;
 
+    const int32_t n_threads = std::max<int32_t>(1, n_hw_2 / n_pool);
+
     // a column is as wide as its case name, but never narrower than the widest verdict
     const size_t verdict_width = std::max({ test_status_width(test_status::PASS),
                                             test_status_width(test_status::FAIL),

tests/test-server-component.cpp

@@ -2,6 +2,7 @@
 
 #include "../tools/server/server-common.h"
 #include "../tools/server/server-decision.h"
+#include "../tools/server/server-schema.h"
 
 #include <cstdlib>
 #include <regex>
@@ -490,6 +491,68 @@ static void test_decision_depth(testing & t) {
     });
 }
 
+// DRY sequence breakers
+//
+
+static void test_dry_sequence_breakers(testing & t) {
+    common_params params_base;
+
+    const auto eval = [&](const json & data) {
+        return server_schema::eval_llama_cmpl_schema(nullptr, params_base, {}, data);
+    };
+
+    {
+        json data;
+        data["dry_sequence_breakers"] = json::array({"ab", "ab", "cd", "ab"});
+        const auto params = eval(data);
+        const std::vector<std::string> expected = {"ab", "cd"};
+        t.assert_true("duplicates removed, order kept", params.sampling.dry_sequence_breakers == expected);
+    }
+
+    {
+        std::vector<std::string> many;
+        for (int i = 0; i <= LLAMA_DRY_MAX_SEQ_BREAKERS; ++i) {
+            many.push_back("b" + std::to_string(i));
+        }
+        json data;
+        data["dry_sequence_breakers"] = many;
+
+        bool thrown = false;
+        try {
+            eval(data);
+        } catch (const std::exception &) {
+            thrown = true;
+        }
+        t.assert_true("more than the cap is rejected", thrown);
+    }
+}
+
+//
+// glob_match
+//
+
+static void test_glob_match(testing & t) {
+    t.test("patterns", [](testing & t) {
+        t.assert_true("literal", glob_match("src/main.cpp", "src/main.cpp"));
+        t.assert_true("star", glob_match("src/*.cpp", "src/main.cpp"));
+        t.assert_true("star stops at /", !glob_match("*.cpp", "src/main.cpp"));
+        t.assert_true("double star", glob_match("**/*.cpp", "a/b/c/main.cpp"));
+        t.assert_true("double star at end", glob_match("src/**", "src/a/b"));
+        t.assert_true("question mark", glob_match("ma?n.cpp", "main.cpp"));
+        t.assert_true("question mark not /", !glob_match("a?b", "a/b"));
+        t.assert_true("class", glob_match("[a-c]x", "bx"));
+        t.assert_true("negated class", !glob_match("[!a-c]x", "bx"));
+        t.assert_true("unclosed class", glob_match("[ab", "[ab"));
+        t.assert_true("mismatch", !glob_match("main.cpp", "main.c"));
+    });
+    t.test("long input", [](testing & t) {
+        const std::string path(100000, 'a');
+        t.assert_true("long literal", glob_match(path, path));
+        t.assert_true("long path", glob_match("**", path));
+        t.assert_true("too many wildcards", !glob_match(std::string(1000, '*'), "a"));
+    });
+}
+
 int main(int argc, char ** argv) {
     testing t;
 
@@ -511,6 +574,8 @@ int main(int argc, char ** argv) {
     t.test("decision_escape", test_decision_escape);
     t.test("rerank_template", test_rerank_template);
     t.test("decision_depth",  test_decision_depth);
+    t.test("glob_match",      test_glob_match);
+    t.test("dry_sequence_breakers", test_dry_sequence_breakers);
 
     return t.summary();
 }

tools/server/README.md

@@ -563,7 +563,7 @@ These words will not be included in the completion, so make sure to add them to
 
 `dry_penalty_last_n`: How many tokens to scan for repetitions. Default: `64`, where `0` is disabled.
 
-`dry_sequence_breakers`: Specify an array of sequence breakers for DRY sampling. Only a JSON array of strings is accepted. Default: `['\n', ':', '"', '*']`
+`dry_sequence_breakers`: Specify an array of sequence breakers for DRY sampling. Only a JSON array of strings is accepted. At most 64 distinct breakers, duplicates are ignored. Default: `['\n', ':', '"', '*']`
 
 `xtc_probability`: Set the chance for token removal via XTC sampler. Default: `0.0`, which is disabled.
 

tools/server/server-schema.cpp

@@ -2,6 +2,8 @@
 
 #include "json-schema-to-grammar.h"
 
+#include <unordered_set>
+
 namespace server_schema {
 
 //
@@ -246,6 +248,23 @@ std::vector<std::unique_ptr<field>> make_llama_cmpl_schema(const common_params &
             if (ctx.params.sampling.dry_sequence_breakers.empty()) {
                 throw std::runtime_error("Error: dry_sequence_breakers must be a non-empty array of strings");
             }
+
+            // each distinct breaker costs a full vocabulary scan at sampler init, so dedupe before counting
+            std::unordered_set<std::string> seen;
+            std::vector<std::string> unique_breakers;
+
+            for (const auto & breaker : ctx.params.sampling.dry_sequence_breakers) {
+                if (seen.insert(breaker).second) {
+                    unique_breakers.push_back(breaker);
+                }
+            }
+
+            if (unique_breakers.size() > LLAMA_DRY_MAX_SEQ_BREAKERS) {
+                throw std::runtime_error(string_format(
+                    "Error: dry_sequence_breakers must contain at most %d distinct entries", LLAMA_DRY_MAX_SEQ_BREAKERS));
+            }
+
+            ctx.params.sampling.dry_sequence_breakers = std::move(unique_breakers);
         }));
 
     // handle both "json_schema" and "grammar"

tools/server/tests/utils.py

@@ -29,7 +29,7 @@ from re import RegexFlag
 import wget
 
 
-DEFAULT_HTTP_TIMEOUT = 60
+DEFAULT_HTTP_TIMEOUT = 600
 
 # per-request timeout, a hung server fails the test instead of stalling the CI for hours
 DEFAULT_REQUEST_TIMEOUT = 600

tools/ui/src/app.css

@@ -35,6 +35,8 @@
 
 :root {
 	--radius: 0.625rem;
+	/* expanded desktop sidebar width, overridden once the user drags its edge */
+	--sidebar-width: 18rem;
 	--background: oklch(1 0 0);
 	--foreground: oklch(0.145 0 0);
 	--card: oklch(1 0 0);
@@ -172,6 +174,23 @@
 		scrollbar-gutter: stable;
 	}
 
+	/* no elastic bounce past the ends of the page: it drags the sticky and
+	   fixed bars (tab bar, chat form) along and they jump back on release */
+	html,
+	body {
+		overscroll-behavior: none;
+	}
+
+	/* same for dialogs, whose panels and lists scroll on their own with sticky
+	   headers and footers. Scroll containers only: Firefox treats an
+	   overflow-hidden box (truncated text) as one, and with `none` it would
+	   keep the wheel from reaching the list around it */
+	:is([role='dialog'], [role='alertdialog']),
+	:is([role='dialog'], [role='alertdialog'])
+		:is(.overflow-auto, .overflow-y-auto, [data-slot='scroll-area-viewport']) {
+		overscroll-behavior: none;
+	}
+
 	/* Global scrollbar styling - visible only on hover */
 	* {
 		scrollbar-width: thin;
@@ -221,6 +240,17 @@
 		scrollbar-width: none;
 	}
 
+	/* Fade at the top edge of the chat, under the tab bar or a sticky header */
+	.chat-top-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%
+		);
+	}
+
 	.shimmer-text {
 		background: linear-gradient(
 			90deg,

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

@@ -96,17 +96,37 @@
 			'.chat-message-user'
 		) as HTMLElement | null;
 
-		if (!userMessageEl) {
+		if (!previousChatMessage || !userMessageEl) {
 			lastUserMessageHeight = 0;
 
 			return;
 		}
 
+		// a first exchange opened only by synthetic rows (e.g. the working
+		// directory notice) must fit the screen as well, so they count as part
+		// of the user turn; real earlier messages still scroll away above it
+		const leadingRows: Element[] = [];
+
+		for (
+			let row = previousChatMessage.previousElementSibling;
+			row;
+			row = row.previousElementSibling
+		) {
+			leadingRows.unshift(row);
+		}
+
+		const leadingRowsAreSynthetic =
+			leadingRows.length > 0 &&
+			leadingRows.every((row) => row.classList.contains('chat-message--synthetic'));
 		const updateHeight = () => {
 			const rect = userMessageEl.getBoundingClientRect();
 			const marginTop = Math.round(parseFloat(getComputedStyle(userMessageEl).marginTop));
+			const leadingHeight = leadingRowsAreSynthetic
+				? previousChatMessage.getBoundingClientRect().top -
+					leadingRows[0].getBoundingClientRect().top
+				: 0;
 
-			lastUserMessageHeight = Math.round(rect.height + marginTop);
+			lastUserMessageHeight = Math.round(rect.height + marginTop + leadingHeight);
 		};
 
 		updateHeight();
@@ -115,6 +135,10 @@
 
 		resizeObserver.observe(userMessageEl);
 
+		if (leadingRowsAreSynthetic) {
+			for (const row of leadingRows) resizeObserver.observe(row);
+		}
+
 		return () => {
 			resizeObserver.disconnect();
 		};

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

@@ -1,6 +1,7 @@
 <script lang="ts">
 	import LazyChatMessage from './LazyChatMessage.svelte';
 	import { ChatMessageUserPending } from '$lib/components/app';
+	import { LANDING_EAGER_ROWS } from '$lib/constants';
 	import { MessageRole } from '$lib/enums';
 	import { agenticStore, chatStore, conversationsStore, settingsStore } from '$lib/stores';
 	import type { ChatMessageActions } from '$lib/types';
@@ -237,14 +238,15 @@
 	});
 </script>
 
-<!-- Re-created per conversation, so the CSS fade-in below plays on every
-     navigation into a chat route. -->
+<!-- Re-created per conversation, so each one starts from fresh rows that
+     mount eagerly at its bottom. -->
 {#key conversationsStore.activeConversation?.id ?? 'new'}
 	<div class="chat-messages">
-		{#each displayMessages as { isLastAssistantMessage, isLastUserMessage, message, nextAssistantMessage, siblingInfo, toolMessages } (message.id)}
+		{#each displayMessages as { isLastAssistantMessage, isLastUserMessage, message, nextAssistantMessage, siblingInfo, toolMessages }, index (message.id)}
 			<LazyChatMessage
 				{chatActions}
 				class="mx-auto mt-12 w-full max-w-3xl"
+				eager={index >= displayMessages.length - LANDING_EAGER_ROWS}
 				{isLastAssistantMessage}
 				{isLastUserMessage}
 				{message}

tools/ui/src/lib/components/app/chat/ChatMessages/LazyChatMessage.svelte

@@ -6,6 +6,8 @@
 	interface Props {
 		chatActions: ChatMessageActions;
 		class?: string;
+		/** Mount right away instead of waiting for the row to near the viewport. */
+		eager?: boolean;
 		isLastAssistantMessage?: boolean;
 		isLastUserMessage?: boolean;
 		message: DatabaseMessage;
@@ -17,6 +19,7 @@
 	let {
 		chatActions,
 		class: className = '',
+		eager = false,
 		isLastAssistantMessage = false,
 		isLastUserMessage = false,
 		message,
@@ -34,8 +37,17 @@
 	let mounted = $state(false);
 	let wrapperEl: HTMLDivElement | undefined = $state();
 
+	// eager rows render in the same pass as the list, so a conversation
+	// landing at its bottom never paints empty placeholders there; latch it,
+	// since a row stops being eager once newer messages push it up
+	let shouldMount = $derived(mounted || eager);
+
 	$effect(() => {
-		if (mounted || !wrapperEl) return;
+		if (eager) mounted = true;
+	});
+
+	$effect(() => {
+		if (shouldMount || !wrapperEl) return;
 
 		const observer = new IntersectionObserver(
 			(entries) => {
@@ -65,10 +77,11 @@
 
 <div
 	bind:this={wrapperEl}
+	class:chat-message--eager={eager}
 	class:chat-message--synthetic={Boolean(message.isSynthetic)}
 	class="chat-message"
 >
-	{#if mounted}
+	{#if shouldMount}
 		<ChatMessage
 			{chatActions}
 			class={className}
@@ -102,4 +115,13 @@
 	.chat-message--synthetic {
 		--chat-message-intrinsic-size: 40px;
 	}
+
+	/*
+	 * A freshly created row is skipped for a frame until its viewport
+	 * proximity is determined, so the rows a conversation lands on would paint
+	 * as empty placeholders. They sit at the bottom anyway, so lay them out.
+	 */
+	.chat-message--eager {
+		content-visibility: visible;
+	}
 </style>

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

@@ -11,7 +11,16 @@
 		ChatScreenStreamResumeStatus,
 		ServerLoadingSplash
 	} from '$lib/components/app';
-	import { LANDING_SETTLE_MAX_MS, LANDING_STABLE_FRAMES, ROUTES } from '$lib/constants';
+	import {
+		CHAT_FORM_MOVE_IN_EASING,
+		CHAT_FORM_MOVE_IN_MS,
+		CHAT_FORM_MOVE_OUT_EASING,
+		CHAT_FORM_MOVE_OUT_MS,
+		CHAT_FORM_MOVE_SHIFT_PX,
+		CHAT_FORM_MOVE_TO_BOTTOM_IN_MS,
+		CHAT_FORM_MOVE_TO_BOTTOM_OUT_MS,
+		ROUTES
+	} from '$lib/constants';
 	import { createAutoScrollController } from '$lib/hooks/use-auto-scroll.svelte';
 	import { useChatScreenActiveModel } from '$lib/hooks/use-chat-screen-active-model.svelte';
 	import { useChatScreenDragAndDrop } from '$lib/hooks/use-chat-screen-drag-and-drop.svelte';
@@ -26,7 +35,7 @@
 		settingsStore
 	} from '$lib/stores';
 	import { parseFilesToMessageExtras } from '$lib/utils/browser-only';
-	import { onDestroy, onMount, tick } from 'svelte';
+	import { onDestroy, onMount, tick, untrack } from 'svelte';
 
 	let { showCenteredEmpty = false } = $props();
 
@@ -43,6 +52,10 @@
 	let isEmpty = $derived(
 		showCenteredEmpty && conversationsStore.activeMessages.length === 0 && !chatStore.isLoading
 	);
+	// where the form is drawn: follows `isEmpty` once the form has faded out
+	let isFormCentered = $state(untrack(() => isEmpty));
+	let isFormMoving = false;
+	let formWrapper: HTMLElement | undefined = $state();
 	let activeErrorDialog = $derived(chatStore.errorDialogState);
 	let isServerLoading = $derived(serverStore.loading);
 	let hasPropsError = $derived(!!serverStore.error);
@@ -133,11 +146,11 @@
 	let lastScrolledConversationId: string | null = null;
 
 	// Lands at the bottom of a conversation the first time its messages
-	// render, whether the route comes from another conversation or from a
-	// non-conversation route. The page keeps growing after the first pin
-	// without DOM mutations (content-visibility size realizations, syntax
-	// highlight passes), so the instant pin repeats every frame until the
-	// height settles, bailing out on user scroll or conversation change.
+	// render, whether the route comes from another conversation, from a
+	// non-conversation route or from a page refresh. Runs before the first
+	// paint of the new rows, and the landing pin keeps the bottom in place
+	// while the rows settle. Independent of the auto-scroll setting, which
+	// only governs following a streaming response.
 	async function handleMessagesReady(messageCount: number) {
 		if (messageCount === 0) return;
 
@@ -147,38 +160,60 @@
 
 		lastScrolledConversationId = id;
 		await tick();
-		autoScroll.scrollToBottom();
 
-		const container = scroll.chatScrollContainer;
+		if (conversationsStore.activeConversation?.id !== id) return;
 
-		if (!container) return;
-
-		const started = performance.now();
-
-		let stableFrames = 0;
-		let lastHeight = container.scrollHeight;
-
-		const settle = () => {
-			if (autoScroll.userScrolledUp) return;
-
-			if (conversationsStore.activeConversation?.id !== id) return;
-
-			autoScroll.scrollToBottom();
-			const height = container.scrollHeight;
-
-			stableFrames = height === lastHeight ? stableFrames + 1 : 0;
-			lastHeight = height;
-
-			if (stableFrames >= LANDING_STABLE_FRAMES) return;
-
-			if (performance.now() - started > LANDING_SETTLE_MAX_MS) return;
-
-			requestAnimationFrame(settle);
-		};
-
-		requestAnimationFrame(settle);
+		scroll.landAtBottom();
 	}
 
+	/**
+	 * Moves the form between its centered and bottom positions: it fades out
+	 * where it is, drifting toward where it goes, and fades
+	 * back in from the other side once it is there. Reduced motion keeps only
+	 * the fade. Re-reads `isEmpty` after each move, so a flip back mid-move
+	 * ends where the screen is.
+	 */
+	async function moveForm() {
+		if (isFormMoving) return;
+
+		isFormMoving = true;
+
+		while (formWrapper && isFormCentered !== isEmpty) {
+			const toCenter = isEmpty;
+			const reduceMotion = matchMedia('(prefers-reduced-motion: reduce)').matches;
+			const shift = reduceMotion ? 0 : (toCenter ? -1 : 1) * CHAT_FORM_MOVE_SHIFT_PX;
+			const shown = { opacity: 1, transform: 'translateY(0)' };
+			// `transform` composes with the `translate` that centers the form
+			const leave = formWrapper.animate(
+				[shown, { opacity: 0, transform: `translateY(${shift}px)` }],
+				{
+					duration: toCenter ? CHAT_FORM_MOVE_OUT_MS : CHAT_FORM_MOVE_TO_BOTTOM_OUT_MS,
+					easing: CHAT_FORM_MOVE_OUT_EASING,
+					fill: 'forwards'
+				}
+			);
+
+			await leave.finished.catch(() => {});
+
+			isFormCentered = toCenter;
+			await tick();
+
+			// cancelled in the same frame the entry starts, so it never flashes
+			leave.cancel();
+			formWrapper?.animate([{ opacity: 0, transform: `translateY(${-shift}px)` }, shown], {
+				duration: toCenter ? CHAT_FORM_MOVE_IN_MS : CHAT_FORM_MOVE_TO_BOTTOM_IN_MS,
+				easing: CHAT_FORM_MOVE_IN_EASING
+			});
+		}
+
+		isFormCentered = isEmpty;
+		isFormMoving = false;
+	}
+
+	$effect(() => {
+		if (isEmpty !== isFormCentered) untrack(() => void moveForm());
+	});
+
 	function handleSendLikeScroll() {
 		if (!deviceStore.isMobile) {
 			autoScroll.enable();
@@ -189,9 +224,12 @@
 
 			if (!container) return;
 
-			const lastUserBubble = container.querySelector(
-				'.chat-message:nth-last-child(2) .chat-message-user .chat-message-user-bubble'
-			) as HTMLElement | null;
+			// the bubble of the message just sent: the last user bubble, whether or
+			// not the answer row below it has rendered yet
+			const userBubbles = container.querySelectorAll<HTMLElement>(
+				'.chat-message-user .chat-message-user-bubble'
+			);
+			const lastUserBubble = userBubbles[userBubbles.length - 1] ?? null;
 
 			if (deviceStore.isMobile) {
 				// Keep the last user message bubble just above the input on mobile
@@ -295,7 +333,7 @@
 {:else}
 	<div
 		style:--chat-form-bottom-position={chatFormBottomPosition}
-		class="chat-screen flex grow flex-col min-h-[calc(100dvh-1rem)] md:min-h-[calc(100dvh-1rem-var(--chat-tabs-offset,0px))] px-4 md:py-0 pt-12 pb-48 md:pb-4"
+		class="chat-screen flex grow flex-col min-h-[calc(100dvh-1rem)] md:min-h-[calc(100dvh-var(--chat-tabs-offset,0px))] px-4 md:py-0 pt-12 pb-48 md:pb-4"
 		ondragenter={dragAndDrop.dragHandlers.dragenter}
 		ondragleave={dragAndDrop.dragHandlers.dragleave}
 		ondragover={dragAndDrop.dragHandlers.dragover}
@@ -313,23 +351,26 @@
 		{/if}
 
 		<div
-			style:padding-top={!isEmpty ? 'var(--chat-form-padding-top)' : undefined}
+			bind:this={formWrapper}
+			style:padding-top="var(--chat-form-padding-top)"
 			class={[
-				// animate the centered->bottomed move with transform, not bottom:
-				// layout-property transitions need the main thread every frame and
-				// stutter while a long conversation loads; transform transitions
-				// run on the compositor and stay smooth
-				'pointer-events-none md:sticky fixed  mt-auto transition-transform duration-200',
+				// moves between the centered (empty chat) and bottomed positions with
+				// a fade, see moveForm. Its height in the page is the same in both
+				// (the padding stays, the desktop greeting floats above the form),
+				// so the messages never shift while it moves
+				'pointer-events-none md:sticky fixed mt-auto',
 				deviceStore.isStandalone
 					? 'bottom-6 right-4 left-4'
 					: deviceStore.isIOSSafari
 						? 'bottom-1 left-2 right-2'
 						: 'bottom-2 right-2 left-2',
 				'md:bottom-4',
-				isEmpty ? 'md:translate-y-[calc(-50dvh+8rem)] 2xl:translate-y-[calc(-50dvh+5rem)]' : ''
+				isFormCentered
+					? 'md:translate-y-[calc(-50dvh+7rem)] 2xl:translate-y-[calc(-50dvh+4rem)]'
+					: ''
 			]}
 		>
-			<ChatScreenGreeting {isEmpty} />
+			<ChatScreenGreeting isEmpty={isFormCentered} />
 
 			<ChatScreenServerError />
 

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { CHAT_GREETING_ENTER_DELAY_MS } from '$lib/constants';
 	import { serverStore } from '$lib/stores';
 
 	interface Props {
@@ -6,18 +7,47 @@
 	}
 
 	let { isEmpty = false }: Props = $props();
+
+	let hasEntered = $state(false);
+
+	// Starts the entry only once the greeting has been painted hidden, so it
+	// plays on a refresh or a first visit too, when the first frames are busy
+	// loading the app, not just after the form moves.
+	$effect(() => {
+		if (!isEmpty) {
+			hasEntered = false;
+
+			return;
+		}
+
+		let frame = requestAnimationFrame(() => {
+			frame = requestAnimationFrame(() => (hasEntered = true));
+		});
+
+		return () => cancelAnimationFrame(frame);
+	});
+
+	let enterClass = $derived([
+		'transition-opacity duration-300 ease-[cubic-bezier(0.23,1,0.32,1)]',
+		!hasEntered && 'opacity-0'
+	]);
 </script>
 
 <div
 	class={[
 		'pointer-events-none mb-4 hidden px-4 text-center text-balance',
-		isEmpty && 'mb-[calc(50dvh-8rem)] md:mb-8 pointer-events-auto block!'
+		isEmpty && 'mb-[calc(50dvh-8rem)] md:mb-8 pointer-events-auto block!',
+		// out of the page flow on desktop, just above the form: inside the form
+		// wrapper's top padding, so showing it never changes the wrapper's height
+		isEmpty && 'md:absolute md:inset-x-0 md:bottom-[calc(100%-var(--chat-form-padding-top))]'
 	]}
 >
-	<h1 class="mb-2 text-2xl font-semibold tracking-tight md:text-3xl">Hello there</h1>
+	<div style:transition-delay="{CHAT_GREETING_ENTER_DELAY_MS}ms" class={enterClass}>
+		<h1 class="mb-2 text-2xl font-semibold tracking-tight md:text-3xl">Hello there</h1>
 
-	<p class="text-muted-foreground md:text-lg">
-		{serverStore.props?.modalities?.audio ? 'Record audio, type a message ' : 'Type a message'} or upload
-		files to get started
-	</p>
+		<p class="text-muted-foreground md:text-lg">
+			{serverStore.props?.modalities?.audio ? 'Record audio, type a message ' : 'Type a message'} or upload
+			files to get started
+		</p>
+	</div>
 </div>

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

@@ -1,6 +1,7 @@
 <script lang="ts">
 	import ChatTabsItem from './ChatTabsItem.svelte';
 	import ChatTabsNewChatButton from './ChatTabsNewChatButton.svelte';
+	import { SquarePen } from '@lucide/svelte';
 	import { page } from '$app/state';
 	import { ScrollCarousel } from '$lib/components/app';
 	import {
@@ -12,30 +13,35 @@
 	} from '$lib/constants';
 	import { useScrollCarousel } from '$lib/hooks/use-scroll-carousel.svelte';
 	import { chatStore, conversationsStore, tabsStore, uiStore } from '$lib/stores';
+	import { resolveTabView } from '$lib/utils';
 	import { tick } from 'svelte';
 
 	const carousel = useScrollCarousel();
 
-	let activeId = $derived(page.params.id ?? NEW_CHAT_TAB_ID);
+	let view = $derived(resolveTabView(page, tabsStore.openTabs));
+
+	let activeId = $derived(view.activeId);
 
 	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)
-		}))
+		view.ids.map((id) =>
+			id === NEW_CHAT_TAB_ID
+				? { href: view.hrefFor(id), icon: SquarePen, id, name: NEW_CHAT_LABEL }
+				: {
+						href: view.hrefFor(id),
+						id,
+						name:
+							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 showNewChatButton = $derived(!view.ids.includes(NEW_CHAT_TAB_ID));
 
 	let loadingIds = $derived(new Set(chatStore.getAllLoadingChats()));
 
 	function handleClose(id: string) {
-		void tabsStore.close(id, activeId);
+		void tabsStore.close(id, id === activeId, view);
 	}
 
 	function handleStop(id: string, event: MouseEvent) {
@@ -82,9 +88,12 @@
 
 <nav
 	aria-label="Open conversations"
-	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}"
+	class={[
+		'group sticky pl-1 top-0 z-10 hidden md:block chat-top-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
+	]}
 >
 	<div class="relative">
 		<ScrollCarousel
@@ -106,7 +115,9 @@
 			{/each}
 
 			{#if showNewChatButton}
-				<ChatTabsNewChatButton onclick={() => void conversationsStore.openNewChat()} />
+				<ChatTabsNewChatButton
+					onclick={() => void tabsStore.openInNewTab(NEW_CHAT_TAB_ID, { activate: true })}
+				/>
 			{/if}
 		</ScrollCarousel>
 
@@ -123,15 +134,3 @@
 		></div>
 	</div>
 </nav>
-
-<style>
-	.chat-tabs-fade {
-		background: linear-gradient(
-			to bottom,
-			color-mix(in srgb, var(--background) 100%, transparent) 25%,
-			color-mix(in srgb, var(--background) 80%, transparent) 50%,
-			color-mix(in srgb, var(--background) 40%, transparent) 75%,
-			transparent 100%
-		);
-	}
-</style>

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

@@ -1,13 +1,22 @@
 <script lang="ts">
-	import { Loader2, Square, SquarePen, X } from '@lucide/svelte';
+	import { Loader2, Square, 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';
+	import {
+		ICON_CLASS_SM,
+		ICON_CLASS_XS,
+		TAB_SURFACE_ACTIVE_CLASS,
+		TAB_SURFACE_CLASS,
+		TAB_SURFACE_IDLE_CLASS,
+		UI_DATA_ATTRS
+	} from '$lib/constants';
+	import type { Component } from 'svelte';
 
 	interface Tab {
+		href: string;
+		/** Marks the new-chat tab. */
+		icon?: Component;
 		id: string;
-		isNewChat: boolean;
 		name: string;
 	}
 
@@ -33,8 +42,6 @@
 
 	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
@@ -59,18 +66,17 @@
 <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',
+		TAB_SURFACE_CLASS,
+		'relative max-w-52 min-w-0 shrink-0 gap-1 pr-1 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'
+		isActive ? TAB_SURFACE_ACTIVE_CLASS : TAB_SURFACE_IDLE_CLASS
 	)}
 >
 	<a
 		aria-current={isActive ? 'page' : undefined}
 		aria-label={tab.name}
 		class="absolute inset-0 z-0 rounded-lg"
-		{href}
+		href={tab.href}
 		onauxclick={(e) => onAuxClick?.(tab.id, e)}
 		onclick={handleActivate}
 	></a>
@@ -102,8 +108,8 @@
 		</Tooltip.Root>
 	{/if}
 
-	{#if tab.isNewChat}
-		<SquarePen
+	{#if tab.icon}
+		<tab.icon
 			class="pointer-events-none {ICON_CLASS_SM} shrink-0 transition-opacity {contentOpacity}"
 		/>
 	{/if}

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

@@ -13,13 +13,21 @@
 		SidebarNavigationActions,
 		SidebarNavigationConversationList
 	} from '$lib/components/app';
-	import { ROUTES } from '$lib/constants';
+	import {
+		ROUTES,
+		SIDEBAR_COLLAPSE_WIDTH,
+		SIDEBAR_COLLAPSED_WIDTH,
+		SIDEBAR_MAX_WIDTH,
+		SIDEBAR_MIN_WIDTH,
+		SIDEBAR_OVERSTRETCH_PX,
+		SIDEBAR_SETTLE_MS
+	} from '$lib/constants';
 	import { TooltipSide } from '$lib/enums';
 	import { useKeyboardShortcuts } from '$lib/hooks/use-keyboard-shortcuts.svelte';
 	import { useMarqueeSelection } from '$lib/hooks/use-marquee-selection.svelte';
 	import { RouterService } from '$lib/services/router.service';
 	import { chatStore, conversationsStore, deviceStore, settingsStore, uiStore } from '$lib/stores';
-	import { buildConversationTree } from '$lib/utils';
+	import { buildConversationTree, remToPx } from '$lib/utils';
 	import { circIn } from 'svelte/easing';
 	import { SvelteSet } from 'svelte/reactivity';
 	import { fade } from 'svelte/transition';
@@ -311,6 +319,120 @@
 		}
 	}
 
+	let asideEl = $state<HTMLElement | null>(null);
+	/** Set briefly after a drag ends past a bound, while the width springs back to it. */
+	let isSettling = $state(false);
+
+	// the tab bar sits outside the sidebar, so the width lives on the root
+	$effect(() => {
+		const width = uiStore.sidebarWidth;
+
+		if (width === null) {
+			document.documentElement.style.removeProperty('--sidebar-width');
+		} else {
+			document.documentElement.style.setProperty('--sidebar-width', `${width}px`);
+		}
+	});
+
+	// labels and the conversation list fade with the live width, so they follow
+	// both the toggle's width transition and a drag below the minimum width
+	$effect(() => {
+		const el = asideEl;
+
+		if (!el) return;
+
+		const collapsedPx = remToPx(SIDEBAR_COLLAPSED_WIDTH);
+		const minPx = remToPx(SIDEBAR_MIN_WIDTH);
+		const observer = new ResizeObserver(([entry]) => {
+			const fade = (entry.contentRect.width - collapsedPx) / (minPx - collapsedPx);
+
+			el.style.setProperty('--sidebar-fade', String(Math.min(1, Math.max(0, fade))));
+		});
+
+		observer.observe(el);
+
+		return () => observer.disconnect();
+	});
+
+	function handleResizePointerDown(event: PointerEvent) {
+		if (!asideEl || isOnMobile) return;
+
+		event.preventDefault();
+
+		const handle = event.currentTarget as HTMLElement;
+		const left = asideEl.getBoundingClientRect().left;
+		const minPx = remToPx(SIDEBAR_MIN_WIDTH);
+		const maxPx = remToPx(SIDEBAR_MAX_WIDTH);
+		const collapsePx = remToPx(SIDEBAR_COLLAPSE_WIDTH);
+		const startWidth = uiStore.sidebarWidth;
+
+		let willCollapse = false;
+
+		handle.setPointerCapture(event.pointerId);
+		uiStore.isSidebarResizing = true;
+
+		/** Rubber-band offset: grows with the distance, ever slower, never past the overstretch. */
+		function stretch(distance: number): number {
+			return SIDEBAR_OVERSTRETCH_PX * (1 - 1 / (distance / SIDEBAR_OVERSTRETCH_PX + 1));
+		}
+
+		function onPointerMove(moveEvent: PointerEvent) {
+			const pointerWidth = moveEvent.clientX - left;
+
+			// past either bound the edge resists harder the further it is pulled,
+			// until a pull below the collapse width snaps it to the pointer
+			willCollapse = pointerWidth < collapsePx;
+
+			const width = willCollapse
+				? Math.max(0, pointerWidth)
+				: pointerWidth < minPx
+					? minPx - stretch(minPx - pointerWidth)
+					: pointerWidth > maxPx
+						? maxPx + stretch(pointerWidth - maxPx)
+						: pointerWidth;
+
+			// persisted once on release, not on every move
+			uiStore.sidebarWidth = Math.round(width);
+		}
+
+		function onPointerUp() {
+			handle.removeEventListener('pointermove', onPointerMove);
+			handle.removeEventListener('pointerup', onPointerUp);
+			handle.removeEventListener('pointercancel', onPointerCancel);
+
+			uiStore.isSidebarResizing = false;
+
+			// it reopens at its width from before the drag, without springing on the way out
+			if (willCollapse) {
+				uiStore.sidebarWidth = startWidth;
+				toggleExpandedMode();
+
+				return;
+			}
+
+			if (uiStore.sidebarWidth === startWidth) return;
+
+			const width = uiStore.sidebarWidth ?? minPx;
+			const clamped = Math.min(maxPx, Math.max(minPx, width));
+
+			if (clamped !== width) {
+				isSettling = true;
+				setTimeout(() => (isSettling = false), SIDEBAR_SETTLE_MS);
+			}
+
+			uiStore.setSidebarWidth(Math.round(clamped));
+		}
+
+		function onPointerCancel() {
+			willCollapse = false;
+			onPointerUp();
+		}
+
+		handle.addEventListener('pointermove', onPointerMove);
+		handle.addEventListener('pointerup', onPointerUp);
+		handle.addEventListener('pointercancel', onPointerCancel);
+	}
+
 	/** A dialog owns the screen on a phone, so the expanded list steps aside for it. */
 	function openDialog(open: () => void): void {
 		if (deviceStore.isMobile) {
@@ -325,6 +447,7 @@
 
 {#if innerWidth > 768 || !page.url.hash.includes(ROUTES.SEARCH)}
 	<aside
+		bind:this={asideEl}
 		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)]',
@@ -337,7 +460,9 @@
 			'rounded-3xl md:rounded-2xl',
 			'flex flex-col justify-between',
 			'md:transition-[width,padding] duration-200 ease-out',
-			isStripExpanded && 'md:w-72 md:bg-muted/60 md:backdrop-blur-xl shadow-md',
+			uiStore.isSidebarResizing && 'md:duration-0',
+			isSettling && 'md:duration-300 md:ease-[cubic-bezier(0.34,1.56,0.64,1)]',
+			isStripExpanded && 'md:w-(--sidebar-width) md:bg-muted/60 md:backdrop-blur-xl shadow-md',
 			!isStripExpanded && 'md:w-12',
 			uiStore.isSidebarExpanded && 'is-expanded'
 		]}
@@ -421,7 +546,7 @@
 			/>
 
 			{#if uiStore.isSidebarExpanded || isOnMobile}
-				<div class="flex min-h-0 flex-1 flex-col overflow-y-auto">
+				<div class="flex min-h-0 flex-1 flex-col overflow-y-auto opacity-(--sidebar-fade)">
 					<SidebarNavigationConversationList
 						{allSelectedArePinned}
 						allVisibleSelected={visibleSelectionStats.visibleCount > 0 &&
@@ -454,6 +579,33 @@
 				</div>
 			{/if}
 		</div>
+
+		{#if uiStore.isSidebarExpanded && !isOnMobile}
+			<!-- svelte-ignore a11y_no_static_element_interactions -->
+			<div
+				class="peer absolute inset-y-3 -right-1 z-20 hidden w-2 cursor-ew-resize touch-none select-none md:block"
+				data-slot="sidebar-resize-handle"
+				ondblclick={() => uiStore.setSidebarWidth(null)}
+				onpointerdown={handleResizePointerDown}
+				title="Drag to resize or hide, double-click to reset"
+			></div>
+
+			<!-- the edge strip is clipped to the sidebar's rounded corners without
+			     clipping the sidebar itself; faint on hover, solid while dragging -->
+			<div
+				class={[
+					'pointer-events-none absolute inset-0 hidden overflow-hidden rounded-[inherit] transition-opacity duration-150 md:block',
+					uiStore.isSidebarResizing ? 'opacity-100' : 'opacity-0 peer-hover:opacity-100'
+				]}
+			>
+				<div
+					class={[
+						'absolute inset-y-0 right-0 w-1.5',
+						uiStore.isSidebarResizing ? 'bg-foreground/15' : 'bg-foreground/6'
+					]}
+				></div>
+			</div>
+		{/if}
 	</aside>
 {/if}
 

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

@@ -175,8 +175,10 @@
 							{@render itemIcon(item.icon)}
 
 							{#if showIcons}
-								<span in:fade={itemTransition} out:fade={itemTransition} class="min-w-0 truncate"
-									>{item.tooltip}</span
+								<span
+									in:fade={itemTransition}
+									out:fade={itemTransition}
+									class="min-w-0 truncate opacity-(--sidebar-fade)">{item.tooltip}</span
 								>
 							{/if}
 						</span>

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

@@ -5,6 +5,7 @@
 		ListChecks,
 		Loader2,
 		MoreHorizontal,
+		PanelTopOpen,
 		Pencil,
 		Pin,
 		PinOff,
@@ -17,7 +18,7 @@
 	import * as Tooltip from '$lib/components/ui/tooltip';
 	import { FORK_TREE_DEPTH_PADDING, ICON_CLASS_DEFAULT, UI_DATA_ATTRS } from '$lib/constants';
 	import { RouterService } from '$lib/services/router.service';
-	import { chatStore, conversationsStore } from '$lib/stores';
+	import { chatStore, conversationsStore, settingsStore, tabsStore } from '$lib/stores';
 	import { onMount } from 'svelte';
 
 	interface Props {
@@ -72,6 +73,20 @@
 		onStop?.(conversation.id);
 	}
 
+	let canOpenInNewTab = $derived(Boolean(settingsStore.config.conversationTabs));
+
+	function openInNewTab() {
+		void tabsStore.openInNewTab(conversation.id);
+	}
+
+	// middle-click opens a background tab, like a link in a browser
+	function handleAuxClick(event: MouseEvent) {
+		if (event.button !== 1 || isSelectionMode || !canOpenInNewTab) return;
+
+		event.preventDefault();
+		openInNewTab();
+	}
+
 	function handleTogglePin() {
 		conversationsStore.toggleConversationPin(conversation.id);
 	}
@@ -104,6 +119,8 @@
 	function handleSelect(event: MouseEvent) {
 		if (isSelectionMode) {
 			onSelectionClick?.(conversation.id, { shiftKey: event.shiftKey });
+		} else if ((event.metaKey || event.ctrlKey) && canOpenInNewTab) {
+			openInNewTab();
 		} else {
 			onSelect?.(conversation.id);
 		}
@@ -161,6 +178,7 @@
 		? 'is-selection-mode'
 		: ''} px-2"
 	{...{ [UI_DATA_ATTRS.CONVERSATION_ROW]: conversation.id }}
+	onauxclick={handleAuxClick}
 	onclick={(e) => handleSelect(e)}
 	onfocusin={handleMouseOver}
 	onfocusout={(e) => {
@@ -245,13 +263,26 @@
 			<DropdownMenuActions
 				bind:open={dropdownOpen}
 				actions={[
+					...(canOpenInNewTab
+						? [
+								{
+									icon: PanelTopOpen,
+									label: 'Open in new tab',
+									onclick: (e: Event) => {
+										e.stopPropagation();
+										openInNewTab();
+									}
+								}
+							]
+						: []),
 					{
 						icon: conversation.pinned ? PinOff : Pin,
 						label: conversation.pinned ? 'Unpin' : 'Pin',
 						onclick: (e: Event) => {
 							e.stopPropagation();
 							handleTogglePin();
-						}
+						},
+						separator: canOpenInNewTab
 					},
 					{
 						icon: Pencil,

tools/ui/src/lib/constants/auto-scroll.constants.ts

@@ -1,10 +1,16 @@
 export const AUTO_SCROLL_INTERVAL = 100;
+// How long after the reader's wheel, touch or scroll key an upward scroll
+// still counts as theirs. Scrolls with no such input (scroll anchoring,
+// blocks re-rendering, programmatic scrolls) never stop auto-scroll.
+export const USER_SCROLL_INTENT_MS = 250;
 // Conversation landing: the page keeps growing after the first bottom pin
-// without DOM mutations (content-visibility size realizations, syntax
-// highlight passes), so the pin repeats every frame until the height holds
-// for this many consecutive frames, bounded by the time cap below.
-export const LANDING_STABLE_FRAMES = 10;
-export const LANDING_SETTLE_MAX_MS = 1000;
+// (lazy rows mounting, content-visibility size realizations, syntax
+// highlight passes), so the pin holds until the page has not resized for
+// this long or the reader interacts.
+export const LANDING_QUIET_MS = 1000;
+// Rows at the end of the list that mount eagerly, so the first painted
+// frame of a conversation already shows real content at the bottom.
+export const LANDING_EAGER_ROWS = 8;
 // Chat main view: tight threshold because scroll-here events come from
 // discrete assistant-message appends.
 export const AUTO_SCROLL_AT_BOTTOM_THRESHOLD = 10;

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

@@ -6,3 +6,25 @@ export const PROMPT_CONTENT_SEPARATOR = '\n\n';
 export const CLIPBOARD_CONTENT_QUOTE_PREFIX = '"';
 export const PROMPT_TRIGGER_PREFIX = '/';
 export const NEW_CHAT_DRAFT_KEY = '__new_chat__';
+
+/**
+ * The chat form moving between its centered (new chat) and bottom positions:
+ * it fades out drifting toward where it goes, then fades back in where it
+ * lands, instead of sliding across the column. The exit uses a gentler curve
+ * so the form stays visible while it travels; the entry settles quickly.
+ */
+export const CHAT_FORM_MOVE_OUT_MS = 140;
+export const CHAT_FORM_MOVE_IN_MS = 180;
+// opening a chat moves the form to the bottom: quicker, so the chat is
+// there right away
+export const CHAT_FORM_MOVE_TO_BOTTOM_OUT_MS = 80;
+export const CHAT_FORM_MOVE_TO_BOTTOM_IN_MS = 120;
+export const CHAT_FORM_MOVE_SHIFT_PX = 16;
+export const CHAT_FORM_MOVE_OUT_EASING = 'cubic-bezier(0.33, 1, 0.68, 1)';
+export const CHAT_FORM_MOVE_IN_EASING = 'cubic-bezier(0.23, 1, 0.32, 1)';
+
+/**
+ * The new-chat greeting fades in as the form settles: it starts
+ * halfway through the form's fade-in so both read as one motion.
+ */
+export const CHAT_GREETING_ENTER_DELAY_MS = CHAT_FORM_MOVE_IN_MS / 2;

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

@@ -9,10 +9,11 @@ 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.
+ * is `--sidebar-width` wide (user resizable) and the collapsed one `md:w-12`;
+ * these hold the fully tuned `max-w-[calc(100vw-?)]` 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)]'
+	EXPANDED_SIDEBAR: 'max-w-[calc(100vw-var(--sidebar-width)-2rem)]'
 } as const;

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

@@ -25,6 +25,18 @@ export const DIALOG_SUBMENU_CONTENT = 'w-60';
 export const PANEL_CLOSE_MOBILE_CLASS =
 	'max-md:grid max-md:size-10 max-md:place-items-center max-md:rounded-full max-md:text-muted-foreground max-md:opacity-100 max-md:hover:bg-accent max-md:hover:text-accent-foreground';
 
+/** Shape of a chat tab, shared with the breadcrumb items next to the project tabs. */
+export const TAB_SURFACE_CLASS =
+	'flex h-8 items-center rounded-lg border text-sm whitespace-nowrap backdrop-blur-xl';
+
+/** Tab surface at rest: invisible until hovered. */
+export const TAB_SURFACE_IDLE_CLASS =
+	'border-transparent hover:bg-primary/10 hover:border-border/10 hover:shadow-sm';
+
+/** Tab surface of the open tab. */
+export const TAB_SURFACE_ACTIVE_CLASS =
+	'bg-muted/60 border-border/10 shadow-sm text-accent-foreground hover:bg-primary/15';
+
 /** Selects the focused chat-form input (either renderer) to restore focus after model actions. */
 export const CHAT_INPUT_FOCUS_SELECTOR =
 	'[data-slot="input-area"] textarea, [data-slot="input-area"] [contenteditable="true"]';

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

@@ -36,6 +36,8 @@ export const RECENT_MODELS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.recentModels`
 export const HIDDEN_MODELS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.hiddenModels`;
 export const REASONING_EFFORT_DEFAULT_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.reasoningEffortDefault`;
 export const CONVERSATION_TABS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.conversationTabs`;
+export const SIDEBAR_EXPANDED_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.sidebarExpanded`;
+export const SIDEBAR_WIDTH_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.sidebarWidth`;
 export const USER_OVERRIDES_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.userOverrides`;
 export const DISMISSED_RECOMMENDED_MCP_SERVERS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.dismissedRecommendedMcpServers`;
 

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

@@ -68,6 +68,18 @@ export const RECENT_MODEL_LIMIT = 20;
 export const ICON_STRIP_TRANSITION_DURATION = 150;
 export const ICON_STRIP_TRANSITION_DELAY_MULTIPLIER = 50;
 
+/** Collapsed desktop sidebar strip width, matching its `md:w-12`. */
+export const SIDEBAR_COLLAPSED_WIDTH = '3rem';
+/** Bounds the expanded desktop sidebar can be dragged between. */
+export const SIDEBAR_MIN_WIDTH = '18rem';
+export const SIDEBAR_MAX_WIDTH = '32rem';
+/** Dragging the sidebar edge narrower than this snaps it to the pointer; letting go hides it. */
+export const SIDEBAR_COLLAPSE_WIDTH = '9rem';
+/** How far past its min or max width a drag can stretch the sidebar before it resists fully. */
+export const SIDEBAR_OVERSTRETCH_PX = 80;
+/** Spring-back duration when a drag ends outside the width bounds. */
+export const SIDEBAR_SETTLE_MS = 300;
+
 /** Max height for tool-result code blocks (json / source / diff / streaming code). */
 export const MAX_HEIGHT_CODE_BLOCK = '22rem';
 

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

@@ -14,9 +14,11 @@ export enum KeyboardKey {
 	E_UPPER = 'E',
 	ENTER = 'Enter',
 	ESCAPE = 'Escape',
+	HOME = 'Home',
 	K_LOWER = 'k',
 	O_LOWER = 'o',
 	O_UPPER = 'O',
+	PAGE_UP = 'PageUp',
 	SPACE = ' ',
 	TAB = 'Tab',
 	X_LOWER = 'x',

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

@@ -1,4 +1,11 @@
-import { AUTO_SCROLL_AT_BOTTOM_THRESHOLD, AUTO_SCROLL_INTERVAL } from '$lib/constants';
+import {
+	AUTO_SCROLL_AT_BOTTOM_THRESHOLD,
+	AUTO_SCROLL_INTERVAL,
+	USER_SCROLL_INTENT_MS
+} from '$lib/constants';
+import { KeyboardKey } from '$lib/enums';
+
+const SCROLL_UP_KEYS: string[] = [KeyboardKey.ARROW_UP, KeyboardKey.HOME, KeyboardKey.PAGE_UP];
 
 export interface AutoScrollOptions {
 	disabled?: boolean;
@@ -11,17 +18,26 @@ export interface AutoScrollOptions {
  * - Auto-scrolls to bottom during streaming/loading
  * - Stops auto-scroll when user manually scrolls up
  * - Resumes auto-scroll when user scrolls back to bottom
+ *
+ * Only the reader's own input (wheel, touch, scroll keys, scrollbar drag)
+ * counts as scrolling up: the page also moves up on its own while a
+ * response streams (scroll anchoring, blocks re-rendering shorter, the
+ * send animation), and that must not stop following the response.
  */
 export class AutoScrollController {
 	private _autoScrollEnabled = $state(true);
 	private _container: HTMLElement | undefined;
 	private _disabled: boolean;
+	private _draggingScrollbar = false;
 	private _lastScrollTop = $state(0);
+	private _lastUserIntentAt = 0;
 	private _mutationObserver: MutationObserver | null = null;
 	private _observerEnabled = false;
 	private _rafPending = false;
 	private _scrollInterval: ReturnType<typeof setInterval> | undefined;
+	private _unbindUserIntent: (() => void) | null = null;
 	private _userScrolledUp = $state(false);
+
 	get autoScrollEnabled(): boolean {
 		return this._autoScrollEnabled;
 	}
@@ -40,6 +56,8 @@ export class AutoScrollController {
 	destroy(): void {
 		this.stopInterval();
 		this._doStopObserving();
+		this._unbindUserIntent?.();
+		this._unbindUserIntent = null;
 	}
 
 	/**
@@ -62,8 +80,10 @@ export class AutoScrollController {
 		const distanceFromBottom = scrollHeight - clientHeight - scrollTop;
 		const isScrollingUp = scrollTop < this._lastScrollTop;
 		const isAtBottom = distanceFromBottom < AUTO_SCROLL_AT_BOTTOM_THRESHOLD;
+		const isUserScrolling =
+			this._draggingScrollbar || performance.now() - this._lastUserIntentAt < USER_SCROLL_INTENT_MS;
 
-		if (isScrollingUp && !isAtBottom) {
+		if (isScrollingUp && !isAtBottom && isUserScrolling) {
 			this._userScrolledUp = true;
 			this._autoScrollEnabled = false;
 		} else if (isAtBottom && this._userScrolledUp) {
@@ -100,8 +120,14 @@ export class AutoScrollController {
 	 */
 	setContainer(container: HTMLElement | undefined): void {
 		this._doStopObserving();
+		this._unbindUserIntent?.();
+		this._unbindUserIntent = null;
 		this._container = container;
 
+		if (container) {
+			this._unbindUserIntent = this._bindUserIntent(container);
+		}
+
 		if (this._observerEnabled && container && !this._disabled) {
 			this._doStartObserving();
 		}
@@ -185,6 +211,52 @@ export class AutoScrollController {
 		}
 	}
 
+	/**
+	 * Records the reader's own scroll input, so `handleScroll` can tell their
+	 * scrolling from the page moving on its own.
+	 */
+	private _bindUserIntent(container: HTMLElement): () => void {
+		const isPage = container === document.documentElement;
+		const target: HTMLElement | Window = isPage ? window : container;
+		const markIntent = () => {
+			this._lastUserIntentAt = performance.now();
+		};
+		const onWheel = (event: WheelEvent) => {
+			if (event.deltaY < 0) markIntent();
+		};
+		const onKeyDown = (event: KeyboardEvent) => {
+			if (
+				SCROLL_UP_KEYS.includes(event.key) ||
+				(event.key === KeyboardKey.SPACE && event.shiftKey)
+			) {
+				markIntent();
+			}
+		};
+		// a press right of the content box lands on the scrollbar
+		const onPointerDown = (event: PointerEvent) => {
+			const left = isPage ? 0 : container.getBoundingClientRect().left;
+
+			this._draggingScrollbar = event.clientX >= left + container.clientWidth;
+		};
+		const onPointerUp = () => {
+			this._draggingScrollbar = false;
+		};
+
+		target.addEventListener('wheel', onWheel as EventListener, { passive: true });
+		target.addEventListener('touchmove', markIntent, { passive: true });
+		target.addEventListener('keydown', onKeyDown as EventListener);
+		target.addEventListener('pointerdown', onPointerDown as EventListener);
+		window.addEventListener('pointerup', onPointerUp);
+
+		return () => {
+			target.removeEventListener('wheel', onWheel as EventListener);
+			target.removeEventListener('touchmove', markIntent);
+			target.removeEventListener('keydown', onKeyDown as EventListener);
+			target.removeEventListener('pointerdown', onPointerDown as EventListener);
+			window.removeEventListener('pointerup', onPointerUp);
+		};
+	}
+
 	private _doStartObserving(): void {
 		if (!this._container || this._mutationObserver) return;
 

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

@@ -8,25 +8,83 @@
  */
 
 import type { AutoScrollController } from './use-auto-scroll.svelte';
-import { afterNavigate, beforeNavigate } from '$app/navigation';
+import { afterNavigate, beforeNavigate, disableScrollHandling, onNavigate } from '$app/navigation';
+import { LANDING_QUIET_MS } from '$lib/constants';
+import { onDestroy } from 'svelte';
+
+const LANDING_INTENT_EVENTS = ['keydown', 'pointerdown', 'touchstart', 'wheel'] as const;
 
 export function useChatScreenScroll(autoScroll: AutoScrollController) {
 	let chatScrollContainer: HTMLElement | undefined = $state();
 	let isNavigating = $state(false);
+	let stopLanding: (() => void) | null = null;
 
 	function handleScroll(event: UIEvent) {
-		// Ignore scroll events caused by navigation layout changes or by our own
-		// programmatic scrolls so they don't accidentally disable auto-scroll.
+		// Ignore scroll events caused by navigation layout changes. Programmatic
+		// and layout scrolls cannot stop auto-scroll: only the reader's own
+		// input counts as scrolling up (see AutoScrollController.handleScroll).
 		if (isNavigating || !event.isTrusted) return;
 
 		autoScroll.handleScroll();
 	}
 
+	/**
+	 * Pins the page to its bottom while a freshly opened conversation settles.
+	 * ResizeObserver callbacks run after layout and before paint, so every
+	 * height change (lazy rows mounting, content-visibility realizing a row's
+	 * size, highlight passes) is corrected in the frame that caused it and no
+	 * frame is ever painted off the bottom. Stops on the reader's first input
+	 * or once the page has not resized for `LANDING_QUIET_MS`.
+	 */
+	function landAtBottom() {
+		stopLanding?.();
+
+		const container = chatScrollContainer;
+
+		if (!container) return;
+
+		let quietTimer: ReturnType<typeof setTimeout> | undefined;
+
+		const pin = () => {
+			container.scrollTop = container.scrollHeight;
+			clearTimeout(quietTimer);
+			quietTimer = setTimeout(stop, LANDING_QUIET_MS);
+		};
+		const observer = new ResizeObserver(pin);
+		const stop = () => {
+			observer.disconnect();
+			clearTimeout(quietTimer);
+
+			for (const type of LANDING_INTENT_EVENTS) {
+				window.removeEventListener(type, stop, { capture: true });
+			}
+
+			if (stopLanding === stop) stopLanding = null;
+		};
+
+		stopLanding = stop;
+
+		for (const type of LANDING_INTENT_EVENTS) {
+			window.addEventListener(type, stop, { capture: true, passive: true });
+		}
+
+		pin();
+		observer.observe(document.body);
+	}
+
 	beforeNavigate(() => {
 		isNavigating = true;
+		stopLanding?.();
 		autoScroll.resetScrollState();
 	});
 
+	// Opening a conversation lands at its bottom once its messages load, so
+	// SvelteKit's own reset (or restore, on back/forward) would only move the
+	// previous conversation, still on screen, to its top for a few frames.
+	onNavigate((navigation) => {
+		if (navigation.to?.params?.id) disableScrollHandling();
+	});
+
 	afterNavigate(() => {
 		setTimeout(() => {
 			isNavigating = false;
@@ -39,10 +97,13 @@ export function useChatScreenScroll(autoScroll: AutoScrollController) {
 		autoScroll.setContainer(chatScrollContainer);
 	});
 
+	onDestroy(() => stopLanding?.());
+
 	return {
 		get chatScrollContainer() {
 			return chatScrollContainer;
 		},
-		handleScroll
+		handleScroll,
+		landAtBottom
 	};
 }

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

@@ -63,7 +63,7 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) {
 				page.params.id ?? (page.route.id === '/(chat)' ? NEW_CHAT_TAB_ID : undefined);
 
 			if (activeId) {
-				void tabsStore.close(activeId, activeId);
+				void tabsStore.close(activeId, true);
 			}
 		}
 

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

@@ -156,7 +156,7 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 				tabsStore.removeTabs([...idsToRemove].filter((id) => id !== activeId));
 				this.clearActiveConversation();
-				await tabsStore.close(activeId, activeId);
+				await tabsStore.closeConversation(activeId);
 			} else {
 				tabsStore.removeTabs([...idsToRemove]);
 			}
@@ -340,7 +340,7 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 					tabsStore.removeTabs([...idsToRemove].filter((id) => id !== activeId));
 					this.clearActiveConversation();
-					await tabsStore.close(activeId, activeId);
+					await tabsStore.closeConversation(activeId);
 				} else {
 					tabsStore.removeTabs([...idsToRemove]);
 				}
@@ -361,7 +361,7 @@ class ConversationsStore implements ConversationsPreferencesHost {
 
 				if (this.activeConversation?.id === convId) {
 					this.clearActiveConversation();
-					await tabsStore.close(convId, convId);
+					await tabsStore.closeConversation(convId);
 				} else {
 					tabsStore.removeTabs([convId]);
 				}

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

@@ -1,16 +1,18 @@
 /**
  * tabsStore - Reactive State Store for Browser-Style Conversation Tabs
  *
- * Tracks which conversations and the new-chat screen are open as tabs in
- * the chat layout, in order. Real conversation tabs are `#/chat/<id>`
- * routes; the new-chat tab is the bare `#/` route, represented here by the
- * `NEW_CHAT_TAB_ID` sentinel (see {@link NEW_CHAT_TAB_ID}).
+ * Tracks which conversations and the new-chat screen are open as tabs in the
+ * chat layout, in order. Conversation tabs hold the conversation id; the
+ * new-chat tab of the bare `#/` route is the `NEW_CHAT_TAB_ID` sentinel.
+ *
+ * Tabs open like in a browser: a plain navigation shows the new target in
+ * the current tab, and only `openInNewTab()` adds a tab.
  *
  * **Architecture & Relationships:**
  * - **conversationsStore**: owns conversation data; calls `removeTabs()` /
- *   `close()` when conversations are deleted. This store never imports it,
- *   so there is no circular dependency - tab names are resolved by the
- *   ChatTabs component from conversationsStore.
+ *   `closeConversation()` when conversations are deleted. This store never
+ *   imports it, so there is no circular dependency - tab names are resolved
+ *   by the ChatTabs component from conversationsStore.
  * - Tab order persists to localStorage and is pruned against the loaded
  *   conversation list on init. The new-chat tab is kept across reloads.
  */
@@ -19,42 +21,57 @@ import { browser } from '$app/environment';
 import { goto } from '$app/navigation';
 import { CONVERSATION_TABS_LOCALSTORAGE_KEY, NEW_CHAT_TAB_ID, ROUTES } from '$lib/constants';
 import { RouterService } from '$lib/services/router.service';
+import type { TabView } from '$lib/utils/chat-tabs';
 import { untrack } from 'svelte';
 
-class TabsStore {
+export class TabsStore {
 	/** Ordered tab ids: conversation ids and the `NEW_CHAT_TAB_ID` sentinel */
 	openTabs = $state<string[]>([]);
 
+	/** Target of the last route the chat layout reported, the tab a plain navigation reuses */
+	private current: string | null = null;
+
 	/** False until init() has read the persisted tabs; save() is a no-op before that */
 	private initialized = false;
 
+	/** Route synced before init, applied once the persisted tabs are loaded */
+	private pendingTarget: string | null = null;
+
 	/** Navigate to a tab (the new-chat sentinel maps to the bare `#/` route) */
 	async activate(id: string): Promise<void> {
-		await goto(id === NEW_CHAT_TAB_ID ? ROUTES.START : RouterService.chat(id));
+		await goto(this.hrefFor(id));
 	}
 
 	/** Remove all tabs (e.g. after deleting all conversations) */
 	clear(): void {
 		this.openTabs = [];
+		this.current = null;
 		this.save();
 	}
 
 	/**
 	 * Close a tab. When it belongs to the active route, navigate to the left
-	 * neighbor (or the right one when the closed tab was leftmost), falling
-	 * back to the new-chat screen when no tabs remain.
+	 * neighbor (or the right one when the closed tab was leftmost) among the
+	 * tabs the view shows, falling back to the view's start when none remain.
 	 * @param id - Tab id to close
-	 * @param activeTabId - Tab id of the current route, if any
+	 * @param isActive - Whether the tab is the current route
+	 * @param view - Visible tabs and their links; all tabs when omitted
 	 */
-	async close(id: string, activeTabId: string | null): Promise<void> {
-		const idx = this.openTabs.indexOf(id);
+	async close(
+		id: string,
+		isActive: boolean,
+		view: Pick<TabView, 'fallbackHref' | 'hrefFor' | 'ids'> = {
+			fallbackHref: ROUTES.START,
+			hrefFor: (tabId) => this.hrefFor(tabId),
+			ids: this.openTabs
+		}
+	): Promise<void> {
+		const idx = view.ids.indexOf(id);
 
-		if (idx === -1) {
-			// tab not tracked (e.g. Conversation tabs are off); still fall back to
-			// the new-chat screen when closing the active conversation
-			if (id === activeTabId) {
-				await goto(ROUTES.START);
-			}
+		if (!this.openTabs.includes(id)) {
+			// tab not tracked (e.g. Conversation tabs are off); still leave the
+			// closed route for the view's start
+			if (isActive) await goto(view.fallbackHref);
 
 			return;
 		}
@@ -62,23 +79,27 @@ class TabsStore {
 		this.openTabs = this.openTabs.filter((tabId) => tabId !== id);
 		this.save();
 
-		if (id !== activeTabId) return;
+		if (!isActive) return;
 
-		const target = (idx > 0 ? this.openTabs[idx - 1] : this.openTabs[0]) ?? null;
+		const visible = view.ids.filter((tabId) => tabId !== id);
+		const next = (idx > 0 ? visible[idx - 1] : visible[0]) ?? null;
 
-		if (target) {
-			await goto(target === NEW_CHAT_TAB_ID ? ROUTES.START : RouterService.chat(target));
-		} else {
-			await goto(ROUTES.START);
-		}
+		await goto(next ? view.hrefFor(next) : view.fallbackHref);
+	}
+
+	/** Close the tab of the conversation open in the route, e.g. after deleting it */
+	async closeConversation(id: string): Promise<void> {
+		await this.close(id, true);
+	}
+
+	hrefFor(id: string): string {
+		return id === NEW_CHAT_TAB_ID ? ROUTES.START : RouterService.chat(id);
 	}
 
 	/**
-	 * Load persisted tabs, dropping conversation ids that no longer exist.
+	 * Load persisted tabs, dropping conversation ids that no longer exist,
+	 * then apply the route synced while loading.
 	 * Called once from initStores() after conversations are loaded.
-	 * Merges with (rather than replaces) current openTabs: the chat layout
-	 * syncs the route's tab before this async init completes, and replacing
-	 * here would drop it.
 	 * @param validIds - Ids of conversations present in the database
 	 */
 	init(validIds: string[]): void {
@@ -86,16 +107,29 @@ class TabsStore {
 
 		// the new-chat sentinel is a pseudo-tab, not a conversation, but it is
 		// still kept so a reload on `#/` does not drop the tab the user is on
-		const isLive = (id: string) => validIds.includes(id) || id === NEW_CHAT_TAB_ID;
-		const persisted = this.load().filter(isLive);
-		// tabs already in openTabs come from the live route, so they stay as they
-		// are: `validIds` is a snapshot and a conversation created while the list
-		// was loading is not in it
-		const extras = this.openTabs.filter((id) => !persisted.includes(id));
-
-		this.openTabs = [...persisted, ...extras];
+		this.openTabs = this.load().filter((id) => id === NEW_CHAT_TAB_ID || validIds.includes(id));
 		this.initialized = true;
 		this.save();
+
+		if (this.pendingTarget) {
+			this.syncWithRoute(this.pendingTarget);
+			this.pendingTarget = null;
+		}
+	}
+
+	/**
+	 * Add a tab at the end. Without `activate` the tab opens in the
+	 * background, like a middle-click in a browser.
+	 * @param id - What the new tab shows
+	 * @param options.activate - Also navigate to the tab
+	 */
+	async openInNewTab(id: string, { activate = false } = {}): Promise<void> {
+		if (!this.openTabs.includes(id)) {
+			this.openTabs = [...this.openTabs, id];
+			this.save();
+		}
+
+		if (activate) await this.activate(id);
 	}
 
 	/**
@@ -104,8 +138,7 @@ class TabsStore {
 	 * @param ids - Tab ids to drop
 	 */
 	removeTabs(ids: string[]): void {
-		const removed = new Set(ids);
-		const next = this.openTabs.filter((id) => !removed.has(id));
+		const next = this.openTabs.filter((id) => !ids.includes(id));
 
 		if (next.length !== this.openTabs.length) {
 			this.openTabs = next;
@@ -115,19 +148,35 @@ class TabsStore {
 
 	/**
 	 * Sync the tab strip with the route. Called from the chat layout on every
-	 * navigation, so any way of reaching a conversation or new-chat tab opens
-	 * a tab for it.
-	 * @param id - The conversation (or temporary new-chat) id of the route
+	 * navigation. A target already open activates its tab; otherwise the tab
+	 * of the previous route now shows it, and only a navigation from a page
+	 * without a tab adds one.
+	 * @param target - Conversation id or `NEW_CHAT_TAB_ID`
 	 */
-	syncWithRoute(id: string): void {
+	syncWithRoute(target: string): void {
 		// untrack: callers invoke this from an effect keyed on the route, and
 		// reading openTabs here would subscribe that effect to openTabs too -
 		// closing the active tab would then re-run the effect and re-add the tab
 		untrack(() => {
-			if (!this.openTabs.includes(id)) {
-				this.openTabs = [...this.openTabs, id];
-				this.save();
+			if (!this.initialized) {
+				this.pendingTarget = target;
+
+				return;
 			}
+
+			const previous = this.current;
+
+			this.current = target;
+
+			if (this.openTabs.includes(target)) return;
+
+			const idx = previous ? this.openTabs.indexOf(previous) : -1;
+
+			this.openTabs =
+				idx === -1
+					? [...this.openTabs, target]
+					: this.openTabs.map((id, i) => (i === idx ? target : id));
+			this.save();
 		});
 	}
 
@@ -144,7 +193,7 @@ class TabsStore {
 
 	private save(): void {
 		// never write before init has read the persisted tabs, or an early
-		// route sync (layout effect runs before async init) would clobber them
+		// write would clobber them
 		if (!browser || !this.initialized) return;
 
 		localStorage.setItem(CONVERSATION_TABS_LOCALSTORAGE_KEY, JSON.stringify(this.openTabs));

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

@@ -6,11 +6,22 @@
  * controls and the chat tab bar reacts to).
  */
 
+import { browser } from '$app/environment';
 import { goto } from '$app/navigation';
 import { page } from '$app/state';
-import { MODELS_ROUTE_ID, ROUTES, SETTINGS_KEYS } from '$lib/constants';
+import {
+	MODELS_ROUTE_ID,
+	ROUTES,
+	SETTINGS_KEYS,
+	SIDEBAR_EXPANDED_LOCALSTORAGE_KEY,
+	SIDEBAR_MAX_WIDTH,
+	SIDEBAR_MIN_WIDTH,
+	SIDEBAR_WIDTH_LOCALSTORAGE_KEY
+} from '$lib/constants';
+import { deviceStore } from '$lib/stores/device.svelte';
 import { settingsStore } from '$lib/stores/settings/index.svelte';
 import type { ModelOption } from '$lib/types/models';
+import { remToPx } from '$lib/utils/css';
 
 class UiStore {
 	/** Set when a flow outside the chat wants the composer focused once it can take it. */
@@ -21,14 +32,19 @@ class UiStore {
 	/** Open state of the Discover dialog. */
 	discoverModelsOpen = $state(false);
 
-	/** Whether the desktop sidebar is expanded (open). */
-	isSidebarExpanded = $state(false);
+	/** Set while the user drags the sidebar edge, so width transitions can step aside. */
+	isSidebarResizing = $state(false);
+
 	/** Model the manager reveals when it opens, a qualified id or a raw model name. */
 	manageModelFocus = $state<string | null>(null);
 	/** Model the information dialog shows; a phone opens it instead of the manager. */
 	modelInformation = $state<ModelOption | null>(null);
 	/** Route the models page returns to: the last page visited before it. */
 	modelsReturnRoute: string = ROUTES.START;
+	/** Expanded desktop sidebar width in px; null keeps the stylesheet default. */
+	sidebarWidth = $state<number | null>(this.loadSidebarWidth());
+
+	#isSidebarExpanded = $state(this.loadSidebarExpanded());
 
 	/** Whether Discover is turned on; every Discover entry point hides without it. */
 	get isDiscoverEnabled(): boolean {
@@ -40,6 +56,20 @@ class UiStore {
 		return page.route.id === MODELS_ROUTE_ID;
 	}
 
+	/** Whether the desktop sidebar is expanded (open); remembered across reloads on desktop. */
+	get isSidebarExpanded(): boolean {
+		return this.#isSidebarExpanded;
+	}
+
+	set isSidebarExpanded(expanded: boolean) {
+		this.#isSidebarExpanded = expanded;
+
+		// a phone opens the sidebar as an overlay, which is no layout to restore
+		if (!browser || deviceStore.isMobile) return;
+
+		localStorage.setItem(SIDEBAR_EXPANDED_LOCALSTORAGE_KEY, String(expanded));
+	}
+
 	/** Leave the models page for the route it was opened from. */
 	async closeModelsManager(): Promise<void> {
 		if (!this.isModelsPageOpen) return;
@@ -80,6 +110,41 @@ class UiStore {
 	requestComposerFocus(): void {
 		this.composerFocusRequested = true;
 	}
+
+	/** Set the expanded sidebar width in px, or null to go back to the default. */
+	setSidebarWidth(width: number | null): void {
+		this.sidebarWidth = width;
+
+		if (!browser) return;
+
+		if (width === null) {
+			localStorage.removeItem(SIDEBAR_WIDTH_LOCALSTORAGE_KEY);
+		} else {
+			localStorage.setItem(SIDEBAR_WIDTH_LOCALSTORAGE_KEY, String(width));
+		}
+	}
+
+	private loadSidebarExpanded(): boolean {
+		if (!browser || deviceStore.isMobile) return false;
+
+		return localStorage.getItem(SIDEBAR_EXPANDED_LOCALSTORAGE_KEY) === 'true';
+	}
+
+	private loadSidebarWidth(): number | null {
+		if (!browser) return null;
+
+		const width = Number(localStorage.getItem(SIDEBAR_WIDTH_LOCALSTORAGE_KEY));
+
+		if (!Number.isFinite(width) || width <= 0) return null;
+
+		// a width saved under older bounds is pulled back inside the current ones
+		const clamped = Math.min(
+			remToPx(SIDEBAR_MAX_WIDTH),
+			Math.max(remToPx(SIDEBAR_MIN_WIDTH), width)
+		);
+
+		return Math.round(clamped);
+	}
 }
 
 export const uiStore = new UiStore();

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

@@ -0,0 +1,38 @@
+import { NEW_CHAT_TAB_ID, ROUTES } from '$lib/constants';
+import { RouterService } from '$lib/services/router.service';
+
+interface TabRoute {
+	params: Partial<Record<string, string>>;
+	route: { id: string | null };
+}
+
+/** Which tabs a route shows and where they lead. */
+export interface TabView {
+	/** Tab of the current route; null on a page without one. */
+	activeId: string | null;
+	/** Where to go when the last visible tab closes. */
+	fallbackHref: string;
+	hrefFor: (id: string) => string;
+	/** Visible tab ids, in strip order. */
+	ids: string[];
+}
+
+/**
+ * Tab target of the current route: a conversation id or the new-chat
+ * screen. Null on routes without tabs.
+ */
+export function resolveTabTarget(page: TabRoute): string | null {
+	if (page.params.id) return page.params.id;
+
+	return page.route.id === '/(chat)' ? NEW_CHAT_TAB_ID : null;
+}
+
+/** Tabs of the current route: every open tab, linked to its chat. */
+export function resolveTabView(page: TabRoute, openTabs: string[]): TabView {
+	return {
+		activeId: resolveTabTarget(page),
+		fallbackHref: ROUTES.START,
+		hrefFor: (id) => (id === NEW_CHAT_TAB_ID ? ROUTES.START : RouterService.chat(id)),
+		ids: openTabs
+	};
+}

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

@@ -215,6 +215,8 @@ export {
 	type PathQuery
 } from './working-directory';
 
+export { resolveTabTarget, resolveTabView, type TabView } from './chat-tabs';
+
 // Shared `file_glob_search` runner with a short-lived result cache
 export { runGlobSearch, runGlobSearchWithChildren } from './glob-search';
 

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

@@ -23,7 +23,9 @@
 	});
 </script>
 
-<div class={showTabs ? 'md:[--chat-tabs-offset:1.25rem]' : ''}>
+<!-- the offset is the tab bar's full height (pt-3.25 + h-10), so the chat
+     column still fits the viewport and only scrolls when the messages overflow -->
+<div class={showTabs ? 'md:[--chat-tabs-offset:3.3125rem]' : ''}>
 	{#if showTabs}
 		<ChatTabs />
 	{/if}

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

@@ -125,20 +125,6 @@
 		}
 	}
 
-	// navigating away from the new-chat screen drops its tab, so it does not
-	// linger once the user moves to a real conversation or another route
-	let previousChatId = $state<string | undefined>(undefined);
-
-	$effect(() => {
-		const id = page.params.id ?? (page.route.id === '/(chat)' ? NEW_CHAT_TAB_ID : undefined);
-		const prev = untrack(() => previousChatId);
-
-		previousChatId = id;
-
-		if (id !== prev && prev && settingsStore.config.conversationTabs && prev === NEW_CHAT_TAB_ID) {
-			untrack(() => tabsStore.removeTabs([NEW_CHAT_TAB_ID]));
-		}
-	});
 	// the models page returns to wherever it was reached from, by any kind of navigation
 	afterNavigate(() => uiStore.rememberModelsReturnRoute());
 

tools/ui/tests/client/chat-tabs-store.svelte.test.ts

@@ -0,0 +1,117 @@
+import { beforeEach, describe, expect, it, vi } from 'vitest';
+
+vi.mock('$app/navigation', () => ({ goto: vi.fn() }));
+
+import { goto } from '$app/navigation';
+import { CONVERSATION_TABS_LOCALSTORAGE_KEY, NEW_CHAT_TAB_ID } from '$lib/constants';
+import { TabsStore } from '$lib/stores/tabs.svelte';
+
+const gotoMock = vi.mocked(goto);
+
+function freshStore(persisted?: unknown, validIds: string[] = []) {
+	localStorage.removeItem(CONVERSATION_TABS_LOCALSTORAGE_KEY);
+
+	if (persisted !== undefined) {
+		localStorage.setItem(CONVERSATION_TABS_LOCALSTORAGE_KEY, JSON.stringify(persisted));
+	}
+
+	const tabs = new TabsStore();
+
+	tabs.init(validIds);
+
+	return tabs;
+}
+
+beforeEach(() => {
+	gotoMock.mockClear();
+});
+
+/**
+ * A plain navigation shows the new target in the current tab; only an
+ * explicit new-tab action grows the strip.
+ */
+describe('tabsStore route sync', () => {
+	it('creates the first tab, then navigates it in place', () => {
+		const tabs = freshStore();
+
+		tabs.syncWithRoute('a');
+		tabs.syncWithRoute('b');
+		tabs.syncWithRoute('c');
+
+		expect(tabs.openTabs).toEqual(['c']);
+	});
+
+	it('activates an open tab instead of replacing the current one', () => {
+		const tabs = freshStore(['a', 'b'], ['a', 'b']);
+
+		tabs.syncWithRoute('b');
+		tabs.syncWithRoute('a');
+
+		expect(tabs.openTabs).toEqual(['a', 'b']);
+	});
+
+	it('turns the new-chat tab into the conversation it creates', async () => {
+		const tabs = freshStore();
+
+		tabs.syncWithRoute('a');
+		await tabs.openInNewTab(NEW_CHAT_TAB_ID, { activate: true });
+		tabs.syncWithRoute(NEW_CHAT_TAB_ID);
+		tabs.syncWithRoute('b');
+
+		expect(tabs.openTabs).toEqual(['a', 'b']);
+	});
+
+	it('waits for init before applying the route', () => {
+		localStorage.setItem(CONVERSATION_TABS_LOCALSTORAGE_KEY, JSON.stringify(['a', 'b']));
+
+		const tabs = new TabsStore();
+
+		tabs.syncWithRoute('b');
+		tabs.init(['a', 'b']);
+
+		expect(tabs.openTabs).toEqual(['a', 'b']);
+	});
+});
+
+describe('tabsStore explicit tabs', () => {
+	it('opens new tabs at the end, in the background, without duplicates', async () => {
+		const tabs = freshStore();
+
+		tabs.syncWithRoute('a');
+		await tabs.openInNewTab('b');
+		await tabs.openInNewTab('c');
+		await tabs.openInNewTab('b');
+
+		expect(tabs.openTabs).toEqual(['a', 'b', 'c']);
+		expect(gotoMock).not.toHaveBeenCalled();
+	});
+
+	it('closes the active tab to its left neighbour', async () => {
+		const tabs = freshStore();
+
+		tabs.syncWithRoute('a');
+		await tabs.openInNewTab('b');
+		tabs.syncWithRoute('b');
+		await tabs.close('b', true);
+
+		expect(tabs.openTabs).toEqual(['a']);
+		expect(gotoMock).toHaveBeenLastCalledWith('#/chat/a');
+	});
+
+	it('falls back to the new-chat screen when the last tab closes', async () => {
+		const tabs = freshStore();
+
+		tabs.syncWithRoute('a');
+		await tabs.close('a', true);
+
+		expect(gotoMock).toHaveBeenLastCalledWith('#/');
+	});
+});
+
+describe('tabsStore persistence', () => {
+	it('prunes deleted conversations and keeps the new-chat tab', () => {
+		const tabs = freshStore(['a', NEW_CHAT_TAB_ID, 'gone'], ['a']);
+
+		expect(tabs.openTabs).toEqual(['a', NEW_CHAT_TAB_ID]);
+	});
+});