Merge remote-tracking branch 'origin/upstream' into concedo_experimental

# Conflicts:
#	examples/save-load-state/save-load-state.cpp
#	ggml/CMakeLists.txt
#	ggml/src/ggml-cpu/CMakeLists.txt
#	ggml/src/ggml-hexagon/htp/hmx-matmul-ops.c
#	ggml/src/ggml-hexagon/htp/matmul-ops.c
#	ggml/src/ggml-opencl/CMakeLists.txt
#	ggml/src/ggml-opencl/ggml-opencl.cpp
#	ggml/src/ggml-opencl/kernels/gemm_noshuffle_q4_0_f32.cl
#	ggml/src/ggml-opencl/kernels/gemm_noshuffle_q8_0_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32_spec.cl
#	ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl
#	ggml/src/ggml-rpc/ggml-rpc.cpp
#	scripts/sync-ggml.last
#	scripts/sync_vendor.py
#	src/llama-graph.cpp
#	tests/test-backend-ops.cpp
#	tests/test-state-restore-fragmented.cpp

common/arg.cpp

@@ -250,6 +250,8 @@ std::vector<std::string> common_arg::get_env() const {
 
 // Helper function to parse tensor buffer override strings
 static void parse_tensor_buffer_overrides(const std::string & value, std::vector<llama_model_tensor_buft_override> & overrides) {
+    ggml_backend_load_all();
+
     std::map<std::string, ggml_backend_buffer_type_t> buft_list;
     for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
         auto * dev = ggml_backend_dev_get(i);
@@ -427,6 +429,10 @@ static bool parse_bool_value(const std::string & value) {
     }
 }
 
+[[noreturn]] static void arg_removed(const std::string & msg) {
+    throw std::invalid_argument("the argument has been removed. " + msg);
+}
+
 //
 // CLI argument parsing functions
 //
@@ -805,6 +811,7 @@ static std::vector<ggml_backend_dev_t> parse_device_list(const std::string & val
     if (dev_names.size() == 1 && dev_names[0] == "none") {
         devices.push_back(nullptr);
     } else {
+        ggml_backend_load_all();
         for (const auto & device : dev_names) {
             auto * dev = ggml_backend_dev_by_name(device.c_str());
             if (!dev || ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_CPU) {
@@ -822,6 +829,7 @@ static void add_rpc_devices(const std::string & servers) {
     if (rpc_servers.empty()) {
         throw std::invalid_argument("no RPC servers specified");
     }
+    ggml_backend_load_all();
     ggml_backend_reg_t rpc_reg = ggml_backend_reg_by_name("RPC");
     if (!rpc_reg) {
         throw std::invalid_argument("failed to find RPC backend");
@@ -1018,9 +1026,6 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
 
     params.use_color = tty_can_use_colors();
 
-    // load dynamic backends
-    ggml_backend_load_all();
-
     common_params_context ctx_arg(params);
     ctx_arg.print_usage = print_usage;
     ctx_arg.ex          = ex;
@@ -2277,6 +2282,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
         {"--list-devices"},
         "print list of available devices and exit",
         [](common_params &) {
+            ggml_backend_load_all();
             std::vector<ggml_backend_dev_t> devices;
             for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
                 auto * dev = ggml_backend_dev_get(i);
@@ -3717,35 +3723,35 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
         {"--draft", "--draft-n", "--draft-max"}, "N",
         "the argument has been removed. use --spec-draft-n-max or --spec-ngram-mod-n-max",
         [](common_params & /*params*/, int /*value*/) {
-            throw std::invalid_argument("the argument has been removed. use --spec-draft-n-max or --spec-ngram-mod-n-max");
+            arg_removed("use --spec-draft-n-max or --spec-ngram-mod-n-max");
         }
     ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_LOOKUP, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_DRAFT_MAX"));
     add_opt(common_arg(
         {"--draft-min", "--draft-n-min"}, "N",
         "the argument has been removed. use --spec-draft-n-min or --spec-ngram-mod-n-min",
         [](common_params & /*params*/, int /*value*/) {
-            throw std::invalid_argument("the argument has been removed. use --spec-draft-n-min or --spec-ngram-mod-n-min");
+            arg_removed("use --spec-draft-n-min or --spec-ngram-mod-n-min");
         }
     ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_LOOKUP, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_DRAFT_MIN"));
     add_opt(common_arg(
         {"--spec-ngram-size-n"}, "N",
         "the argument has been removed. use the respective --spec-ngram-*-size-n or --spec-ngram-mod-n-match",
         [](common_params & /*params*/, int /*value*/) {
-            throw std::invalid_argument("the argument has been removed. use the respective --spec-ngram-*-size-n");
+            arg_removed("use the respective --spec-ngram-*-size-n");
         }
     ).set_spec().set_examples({LLAMA_EXAMPLE_SERVER}));
     add_opt(common_arg(
         {"--spec-ngram-size-m"}, "N",
         "the argument has been removed. use the respective --spec-ngram-*-size-m",
         [](common_params & /*params*/, int /*value*/) {
-            throw std::invalid_argument("the argument has been removed. use the respective --spec-ngram-*-size-m");
+            arg_removed("use the respective --spec-ngram-*-size-m");
         }
     ).set_spec().set_examples({LLAMA_EXAMPLE_SERVER}));
     add_opt(common_arg(
         {"--spec-ngram-min-hits"}, "N",
         "the argument has been removed. use the respective --spec-ngram-*-min-hits",
         [](common_params & /*params*/, int /*value*/) {
-            throw std::invalid_argument("the argument has been removed. use the respective --spec-ngram-*-min-hits");
+            arg_removed("use the respective --spec-ngram-*-min-hits");
         }
     ).set_spec().set_examples({LLAMA_EXAMPLE_SERVER}));
 

common/speculative.cpp

@@ -252,14 +252,14 @@ struct common_speculative_state_draft : public common_speculative_state {
 
     size_t create_checkpoint(int n_tokens_prompt) {
         int slot_id = 0;
-        const size_t checkpoint_size = llama_state_seq_get_size_ext(ctx_dft, slot_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+        const size_t checkpoint_size = llama_state_seq_get_size_ext(ctx_dft, slot_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE);
 
         ckpt.pos_min  = llama_memory_seq_pos_min(llama_get_memory(ctx_dft), slot_id);
         ckpt.pos_max  = llama_memory_seq_pos_max(llama_get_memory(ctx_dft), slot_id);
         ckpt.n_tokens = n_tokens_prompt;
         ckpt.data.resize(checkpoint_size);
 
-        const size_t n = llama_state_seq_get_data_ext(ctx_dft, ckpt.data.data(), checkpoint_size, slot_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+        const size_t n = llama_state_seq_get_data_ext(ctx_dft, ckpt.data.data(), checkpoint_size, slot_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE);
         if (n != checkpoint_size) {
             GGML_ABORT("checkpoint size mismatch: expected %zu, got %zu\n", checkpoint_size, n);
         }
@@ -272,7 +272,7 @@ struct common_speculative_state_draft : public common_speculative_state {
     size_t restore_checkpoint() {
         int slot_id = 0;
         LOG_DBG("%s: pos_min = %d, pos_max = %d\n", __func__, ckpt.pos_min, ckpt.pos_max);
-        const size_t n = llama_state_seq_set_data_ext(ctx_dft, ckpt.data.data(), ckpt.size(), slot_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+        const size_t n = llama_state_seq_set_data_ext(ctx_dft, ckpt.data.data(), ckpt.size(), slot_id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE);
         if (n != ckpt.size()) {
             GGML_ABORT("%s: failed to restore context checkpoint (pos_min=%d, pos_max=%d, size=%zu",
                         __func__, ckpt.pos_min, ckpt.pos_max, ckpt.size());

ggml/include/ggml.h

@@ -444,6 +444,12 @@ extern "C" {
         GGML_PREC_F32     = 10,
     };
 
+    // op hint
+    enum ggml_op_hint {
+        GGML_HINT_NONE             = 0,
+        GGML_HINT_SRC0_IS_HADAMARD = 1,
+    };
+
     // model file types
     enum ggml_ftype {
         GGML_FTYPE_UNKNOWN        = -1,
@@ -1437,6 +1443,11 @@ extern "C" {
             struct ggml_tensor * a,
             enum ggml_prec       prec);
 
+    // change the hint of a matrix multiplication
+    GGML_API void ggml_mul_mat_set_hint(
+            struct ggml_tensor * a,
+            enum ggml_op_hint    hint);
+
     // indirect matrix multiplication
     GGML_API struct ggml_tensor * ggml_mul_mat_id(
             struct ggml_context * ctx,

ggml/src/ggml-cpu/ggml-cpu.c

@@ -1247,6 +1247,12 @@ void ggml_compute_forward_mul_mat(
     const struct ggml_tensor * src0 = dst->src[0];
     const struct ggml_tensor * src1 = dst->src[1];
 
+    const int32_t hint = ggml_get_op_params_i32(dst, 1);
+    if (hint == GGML_HINT_SRC0_IS_HADAMARD && !params->use_ref) {
+        ggml_compute_forward_fwht(params, dst);
+        return;
+    }
+
     GGML_TENSOR_BINARY_OP_LOCALS
 
     const int ith = params->ith;
@@ -3808,6 +3814,45 @@ struct ggml_cplan ggml_graph_plan(
     return cplan;
 }
 
+
+// Try to fuse the current node with subsequent nodes for better performance.
+// Returns the number of nodes skipped by fusion (>=1), or 0 if no fusion was applied.
+static bool ggml_cpu_disable_fusion = false;  // initialized once in ggml_cpu_init(), read-only afterwards
+
+static int ggml_cpu_try_fuse_ops(
+        const struct ggml_cgraph * cgraph,
+        const int node_n,
+        const struct ggml_compute_params * params,
+        const struct ggml_cplan * cplan) {
+
+    if (ggml_cpu_disable_fusion || cplan->use_ref) {
+        return 0;
+    }
+
+    struct ggml_tensor * node = cgraph->nodes[node_n];
+
+    if (node->op == GGML_OP_RMS_NORM) {
+        // RMS_NORM + MUL fusion
+        const enum ggml_op fuse_ops[] = { GGML_OP_RMS_NORM, GGML_OP_MUL };
+        if (ggml_can_fuse(cgraph, node_n, fuse_ops, 2)) {
+            struct ggml_tensor * mul_node = cgraph->nodes[node_n + 1];
+            const struct ggml_tensor * mul_w = (mul_node->src[0] == node)
+                ? mul_node->src[1] : mul_node->src[0];
+            if (node->src[0]->type  == GGML_TYPE_F32 &&
+                mul_node->type      == GGML_TYPE_F32 &&
+                mul_w->type         == GGML_TYPE_F32 &&
+                mul_w->ne[0]        == node->ne[0]   &&
+                mul_w->nb[0]        == sizeof(float)) {
+
+                ggml_compute_forward_rms_norm_mul_fused(params, node, mul_node);
+                return 1;
+            }
+        }
+    }
+
+    return 0;
+}
+
 static thread_ret_t ggml_graph_compute_thread(void * data) {
     struct ggml_compute_state * state = (struct ggml_compute_state *) data;
     struct ggml_threadpool    * tp    = state->threadpool;
@@ -3844,7 +3889,14 @@ static thread_ret_t ggml_graph_compute_thread(void * data) {
             continue;
         }
 
-        ggml_compute_forward(&params, node);
+        // TODO: move fused-op detection into ggml_graph_plan so fusion decisions are made once at planning time
+        // Try fused ops, fall back to normal compute
+        const int n_fused = ggml_cpu_try_fuse_ops(cgraph, node_n, &params, cplan);
+        if (n_fused > 0) {
+            node_n += n_fused;
+        } else {
+            ggml_compute_forward(&params, node);
+        }
 
         if (state->ith == 0 && cplan->abort_callback &&
                 cplan->abort_callback(cplan->abort_callback_data)) {
@@ -4606,6 +4658,11 @@ void ggml_cpu_init(void) {
         ggml_init_riscv_arch_features();
 #endif
 
+        {
+            const char * env = getenv("GGML_CPU_DISABLE_FUSION");
+            ggml_cpu_disable_fusion = (env != NULL && atoi(env) == 1);
+        }
+
         is_first_call = false;
     }
 

ggml/src/ggml-cpu/ops.cpp

@@ -3713,11 +3713,27 @@ void ggml_compute_forward_norm(
 
 // ggml_compute_forward_group_rms_norm
 
+// fusion kinds that can be combined with the rms_norm computation in a single pass.
+// extend this enum when adding new fused variants (e.g. FUSE_ADD, FUSE_MUL_ADD, ...).
+enum ggml_rms_norm_fuse_op {
+    GGML_RMS_NORM_FUSE_OP_NONE,
+    GGML_RMS_NORM_FUSE_OP_MUL,
+};
+
+template <ggml_rms_norm_fuse_op FUSE_OP>
 static void ggml_compute_forward_rms_norm_f32(
         const ggml_compute_params * params,
-        ggml_tensor * dst) {
+        ggml_tensor * dst_rms_norm,
+        ggml_tensor * dst_fused = nullptr) {
 
-    const ggml_tensor * src0 = dst->src[0];
+    const ggml_tensor * src0 = dst_rms_norm->src[0];
+    const ggml_tensor * src1 = nullptr;
+    ggml_tensor       * dst  = dst_rms_norm;
+
+    if constexpr (FUSE_OP == GGML_RMS_NORM_FUSE_OP_MUL) {
+        src1 = (dst_fused->src[0] == dst_rms_norm) ? dst_fused->src[1] : dst_fused->src[0];
+        dst  = dst_fused;
+    }
 
     GGML_ASSERT(ggml_are_same_shape(src0, dst));
 
@@ -3726,11 +3742,10 @@ static void ggml_compute_forward_rms_norm_f32(
     const int ith = params->ith;
     const int nth = params->nth;
 
-    GGML_TENSOR_UNARY_OP_LOCALS
+    GGML_TENSOR_BINARY_OP_LOCALS
 
     float eps;
-    memcpy(&eps, dst->op_params, sizeof(float));
-
+    memcpy(&eps, dst_rms_norm->op_params, sizeof(float));
     GGML_ASSERT(eps >= 0.0f);
 
     // TODO: optimize
@@ -3740,25 +3755,32 @@ static void ggml_compute_forward_rms_norm_f32(
                 const float * x = (float *) ((char *) src0->data + i01*nb01 + i02*nb02 + i03*nb03);
 
                 ggml_float sum = 0.0;
+                // worth switching to explicit SIMD?
                 for (int64_t i00 = 0; i00 < ne00; i00++) {
                     sum += (ggml_float)(x[i00] * x[i00]);
                 }
 
-                const float mean = sum/ne00;
-
-                float * y = (float *) ((char *) dst->data + i01*nb1 + i02*nb2 + i03*nb3);
-
-                memcpy(y, x, ne00 * sizeof(float));
-                // for (int i00 = 0; i00 < ne00; i00++) {
-                //     y[i00] = x[i00];
-                // }
-
+                const float mean  = sum/ne00;
                 const float scale = 1.0f/sqrtf(mean + eps);
 
                 // if you hit this, likely you got an inf somewhere earlier
                 assert(scale > 0.0f);
 
-                ggml_vec_scale_f32(ne00, y, scale);
+                float * y = (float *) ((char *) dst->data + i01*nb1 + i02*nb2 + i03*nb3);
+
+                if constexpr (FUSE_OP == GGML_RMS_NORM_FUSE_OP_MUL) {
+                    const int64_t i11 = i01 % ne11;
+                    const int64_t i12 = i02 % ne12;
+                    const int64_t i13 = i03 % ne13;
+                    const float * w = (float *) ((char *) src1->data + i11*nb11 + i12*nb12 + i13*nb13);
+
+                    for (int64_t i00 = 0; i00 < ne00; i00++) {
+                        y[i00] = x[i00] * scale * w[i00];
+                    }
+                } else {
+                    memcpy(y, x, ne00 * sizeof(float));
+                    ggml_vec_scale_f32(ne00, y, scale);
+                }
             }
         }
     }
@@ -3773,7 +3795,31 @@ void ggml_compute_forward_rms_norm(
     switch (src0->type) {
         case GGML_TYPE_F32:
             {
-                ggml_compute_forward_rms_norm_f32(params, dst);
+                ggml_compute_forward_rms_norm_f32<GGML_RMS_NORM_FUSE_OP_NONE>(params, dst);
+            } break;
+        default:
+            {
+                GGML_ABORT("fatal error");
+            }
+    }
+}
+
+// Fused RMS_NORM + MUL: computes dst = rms_norm(src0) * src1 in a single pass.
+// This avoids materializing the intermediate rms_norm result in memory.
+void ggml_compute_forward_rms_norm_mul_fused(
+        const ggml_compute_params * params,
+        ggml_tensor * dst_rms_norm,
+        ggml_tensor * dst_mul) {
+
+    GGML_ASSERT(dst_mul != nullptr);
+    GGML_ASSERT(dst_mul->src[0] == dst_rms_norm || dst_mul->src[1] == dst_rms_norm);
+
+    const ggml_tensor * src0 = dst_rms_norm->src[0];
+
+    switch (src0->type) {
+        case GGML_TYPE_F32:
+            {
+                ggml_compute_forward_rms_norm_f32<GGML_RMS_NORM_FUSE_OP_MUL>(params, dst_rms_norm, dst_mul);
             } break;
         default:
             {
@@ -11212,3 +11258,91 @@ void ggml_compute_forward_opt_step_sgd(const ggml_compute_params * params, ggml_
             }
     }
 }
+
+static void ggml_compute_forward_fwht_f32(const ggml_compute_params * params, ggml_tensor * dst) {
+    const ggml_tensor * src0 = dst->src[0];
+    const ggml_tensor * src1 = dst->src[1];
+
+    GGML_ASSERT(src1->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+
+    GGML_TENSOR_BINARY_OP_LOCALS
+
+    const int ith = params->ith;
+    const int nth = params->nth;
+
+    const int64_t n = ne10;
+    GGML_ASSERT((n & (n - 1)) == 0); // must be power of 2
+
+    const int64_t nr = ne11 * ne12 * ne13;
+    const int64_t rows_per_thread = (nr + nth - 1) / nth;
+    const int64_t start_row = ith * rows_per_thread;
+    const int64_t end_row = MIN(start_row + rows_per_thread, nr);
+
+    const float scale = 1.0f / sqrtf((float)n);
+
+#if defined(GGML_SIMD)
+    const GGML_F32_VEC v_minus_one = GGML_F32_VEC_SET1(-1.0f);
+#endif
+
+    for (int64_t r = start_row; r < end_row; r++) {
+        const int64_t i13 = r / (ne11 * ne12);
+        const int64_t i12 = (r - i13 * ne11 * ne12) / ne11;
+        const int64_t i11 = r - i13 * ne11 * ne12 - i12 * ne11;
+
+        const float * src_row = (const float *) ((const char *) src1->data + i11 * nb11 + i12 * nb12 + i13 * nb13);
+        float * dst_row = (float *) ((char *) dst->data + i11 * nb1 + i12 * nb2 + i13 * nb3);
+
+        for (int64_t j = 0; j < n; j++) {
+            dst_row[j] = src_row[j] * scale;
+        }
+
+        // Scalar passes
+#if defined(GGML_SIMD)
+        const int step = GGML_F32_EPR;
+#else
+        const int step = n;
+#endif
+        for (int64_t len = 1; len < step && len < n; len <<= 1) {
+            for (int64_t i = 0; i < n; i += 2 * len) {
+                for (int64_t j = 0; j < len; j++) {
+                    float u = dst_row[i + j];
+                    float v = dst_row[i + len + j];
+                    dst_row[i + j] = u + v;
+                    dst_row[i + len + j] = u - v;
+                }
+            }
+        }
+
+        // SIMD passes using GGML_F32_VEC_* macros for multi-architecture support
+#if defined(GGML_SIMD)
+        for (int64_t len = step; len < n; len <<= 1) {
+            for (int64_t i = 0; i < n; i += 2 * len) {
+                for (int64_t j = 0; j < len; j += step) {
+                    GGML_F32_VEC u = GGML_F32_VEC_LOAD(dst_row + i + j);
+                    GGML_F32_VEC v = GGML_F32_VEC_LOAD(dst_row + i + len + j);
+
+                    GGML_F32_VEC_STORE(dst_row + i + j,       GGML_F32_VEC_ADD(u, v));
+                    GGML_F32_VEC_STORE(dst_row + i + len + j, GGML_F32_VEC_FMA(u, v, v_minus_one));
+                }
+            }
+        }
+#endif
+    }
+}
+
+void ggml_compute_forward_fwht(const ggml_compute_params * params, ggml_tensor * dst) {
+    const ggml_tensor * src1 = dst->src[1];
+
+    switch (src1->type) {
+        case GGML_TYPE_F32:
+            {
+                ggml_compute_forward_fwht_f32(params, dst);
+            }
+            break;
+        default:
+            {
+                GGML_ABORT("fatal error - fwht is F32 only");
+            }
+    }
+}

ggml/src/ggml-cpu/ops.h

@@ -44,6 +44,7 @@ void ggml_compute_forward_concat(const struct ggml_compute_params * params, stru
 void ggml_compute_forward_silu_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_norm(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_rms_norm(const struct ggml_compute_params * params, struct ggml_tensor * dst);
+void ggml_compute_forward_rms_norm_mul_fused(const struct ggml_compute_params * params, struct ggml_tensor * dst_rms_norm, struct ggml_tensor * dst_mul);
 void ggml_compute_forward_rms_norm_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_group_norm(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_l2_norm(const struct ggml_compute_params * params, struct ggml_tensor * dst);
@@ -111,6 +112,7 @@ void ggml_compute_forward_cross_entropy_loss(const struct ggml_compute_params *
 void ggml_compute_forward_cross_entropy_loss_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_opt_step_adamw(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_mul_mat(const struct ggml_compute_params * params, struct ggml_tensor * dst);
+void ggml_compute_forward_fwht(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_opt_step_sgd(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 #ifdef __cplusplus
 }

ggml/src/ggml-cuda/getrows.cu

@@ -6,17 +6,18 @@ template<int qk, int qr, dequantize_kernel_t dequantize_kernel, typename dst_t>
 static __global__ void k_get_rows(
         const void * __restrict__ src0, const int32_t * __restrict__ src1, dst_t * __restrict__ dst,
         const int64_t ne00, /*const int64_t ne01, const int64_t ne02, const int64_t ne03,*/
-        /*const int64_t ne10,*/ const int64_t ne11, const int64_t ne12, /*const int64_t ne13,*/
+        /*const int64_t ne10,*/ const int64_t ne11, const uint3 ne12_fdv, /*const int64_t ne13,*/
         /*const size_t s0,*/ const size_t s1, const size_t s2, const size_t s3,
         /*const size_t nb00,*/ const size_t nb01, const size_t nb02, const size_t nb03,
         const size_t s10, const size_t s11, const size_t s12/*, const size_t s13*/) {
 
-    for (int64_t z = blockIdx.z; z < ne11*ne12; z += gridDim.z) {
+    for (int64_t z = blockIdx.z; z < ne11*(int64_t)ne12_fdv.z; z += gridDim.z) {
         for (int64_t i00 = 2*(blockIdx.y*blockDim.x + threadIdx.x); i00 < ne00; i00 += gridDim.y*blockDim.x) {
             // The x and y dimensions of the grid are swapped because the maximum allowed grid size for x is higher.
             const int i10 =  blockIdx.x;
-            const int i11 =  z / ne12; // TODO fastdiv
-            const int i12 =  z % ne12;
+            const uint2 dm  = fast_div_modulo((uint32_t)z, ne12_fdv);
+            const int i11 =  dm.x;
+            const int i12 =  dm.y;
 
             const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
 
@@ -42,17 +43,18 @@ template<typename src0_t, typename dst_t>
 static __global__ void k_get_rows_float(
         const src0_t * __restrict__ src0, const int32_t * __restrict__ src1, dst_t * __restrict__ dst,
         const int64_t ne00, /*const int64_t ne01, const int64_t ne02, const int64_t ne03,*/
-        /*const int64_t ne10,*/ const int64_t ne11, const int64_t ne12, /*const int64_t ne13,*/
+        /*const int64_t ne10,*/ const int64_t ne11, const uint3 ne12_fdv, /*const int64_t ne13,*/
         /*const size_t s0,*/ const size_t s1, const size_t s2, const size_t s3,
         /*const size_t nb00,*/ const size_t nb01, const size_t nb02, const size_t nb03,
         const size_t s10, const size_t s11, const size_t s12/*, const size_t s13*/) {
 
-    for (int64_t z = blockIdx.z; z < ne11*ne12; z += gridDim.z) {
+    for (int64_t z = blockIdx.z; z < ne11*(int64_t)ne12_fdv.z; z += gridDim.z) {
         for (int64_t i00 = blockIdx.y*blockDim.x + threadIdx.x; i00 < ne00; i00 += gridDim.y*blockDim.x) {
             // The x and y dimensions of the grid are swapped because the maximum allowed grid size for x is higher.
             const int i10 = blockIdx.x;
-            const int i11 = z / ne12; // TODO fastdiv
-            const int i12 = z % ne12;
+            const uint2 dm = fast_div_modulo((uint32_t)z, ne12_fdv);
+            const int i11 = dm.x;
+            const int i12 = dm.y;
 
             if (i00 >= ne00) {
                 return;
@@ -115,10 +117,14 @@ static void get_rows_cuda_q(
 
     GGML_ASSERT(ne00 % 2 == 0);
 
+    GGML_ASSERT(ne12 > 0);
+    GGML_ASSERT(ne11 <= std::numeric_limits<uint32_t>::max() / ne12);
+    const uint3 ne12_fdv = init_fastdiv_values(ne12);
+
     k_get_rows<qk, qr, dq><<<block_nums, block_dims, 0, stream>>>(
         src0_d, src1_d, dst_d,
         ne00, /*ne01, ne02, ne03,*/
-        /*ne10,*/ ne11, ne12, /*ne13,*/
+        /*ne10,*/ ne11, ne12_fdv, /*ne13,*/
         /* s0,*/ s1, s2, s3,
         /* nb00,*/ nb01, nb02, nb03,
         s10, s11, s12/*, s13*/);
@@ -146,10 +152,14 @@ static void get_rows_cuda_float(
     const size_t s12 = nb12 / sizeof(int32_t);
     // const size_t s13 = nb13 / sizeof(int32_t);
 
+    GGML_ASSERT(ne12 > 0);
+    GGML_ASSERT(ne11 <= std::numeric_limits<uint32_t>::max() / ne12);
+    const uint3 ne12_fdv = init_fastdiv_values(ne12);
+
     k_get_rows_float<<<block_nums, block_dims, 0, stream>>>(
         src0_d, src1_d, dst_d,
         ne00, /*ne01, ne02, ne03,*/
-        /*ne10,*/ ne11, ne12, /*ne13,*/
+        /*ne10,*/ ne11, ne12_fdv, /*ne13,*/
         /* s0,*/ s1, s2, s3,
         /* nb00,*/ nb01, nb02, nb03,
         s10, s11, s12/*, s13*/);

ggml/src/ggml-metal/ggml-metal-device.h

@@ -282,6 +282,7 @@ bool   ggml_metal_buffer_is_shared(ggml_metal_buffer_t buf);
 void   ggml_metal_buffer_memset_tensor(ggml_metal_buffer_t buf, struct ggml_tensor * tensor, uint8_t value, size_t offset, size_t size);
 void   ggml_metal_buffer_set_tensor   (ggml_metal_buffer_t buf, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size);
 void   ggml_metal_buffer_get_tensor   (ggml_metal_buffer_t buf, const struct ggml_tensor * tensor, void * data, size_t offset, size_t size);
+bool   ggml_metal_buffer_cpy_tensor   (ggml_metal_buffer_t buf, const struct ggml_tensor * src, struct ggml_tensor * dst);
 void   ggml_metal_buffer_clear        (ggml_metal_buffer_t buf, uint8_t value);
 
 // finds the Metal buffer that contains the tensor data on the GPU device

ggml/src/ggml-metal/ggml-metal-device.m

@@ -1,6 +1,7 @@
 #import "ggml-metal-device.h"
 
 #import "ggml-impl.h"
+#import "ggml-backend-impl.h"
 
 #include <Foundation/Foundation.h>
 
@@ -1743,6 +1744,47 @@ void ggml_metal_buffer_get_tensor(ggml_metal_buffer_t buf, const struct ggml_ten
     }
 }
 
+bool ggml_metal_buffer_cpy_tensor(ggml_metal_buffer_t buf_dst, const struct ggml_tensor * src, struct ggml_tensor * dst) {
+    ggml_metal_buffer_t buf_src = (ggml_metal_buffer_t)src->buffer->context;
+
+    const size_t size = ggml_nbytes(src);
+
+    // if both buffers are shared, we can use memcpy directly
+    if (buf_dst->is_shared && buf_src->is_shared) {
+        memcpy(dst->data, src->data, size);
+        return true;
+    }
+
+    // for private buffers, we need to use Metal blit commands
+    @autoreleasepool {
+        struct ggml_metal_buffer_id bid_src = ggml_metal_buffer_get_id(buf_src, src);
+        struct ggml_metal_buffer_id bid_dst = ggml_metal_buffer_get_id(buf_dst, dst);
+
+        if (bid_src.metal == nil || bid_dst.metal == nil) {
+            return false;
+        }
+
+        id<MTLCommandBuffer> cmd_buf = [buf_dst->dev->mtl_queue commandBufferWithUnretainedReferences];
+
+        {
+            id<MTLBlitCommandEncoder> encoder = [cmd_buf blitCommandEncoder];
+
+            [encoder copyFromBuffer:bid_src.metal
+                       sourceOffset:bid_src.offs
+                           toBuffer:bid_dst.metal
+                  destinationOffset:bid_dst.offs
+                               size:size];
+
+            [encoder endEncoding];
+        }
+
+        [cmd_buf commit];
+        [cmd_buf waitUntilCompleted];
+    }
+
+    return true;
+}
+
 void ggml_metal_buffer_clear(ggml_metal_buffer_t buf, uint8_t value) {
     if (buf->is_shared) {
         memset(buf->all_data, value, buf->all_size);

ggml/src/ggml-metal/ggml-metal.cpp

@@ -17,6 +17,9 @@
 // note: can be overridden with GGML_METAL_DEVICES env to simulate virtual devices
 static int g_devices = 1;
 
+// forward declaration
+static bool ggml_backend_buffer_is_metal(ggml_backend_buffer_t buffer);
+
 ////////////////////////////////////////////////////////////////////////////////
 // backend interface
 ////////////////////////////////////////////////////////////////////////////////
@@ -68,11 +71,11 @@ static bool ggml_backend_metal_buffer_shared_cpy_tensor(ggml_backend_buffer_t bu
 
     GGML_ASSERT(ggml_metal_buffer_is_shared(ctx));
 
-    GGML_UNUSED(buffer);
-    GGML_UNUSED(src);
-    GGML_UNUSED(dst);
+    if (!ggml_backend_buffer_is_metal(src->buffer)) {
+        return false;
+    }
 
-    return false;
+    return ggml_metal_buffer_cpy_tensor(ctx, src, dst);
 }
 
 static void ggml_backend_metal_buffer_shared_clear(ggml_backend_buffer_t buffer, uint8_t value) {
@@ -144,11 +147,11 @@ static bool ggml_backend_metal_buffer_private_cpy_tensor(ggml_backend_buffer_t b
 
     GGML_ASSERT(!ggml_metal_buffer_is_shared(ctx));
 
-    GGML_UNUSED(buffer);
-    GGML_UNUSED(src);
-    GGML_UNUSED(dst);
+    if (!ggml_backend_buffer_is_metal(src->buffer)) {
+        return false;
+    }
 
-    return false;
+    return ggml_metal_buffer_cpy_tensor(ctx, src, dst);
 }
 
 static void ggml_backend_metal_buffer_private_clear(ggml_backend_buffer_t buffer, uint8_t value) {

ggml/src/ggml.c

@@ -3280,6 +3280,16 @@ void ggml_mul_mat_set_prec(
     ggml_set_op_params_i32(a, 0, prec_i32);
 }
 
+void ggml_mul_mat_set_hint(
+        struct ggml_tensor * a,
+        enum ggml_op_hint    hint) {
+    GGML_ASSERT(a->op == GGML_OP_MUL_MAT);
+
+    const int32_t hint_i32 = (int32_t) hint;
+
+    ggml_set_op_params_i32(a, 1, hint_i32);
+}
+
 // ggml_mul_mat_id
 
 /*

gguf-py/gguf/tensor_mapping.py

@@ -18,7 +18,6 @@ class TensorNameMap:
             "tok_embeddings",                            # llama-pth
             "embeddings.word_embeddings",                # bert nomic-bert
             "embeddings.tok_embeddings",                 # modern-bert
-            "language_model.embedding.word_embeddings",  # persimmon
             "wte",                                       # gpt2
             "transformer.embd.wte",                      # phi2
             "model.tok_embeddings",                      # internlm2
@@ -32,7 +31,6 @@ class TensorNameMap:
             "rwkv.embeddings",                           # rwkv6
             "model.embeddings",                          # rwkv7
             "model.word_embeddings",                     # bailingmoe
-            "language_model.model.embed_tokens",         # llama4
             "encoder",                                   # neobert
             "model.transformer.wte",                     # llada
             "embed_tokens",                              # qwen3-embedding
@@ -94,7 +92,6 @@ class TensorNameMap:
             "norm",                                    # llama-pth
             "transformer.norm_f",                      # mpt dbrx
             "ln_f",                                    # refact bloom qwen gpt2
-            "language_model.encoder.final_layernorm",  # persimmon
             "model.final_layernorm",                   # persimmon
             "lm_head.ln",                              # phi2
             "model.norm_f",                            # mamba-qbert
@@ -171,7 +168,6 @@ class TensorNameMap:
             "transformer.h.{bid}.ln_mlp",                           # falcon40b
             "model.layers.{bid}.input_layernorm",                   # llama-hf nemotron olmoe phimoe granite-hybrid
             "layers.{bid}.attention_norm",                          # llama-pth
-            "language_model.encoder.layers.{bid}.input_layernorm",  # persimmon
             "model.layers.{bid}.ln1",                               # yi
             "h.{bid}.ln_1",                                         # gpt2
             "transformer.h.{bid}.ln",                               # phi2
@@ -215,7 +211,6 @@ class TensorNameMap:
             "transformer.blocks.{bid}.norm_attn_norm.attn.Wqkv",                   # dbrx
             "transformer.h.{bid}.self_attention.query_key_value",                  # falcon
             "h.{bid}.self_attention.query_key_value",                              # bloom
-            "language_model.encoder.layers.{bid}.self_attention.query_key_value",  # persimmon
             "model.layers.{bid}.self_attn.query_key_value",                        # persimmon
             "model.layers.{bid}.attention.query_key_value",                        # bailingmoe2
             "h.{bid}.attn.c_attn",                                                 # gpt2
@@ -306,7 +301,6 @@ class TensorNameMap:
             "layers.{bid}.attn.Wo",                                         # modern-bert
             "transformer.layer.{bid}.attention.out_lin",                    # distillbert
             "transformer.h.{bid}.attn.out_proj",                            # gpt-j
-            "language_model.encoder.layers.{bid}.self_attention.dense",     # persimmon
             "model.layers.{bid}.self_attn.dense",                           # persimmon
             "model.layers.{bid}.attention.dense",                           # bailingmoe2
             "h.{bid}.attn.c_proj",                                          # gpt2
@@ -373,7 +367,6 @@ class TensorNameMap:
             "transformer.blocks.{bid}.norm_2",                               # mpt
             "model.layers.{bid}.post_attention_layernorm",                   # llama-hf nemotron olmoe phimoe
             "layers.{bid}.ffn_norm",                                         # llama-pth
-            "language_model.encoder.layers.{bid}.post_attention_layernorm",  # persimmon
             "model.layers.{bid}.ln2",                                        # yi
             "h.{bid}.ln_2",                                                  # gpt2
             "model.layers.{bid}.ffn_norm",                                   # internlm2
@@ -475,7 +468,6 @@ class TensorNameMap:
             "transformer.layer.{bid}.ffn.lin1",                       # distillbert
             "transformer.h.{bid}.mlp.fc_in",                          # gpt-j
             "transformer.h.{bid}.mlp.linear_3",                       # refact
-            "language_model.encoder.layers.{bid}.mlp.dense_h_to_4h",  # persimmon
             "model.layers.{bid}.mlp.dense_h_to_4h",                   # persimmon
             "transformer.h.{bid}.mlp.w1",                             # qwen
             "h.{bid}.mlp.c_fc",                                       # gpt2
@@ -608,7 +600,6 @@ class TensorNameMap:
             "layers.{bid}.mlp.Wo",                                    # modern-bert
             "transformer.layer.{bid}.ffn.lin2",                       # distillbert
             "transformer.h.{bid}.mlp.fc_out",                         # gpt-j
-            "language_model.encoder.layers.{bid}.mlp.dense_4h_to_h",  # persimmon
             "model.layers.{bid}.mlp.dense_4h_to_h",                   # persimmon
             "h.{bid}.mlp.c_proj",                                     # gpt2
             "transformer.h.{bid}.mlp.fc2",                            # phi2
@@ -663,7 +654,7 @@ class TensorNameMap:
         ),
 
         MODEL_TENSOR.ATTN_Q_NORM: (
-            "language_model.encoder.layers.{bid}.self_attention.q_layernorm",
+            "encoder.layers.{bid}.self_attention.q_layernorm",
             "model.layers.{bid}.self_attn.q_layernorm",                       # persimmon
             "model.layers.{bid}.self_attn.query_layernorm",                   # hunyuan
             "model.layers.{bid}.attention.query_layernorm",                   # bailingmoe2
@@ -679,7 +670,7 @@ class TensorNameMap:
         ),
 
         MODEL_TENSOR.ATTN_K_NORM: (
-            "language_model.encoder.layers.{bid}.self_attention.k_layernorm",
+            "encoder.layers.{bid}.self_attention.k_layernorm",
             "model.layers.{bid}.self_attn.k_layernorm",                       # persimmon
             "model.layers.{bid}.self_attn.key_layernorm",                     # hunyuan
             "model.layers.{bid}.attention.key_layernorm",                     # bailingmoe2
@@ -695,7 +686,7 @@ class TensorNameMap:
         ),
 
         MODEL_TENSOR.ROPE_FREQS: (
-            "language_model.encoder.layers.{bid}.self_attention.rotary_emb.inv_freq",  # persimmon
+            "encoder.layers.{bid}.self_attention.rotary_emb.inv_freq",  # persimmon
         ),
 
         MODEL_TENSOR.LAYER_OUT_NORM: (

include/llama.h

@@ -867,6 +867,9 @@ extern "C" {
 // work only with partial states, such as SWA KV cache or recurrent cache (e.g. Mamba)
 #define LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY 1
 
+// keeps the tensor data on device buffers (i.e. not accessible in host memory, but faster save/load)
+#define LLAMA_STATE_SEQ_FLAGS_ON_DEVICE 2
+
     typedef uint32_t llama_state_seq_flags;
 
     LLAMA_API size_t llama_state_seq_get_size_ext(

src/llama-context.cpp

@@ -2240,13 +2240,17 @@ llm_graph_cb llama_context::graph_get_cb() const {
 
 class llama_io_write_dummy : public llama_io_write_i {
 public:
-    llama_io_write_dummy() = default;
+    llama_io_write_dummy(bool skip_tensors) : skip_tensors(skip_tensors) {}
 
     void write(const void * /* src */, size_t size) override {
         size_written += size;
     }
 
-    void write_tensor(const ggml_tensor * /* tensor */, size_t /* offset */, size_t size) override {
+    void write_tensor(ggml_tensor * /* tensor */, size_t /* offset */, size_t size) override {
+        if (skip_tensors) {
+            return;
+        }
+
         size_written += size;
     }
 
@@ -2255,34 +2259,21 @@ public:
     }
 
 private:
+    const bool skip_tensors;
+
     size_t size_written = 0;
 };
 
-class llama_io_write_buffer : public llama_io_write_i {
+class llama_io_write_host : public llama_io_write_i {
 public:
-    llama_io_write_buffer(
+    llama_io_write_host(
             uint8_t * p, size_t len) : ptr(p), buf_size(len) {}
 
-    ~llama_io_write_buffer() {
-#if 1
+    ~llama_io_write_host() {
         // TODO: add backend support to batch tensor_get? or some other way to speed this up
-        for (const auto & info : winfos) {
-            ggml_backend_tensor_get(info.tensor, info.ptr, info.offset, info.size);
+        for (const auto & winfo : winfos) {
+            ggml_backend_tensor_get(winfo.tensor, winfo.ptr, winfo.offset, winfo.size);
         }
-#else
-        // flush the writes asynchronously
-        // this helps on Macs, but on other devices - it does not. just an example
-        std::vector<std::future<void>> futures;
-        futures.reserve(winfos.size());
-        for (const auto & info : winfos) {
-            futures.push_back(std::async(std::launch::async, [info]() {
-                ggml_backend_tensor_get(info.tensor, info.ptr, info.offset, info.size);
-            }));
-        }
-        for (auto & f : futures) {
-            f.wait();
-        }
-#endif
     }
 
     void write(const void * src, size_t size) override {
@@ -2295,7 +2286,7 @@ public:
         buf_size -= size;
     }
 
-    void write_tensor(const ggml_tensor * tensor, size_t offset, size_t size) override {
+    void write_tensor(ggml_tensor * tensor, size_t offset, size_t size) override {
         if (size > buf_size) {
             throw std::runtime_error("unexpectedly reached end of buffer");
         }
@@ -2318,7 +2309,7 @@ private:
     size_t size_written = 0;
 
     struct write_info {
-        const ggml_tensor * tensor;
+        ggml_tensor * tensor;
         uint8_t * ptr;
         size_t size;
         size_t offset;
@@ -2326,14 +2317,14 @@ private:
     std::vector<write_info> winfos;
 };
 
-class llama_io_read_buffer : public llama_io_read_i {
+class llama_io_read_host : public llama_io_read_i {
 public:
-    llama_io_read_buffer(const uint8_t * p, size_t len) : ptr(p), buf_size(len) {}
+    llama_io_read_host(const uint8_t * p, size_t len) : ptr(p), buf_size(len) {}
 
-    ~llama_io_read_buffer() {
+    ~llama_io_read_host() {
         // flush the reads
-        for (const auto & info : rinfos) {
-            ggml_backend_tensor_set(info.tensor, info.ptr, info.offset, info.size);
+        for (const auto & rinfo : rinfos) {
+            ggml_backend_tensor_set(rinfo.tensor, rinfo.ptr, rinfo.offset, rinfo.size);
         }
     }
 
@@ -2387,7 +2378,7 @@ public:
         size_written += size;
     }
 
-    void write_tensor(const ggml_tensor * tensor, size_t offset, size_t size) override {
+    void write_tensor(ggml_tensor * tensor, size_t offset, size_t size) override {
         temp_buffer.resize(size);
         ggml_backend_tensor_get(tensor, temp_buffer.data(), offset, size);
         write(temp_buffer.data(), temp_buffer.size());
@@ -2428,8 +2419,162 @@ private:
     std::vector<uint8_t> temp_buffer;
 };
 
+class llama_io_write_device : public llama_io_write_i {
+public:
+    llama_io_write_device(uint8_t * p, size_t len, llama_memory_buffers & mbufs) : ptr(p), buf_size(len), mbufs(mbufs)  {
+    }
+
+    ~llama_io_write_device() {
+        llama_memory_buffers mbufs_new;
+
+        for (const auto & winfo : winfos) {
+            auto * buft = ggml_backend_buffer_get_type(winfo.tensor->buffer);
+
+            mbufs_new[buft].n_tensors++;
+            mbufs_new[buft].total_size += winfo.size;
+        }
+
+        for (auto & [buft, mbuf] : mbufs_new) {
+            ggml_init_params params = {
+                /*.mem_size   =*/ 2*mbuf.n_tensors*ggml_tensor_overhead(),
+                /*.mem_buffer =*/ NULL,
+                /*.no_alloc   =*/ true,
+            };
+
+            mbuf.ctx.reset(ggml_init(params));
+
+            mbuf.org.reserve(mbuf.n_tensors);
+            mbuf.cpy.reserve(mbuf.n_tensors);
+        }
+
+        for (const auto & winfo : winfos) {
+            auto * buft = ggml_backend_buffer_get_type(winfo.tensor->buffer);
+
+            const int64_t n = winfo.size/ggml_element_size(winfo.tensor);
+
+            auto & mbuf = mbufs_new[buft];
+
+            mbuf.org.push_back(ggml_view_1d      (mbuf.ctx.get(), winfo.tensor, n, winfo.offset));
+            mbuf.cpy.push_back(ggml_new_tensor_1d(mbuf.ctx.get(), winfo.tensor->type, n));
+        }
+
+        for (auto & [buft, mbuf] : mbufs_new) {
+            auto & mbuf_cur = mbufs[buft];
+
+            if (!mbuf_cur.buf || mbuf_cur.org.size() != mbuf.org.size() || mbuf_cur.total_size != mbuf.total_size) {
+                mbuf_cur = std::move(mbuf);
+
+                mbuf_cur.buf.reset(ggml_backend_alloc_ctx_tensors_from_buft(mbuf_cur.ctx.get(), buft));
+
+                LLAMA_LOG_INFO("%s: allocated '%s' buffer %.3f MiB\n", __func__, ggml_backend_buft_name(buft), mbuf.total_size/1024.0/1024.0);
+            }
+
+            for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
+                ggml_backend_tensor_copy(mbuf_cur.org[i], mbuf_cur.cpy[i]);
+            }
+        }
+    }
+
+    void write(const void * src, size_t size) override {
+        if (size > buf_size) {
+            throw std::runtime_error("unexpectedly reached end of buffer");
+        }
+        memcpy(ptr, src, size);
+        ptr += size;
+        size_written += size;
+        buf_size -= size;
+    }
+
+    void write_tensor(ggml_tensor * tensor, size_t offset, size_t size) override {
+        // save the write for later during destruction
+        winfos.push_back({tensor, ptr, size, offset});
+    }
+
+    size_t n_bytes() override {
+        return size_written;
+    }
+
+private:
+    uint8_t * ptr;
+    size_t buf_size = 0;
+    size_t size_written = 0;
+
+    struct write_info {
+        ggml_tensor * tensor;
+        uint8_t * ptr;
+        size_t size;
+        size_t offset;
+    };
+    std::vector<write_info> winfos;
+
+    llama_memory_buffers & mbufs;
+};
+
+class llama_io_read_device : public llama_io_read_i {
+public:
+    llama_io_read_device(const uint8_t * p, size_t len, const llama_memory_buffers & mbufs) : ptr(p), buf_size(len), mbufs(mbufs) {
+    }
+
+    ~llama_io_read_device() {
+        llama_memory_buffers mbufs_new;
+
+        for (const auto & rinfo : rinfos) {
+            auto * buft = ggml_backend_buffer_get_type(rinfo.tensor->buffer);
+
+            mbufs_new[buft].n_tensors++;
+            mbufs_new[buft].total_size += rinfo.size;
+        }
+
+        for (auto & [buft, mbuf] : mbufs_new) {
+            const auto & mbuf_cur = mbufs.at(buft);
+
+            if (!mbuf_cur.buf || mbuf_cur.n_tensors != mbuf.n_tensors || mbuf_cur.total_size != mbuf.total_size) {
+                GGML_ABORT("%s: memory buffer mismatch\n", __func__);
+            }
+
+            for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
+                ggml_backend_tensor_copy(mbuf_cur.cpy[i], mbuf_cur.org[i]);
+            }
+        }
+    }
+
+    void read(void * dst, size_t size) override {
+        if (size > buf_size) {
+            throw std::runtime_error("unexpectedly reached end of buffer");
+        }
+        memcpy(dst, ptr, size);
+        ptr += size;
+        size_read += size;
+        buf_size -= size;
+    }
+
+    void read_tensor(ggml_tensor * tensor, size_t offset, size_t size) override {
+        // save for later during destruction
+        rinfos.push_back({tensor, ptr, size, offset});
+    }
+
+    size_t n_bytes() override {
+        return size_read;
+    }
+
+private:
+    const uint8_t * ptr;
+    size_t buf_size = 0;
+    size_t size_read = 0;
+
+    struct read_info {
+        ggml_tensor * tensor;
+        const uint8_t * ptr;
+        size_t size;
+        size_t offset;
+    };
+    std::vector<read_info> rinfos;
+
+    const llama_memory_buffers & mbufs;
+};
+
 size_t llama_context::state_get_size() {
-    llama_io_write_dummy io;
+    llama_io_write_dummy io(false);
     try {
         return state_write_data(io);
     } catch (const std::exception & err) {
@@ -2439,7 +2584,7 @@ size_t llama_context::state_get_size() {
 }
 
 size_t llama_context::state_get_data(uint8_t * dst, size_t size) {
-    llama_io_write_buffer io(dst, size);
+    llama_io_write_host io(dst, size);
     try {
         return state_write_data(io);
     } catch (const std::exception & err) {
@@ -2449,7 +2594,7 @@ size_t llama_context::state_get_data(uint8_t * dst, size_t size) {
 }
 
 size_t llama_context::state_set_data(const uint8_t * src, size_t size) {
-    llama_io_read_buffer io(src, size);
+    llama_io_read_host io(src, size);
     try {
         return state_read_data(io);
     } catch (const std::exception & err) {
@@ -2458,9 +2603,14 @@ size_t llama_context::state_set_data(const uint8_t * src, size_t size) {
     }
 }
 
+static constexpr uint32_t io_magic = 0xaf143cd8;
+
 size_t llama_context::state_seq_get_size(llama_seq_id seq_id, llama_state_seq_flags flags) {
-    llama_io_write_dummy io;
+    llama_io_write_dummy io(flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE);
     try {
+        io.write(&io_magic, sizeof(io_magic));
+        io.write(&seq_id, sizeof(seq_id));
+
         return state_seq_write_data(io, seq_id, flags);
     } catch (const std::exception & err) {
         LLAMA_LOG_ERROR("%s: error getting state size: %s\n", __func__, err.what());
@@ -2469,9 +2619,18 @@ size_t llama_context::state_seq_get_size(llama_seq_id seq_id, llama_state_seq_fl
 }
 
 size_t llama_context::state_seq_get_data(llama_seq_id seq_id, uint8_t * dst, size_t size, llama_state_seq_flags flags) {
-    llama_io_write_buffer io(dst, size);
+    std::unique_ptr<llama_io_write_i> io;
+    if (flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) {
+        io = std::make_unique<llama_io_write_device>(dst, size, mem_storage[seq_id]);
+    } else {
+        io = std::make_unique<llama_io_write_host>(dst, size);
+    }
+
     try {
-        return state_seq_write_data(io, seq_id, flags);
+        io->write(&io_magic, sizeof(io_magic));
+        io->write(&seq_id, sizeof(seq_id));
+
+        return state_seq_write_data(*io, seq_id, flags);
     } catch (const std::exception & err) {
         LLAMA_LOG_ERROR("%s: error saving state: %s\n", __func__, err.what());
         return 0;
@@ -2479,9 +2638,43 @@ size_t llama_context::state_seq_get_data(llama_seq_id seq_id, uint8_t * dst, siz
 }
 
 size_t llama_context::state_seq_set_data(llama_seq_id seq_id, const uint8_t * src, size_t size, llama_state_seq_flags flags) {
-    llama_io_read_buffer io(src, size);
+    std::unique_ptr<llama_io_read_i> io;
+    if (flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) {
+        // create a temporary io to read the magic and the src seq_id
+        io = std::make_unique<llama_io_read_host>(src, size);
+
+        uint32_t magic_read;
+        io->read(&magic_read, sizeof(magic_read));
+        if (io_magic != magic_read) {
+            throw std::runtime_error("wrong sequence state magic");
+        }
+
+        llama_seq_id seq_id_read;
+        io->read(&seq_id_read, sizeof(seq_id_read));
+
+        GGML_ASSERT(mem_storage.find(seq_id_read) != mem_storage.end());
+
+        io = std::make_unique<llama_io_read_device>(src, size, mem_storage[seq_id_read]);
+    } else {
+        io = std::make_unique<llama_io_read_host>(src, size);
+    }
+
     try {
-        return state_seq_read_data(io, seq_id, flags);
+        uint32_t magic_read;
+        io->read(&magic_read, sizeof(magic_read));
+        if (io_magic != magic_read) {
+            throw std::runtime_error("wrong sequence state magic");
+        }
+
+        const bool need_seq_match = (flags & LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+
+        llama_seq_id seq_id_read;
+        io->read(&seq_id_read, sizeof(seq_id_read));
+        if (need_seq_match && seq_id != seq_id_read) {
+            throw std::runtime_error("wrong sequence id");
+        }
+
+        return state_seq_read_data(*io, seq_id, flags);
     } catch (const std::exception & err) {
         LLAMA_LOG_ERROR("%s: error loading state: %s\n", __func__, err.what());
         return 0;
@@ -3472,7 +3665,6 @@ size_t llama_state_seq_get_data_ext(llama_context * ctx, uint8_t * dst, size_t s
 
     return ctx->state_seq_get_data(seq_id, dst, size, flags);
 }
-
 size_t llama_state_seq_set_data_ext(llama_context * ctx, const uint8_t * src, size_t size, llama_seq_id seq_id, llama_state_seq_flags flags) {
     ctx->synchronize();
 

src/llama-context.h

@@ -23,6 +23,21 @@ class llama_io_write_i;
 struct llama_memory_i;
 struct llama_memory_context_i;
 
+// stores copy of the memory in device buffer. used for fast state save/load
+struct llama_memory_buffer {
+    int n_tensors = 0;
+    size_t total_size = 0;
+
+    ggml_backend_buffer_ptr buf;
+
+    ggml_context_ptr ctx;
+
+    std::vector<ggml_tensor *> org;
+    std::vector<ggml_tensor *> cpy;
+};
+
+using llama_memory_buffers = std::map<ggml_backend_buffer_type_t, llama_memory_buffer>;
+
 struct llama_context {
     // init scheduler and compute buffers, reserve worst-case graphs
     llama_context(
@@ -128,6 +143,7 @@ struct llama_context {
     size_t state_set_data(const uint8_t * src, size_t size);
 
     size_t state_seq_get_size(llama_seq_id seq_id, llama_state_seq_flags flags);
+
     size_t state_seq_get_data(llama_seq_id seq_id,       uint8_t * dst, size_t size, llama_state_seq_flags flags);
     size_t state_seq_set_data(llama_seq_id seq_id, const uint8_t * src, size_t size, llama_state_seq_flags flags);
 
@@ -328,6 +344,9 @@ private:
     // host buffer for the model output (logits and embeddings)
     ggml_backend_buffer_ptr buf_output;
 
+    // keep copies of the per-sequence memory on the device
+    std::map<llama_seq_id, llama_memory_buffers> mem_storage;
+
     bool has_evaluated_once = false;
 
     // env: LLAMA_GRAPH_REUSE_DISABLE

src/llama-graph.cpp

@@ -65,8 +65,13 @@ static bool can_reuse_kq_mask(
 
 //     ggml_tensor * res;
 
-//     res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
+//     if (!ggml_is_contiguous(cur)) {
+//         res = ggml_cont_2d   (ctx, cur, n, ggml_nelements(cur)/n);
+//     } else {
+//         res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
+//     }
 //     res = ggml_mul_mat   (ctx, rot, res);
+//     ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
 //     res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
 
 //     return res;

src/llama-io.h

@@ -12,7 +12,7 @@ public:
     virtual ~llama_io_write_i() = default;
 
     virtual void write(const void * src, size_t size) = 0;
-    virtual void write_tensor(const ggml_tensor * tensor, size_t offset, size_t size) = 0;
+    virtual void write_tensor(ggml_tensor * tensor, size_t offset, size_t size) = 0;
 
     // bytes written so far
     virtual size_t n_bytes() = 0;

src/llama-kv-cache.cpp

@@ -65,8 +65,13 @@ static ggml_tensor * ggml_mul_mat_aux(
 
     ggml_tensor * res;
 
-    res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
+    if (!ggml_is_contiguous(cur)) {
+        res = ggml_cont_2d   (ctx, cur, n, ggml_nelements(cur)/n);
+    } else {
+        res = ggml_reshape_2d(ctx, cur, n, ggml_nelements(cur)/n);
+    }
     res = ggml_mul_mat   (ctx, rot, res);
+    ggml_mul_mat_set_hint(res, GGML_HINT_SRC0_IS_HADAMARD);
     res = ggml_reshape_4d(ctx, res, cur->ne[0], cur->ne[1], cur->ne[2], cur->ne[3]);
 
     return res;

src/llama-memory-recurrent.cpp

@@ -784,7 +784,7 @@ void llama_memory_recurrent::state_write_data(llama_io_write_i & io, const std::
     const uint32_t n_layer = hparams.n_layer;
 
     io.write(&s_trans, sizeof(s_trans));
-    io.write(&n_layer,   sizeof(n_layer));
+    io.write(&n_layer, sizeof(n_layer));
 
     // Iterate and write all the R tensors first, each row is a cell
     // Get whole range at a time

src/llama.cpp

@@ -112,6 +112,10 @@ void llama_backend_init(void) {
         struct ggml_context * ctx = ggml_init(params);
         ggml_free(ctx);
     }
+
+    if (!ggml_backend_reg_count()) {
+        ggml_backend_load_all();
+    }
 }
 
 void llama_numa_init(enum ggml_numa_strategy numa) {

tools/mtmd/debug/mtmd-debug.cpp

@@ -68,6 +68,8 @@ int main(int argc, char ** argv) {
         return 1;
     }
 
+    ggml_backend_load_all();
+
     LOG_INF("%s: loading model: %s\n", __func__, params.model.path.c_str());
 
     mtmd::context_ptr ctx_mtmd;

tools/mtmd/mtmd-cli.cpp

@@ -295,6 +295,8 @@ int main(int argc, char ** argv) {
         return 1;
     }
 
+    ggml_backend_load_all();
+
     mtmd_cli_context ctx(params);
     LOG_INF("%s: loading model: %s\n", __func__, params.model.path.c_str());
 

tools/server/server-context.cpp

@@ -36,7 +36,7 @@ using json = nlohmann::ordered_json;
 
 constexpr int HTTP_POLLING_SECONDS = 1;
 
-static void server_prompt_checkpoint_update(server_prompt_checkpoint & ckpt, llama_context * ctx, int id, int64_t n_tokens, llama_pos pos_min = -1, llama_pos pos_max = -1) {
+static void server_prompt_checkpoint_update(server_prompt_checkpoint & ckpt, llama_context * ctx, int id, int64_t n_tokens, bool on_device, llama_pos pos_min = -1, llama_pos pos_max = -1) {
     if (pos_min == -1) {
         pos_min = llama_memory_seq_pos_min(llama_get_memory(ctx), id);
     }
@@ -44,14 +44,19 @@ static void server_prompt_checkpoint_update(server_prompt_checkpoint & ckpt, lla
         pos_max = llama_memory_seq_pos_max(llama_get_memory(ctx), id);
     }
 
-    const size_t checkpoint_size = llama_state_seq_get_size_ext(ctx, id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+    auto flags = LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY;
+    if (on_device) {
+        flags |= LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+    }
+
+    const size_t checkpoint_size = llama_state_seq_get_size_ext(ctx, id, flags);
 
     ckpt.pos_min  = pos_min;
     ckpt.pos_max  = pos_max;
     ckpt.n_tokens = n_tokens;
     ckpt.data.resize(checkpoint_size);
 
-    const size_t n = llama_state_seq_get_data_ext(ctx, ckpt.data.data(), checkpoint_size, id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+    const size_t n = llama_state_seq_get_data_ext(ctx, ckpt.data.data(), checkpoint_size, id, flags);
     if (n != checkpoint_size) {
         GGML_ABORT("checkpoint size mismatch: expected %zu, got %zu\n", checkpoint_size, n);
     }
@@ -362,7 +367,7 @@ struct server_slot {
 
                     //const int64_t t_start = ggml_time_us();
 
-                    server_prompt_checkpoint_update(spec_ckpt, ctx, this->id, n_tokens);
+                    server_prompt_checkpoint_update(spec_ckpt, ctx, this->id, n_tokens, true);
 
                     //const int64_t t_total = ggml_time_us() - t_start;
                     //printf("checkpoint total: %f ms\n", t_total / 1000.0);
@@ -1838,7 +1843,7 @@ private:
         }
 
         auto & cur = slot.prompt.checkpoints.emplace_back();
-        server_prompt_checkpoint_update(cur, ctx, slot.id, slot.prompt.n_tokens() - n_tokens_cur, pos_min, pos_max);
+        server_prompt_checkpoint_update(cur, ctx, slot.id, slot.prompt.n_tokens() - n_tokens_cur, false, pos_min, pos_max);
 
         SLT_WRN(slot,
                 "created context checkpoint %d of %d (pos_min = %d, pos_max = %d, n_tokens = %" PRId64 ", size = %.3f MiB)\n",
@@ -3003,7 +3008,7 @@ private:
                             SLT_DBG(slot, "restoring speculative checkpoint (pos_min = %d, pos_max = %d, size = %zu)\n",
                                     ckpt.pos_min, ckpt.pos_max, ckpt.size());
 
-                            const size_t n = llama_state_seq_set_data_ext(slot.ctx, ckpt.data.data(), ckpt.size(), slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+                            const size_t n = llama_state_seq_set_data_ext(slot.ctx, ckpt.data.data(), ckpt.size(), slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE);
                             if (n != ckpt.size()) {
                                 GGML_ABORT("%s: failed to restore context checkpoint (pos_min=%d, pos_max=%d, size=%zu, get_data_ext->%zu, set_data_ext->%zu",
                                         __func__, ckpt.pos_min, ckpt.pos_max, ckpt.size(), ckpt.size(), n);

tools/server/server-tools.cpp

@@ -10,6 +10,7 @@
 #include <atomic>
 #include <cstring>
 #include <climits>
+#include <algorithm>
 
 namespace fs = std::filesystem;
 
@@ -744,6 +745,24 @@ void server_tools::setup(const std::vector<std::string> & enabled_tools) {
         std::unordered_set<std::string> enabled_set(enabled_tools.begin(), enabled_tools.end());
         auto all_tools = build_tools();
 
+        // collect all known tool names for validation
+        std::vector<std::string> known_names;
+        known_names.reserve(all_tools.size());
+        for (const auto & t : all_tools) {
+            known_names.push_back(t->name);
+        }
+
+        // validate that every requested tool is known
+        for (const auto & name : enabled_tools) {
+            if (name == "all") continue;
+            if (std::find(known_names.begin(), known_names.end(), name) == known_names.end()) {
+                throw std::runtime_error(string_format(
+                    "unknown tool \"%s\". available tools: %s",
+                    name.c_str(),
+                    string_join(known_names, ", ").c_str()));
+            }
+        }
+
         tools.clear();
         for (auto & t : all_tools) {
             if (enabled_set.count(t->name) > 0 || enabled_set.count("all") > 0) {

tools/server/server.cpp

@@ -215,7 +215,12 @@ int main(int argc, char ** argv) {
     }
     // EXPERIMENTAL built-in tools
     if (!params.server_tools.empty()) {
-        tools.setup(params.server_tools);
+        try {
+            tools.setup(params.server_tools);
+        } catch (const std::exception & e) {
+            LOG_ERR("%s: tools setup failed: %s\n", __func__, e.what());
+            return 1;
+        }
         SRV_WRN("%s", "-----------------\n");
         SRV_WRN("%s", "Built-in tools are enabled, do not expose server to untrusted environments\n");
         SRV_WRN("%s", "This feature is EXPERIMENTAL and may be changed in the future\n");

tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddMcpServersSubmenu.svelte

@@ -2,7 +2,7 @@
 	import { Settings, Plus } from '@lucide/svelte';
 	import { Switch } from '$lib/components/ui/switch';
 	import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
-	import { McpLogo, DropdownMenuSearchable } from '$lib/components/app';
+	import { McpLogo, DropdownMenuSearchable, McpServerIdentity } from '$lib/components/app';
 	import { conversationsStore } from '$lib/stores/conversations.svelte';
 	import { mcpStore } from '$lib/stores/mcp.svelte';
 	import { HealthCheckStatus } from '$lib/enums';
@@ -77,6 +77,8 @@
 							{@const healthState = mcpStore.getHealthCheckState(server.id)}
 							{@const hasError = healthState.status === HealthCheckStatus.ERROR}
 							{@const isEnabledForChat = isServerEnabledForChat(server.id)}
+							{@const displayName = getServerLabel(server)}
+							{@const faviconUrl = mcpStore.getServerFavicon(server.id)}
 
 							<button
 								type="button"
@@ -85,18 +87,16 @@
 								disabled={hasError}
 							>
 								<div class="flex min-w-0 flex-1 items-center gap-2">
-									{#if mcpStore.getServerFavicon(server.id)}
-										<img
-											src={mcpStore.getServerFavicon(server.id)}
-											alt=""
-											class="h-4 w-4 shrink-0 rounded-sm"
-											onerror={(e) => {
-												(e.currentTarget as HTMLImageElement).style.display = 'none';
-											}}
+									<div class="min-w-0 flex-1">
+										<McpServerIdentity
+											{displayName}
+											{faviconUrl}
+											iconClass="h-4 w-4"
+											iconRounded="rounded-sm"
+											showVersion={false}
+											nameClass="text-sm"
 										/>
-									{/if}
-
-									<span class="truncate text-sm">{getServerLabel(server)}</span>
+									</div>
 
 									{#if hasError}
 										<span

tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormPickers/ChatFormPicker/ChatFormPickerPopover.svelte

@@ -29,7 +29,11 @@
 		}
 	}}
 >
-	<Popover.Trigger class="pointer-events-none absolute inset-0 opacity-0">
+	<Popover.Trigger
+		class="pointer-events-none absolute inset-0 opacity-0"
+		tabindex={-1}
+		aria-hidden="true"
+	>
 		<span class="sr-only">{srLabel}</span>
 	</Popover.Trigger>
 

tools/server/webui/src/lib/components/app/mcp/McpResourcesBrowser/McpResourcesBrowserServerItem.svelte

@@ -12,6 +12,7 @@
 		sortTreeChildren
 	} from './mcp-resources-browser';
 	import { getDisplayName, getResourceIcon } from '$lib/utils';
+	import { McpServerIdentity } from '$lib/components/app/mcp';
 
 	interface Props {
 		serverName: string;
@@ -43,11 +44,12 @@
 		searchQuery = ''
 	}: Props = $props();
 
+	let serverDisplayName = $derived(mcpStore.getServerDisplayName(serverName));
+	let serverFaviconUrl = $derived(mcpStore.getServerFavicon(serverName));
+
 	const hasResources = $derived(serverRes.resources.length > 0);
 	const hasTemplates = $derived(serverRes.templates.length > 0);
 	const hasContent = $derived(hasResources || hasTemplates);
-	const displayName = $derived(mcpStore.getServerDisplayName(serverName));
-	const favicon = $derived(mcpStore.getServerFavicon(serverName));
 	const resourceTree = $derived(buildResourceTree(serverRes.resources, serverName, searchQuery));
 
 	const templateInfos = $derived<MCPResourceTemplateInfo[]>(
@@ -153,21 +155,15 @@
 			<ChevronRight class="h-3.5 w-3.5" />
 		{/if}
 
-		<span class="inline-flex flex-col items-start text-left">
-			<span class="inline-flex items-center justify-start gap-1.5 font-medium">
-				{#if favicon}
-					<img
-						src={favicon}
-						alt=""
-						class="h-4 w-4 shrink-0 rounded-sm"
-						onerror={(e) => {
-							(e.currentTarget as HTMLImageElement).style.display = 'none';
-						}}
-					/>
-				{/if}
-
-				{displayName}
-			</span>
+		<span class="inline-flex flex-col items-start gap-1 text-left">
+			<div class="inline-flex min-w-0 items-center gap-1.5">
+				<McpServerIdentity
+					displayName={serverDisplayName}
+					faviconUrl={serverFaviconUrl}
+					iconClass="h-4 w-4"
+					showVersion={false}
+				/>
+			</div>
 
 			<span class="text-xs text-muted-foreground">
 				({serverRes.resources.length} resource{serverRes.resources.length !== 1

tools/server/webui/src/lib/components/app/mcp/McpServerCard/McpServerCard.svelte

@@ -17,17 +17,17 @@
 
 	interface Props {
 		server: MCPServerSettingsEntry;
-		faviconUrl: string | null;
 		enabled?: boolean;
 		onToggle: (enabled: boolean) => void;
 		onUpdate: (updates: Partial<MCPServerSettingsEntry>) => void;
 		onDelete: () => void;
 	}
 
-	let { server, faviconUrl, enabled, onToggle, onUpdate, onDelete }: Props = $props();
+	let { server, enabled, onToggle, onUpdate, onDelete }: Props = $props();
 
 	let healthState = $derived<HealthCheckState>(mcpStore.getHealthCheckState(server.id));
 	let displayName = $derived(mcpStore.getServerLabel(server));
+	let faviconUrl = $derived(mcpStore.getServerFavicon(server.id));
 	let isIdle = $derived(healthState.status === HealthCheckStatus.IDLE);
 	let isHealthChecking = $derived(healthState.status === HealthCheckStatus.CONNECTING);
 	let isConnected = $derived(healthState.status === HealthCheckStatus.SUCCESS);

tools/server/webui/src/lib/components/app/mcp/McpServerCard/McpServerCardHeader.svelte

@@ -1,15 +1,14 @@
 <script lang="ts">
-	import { Cable, ExternalLink } from '@lucide/svelte';
 	import { Switch } from '$lib/components/ui/switch';
 	import { Badge } from '$lib/components/ui/badge';
-	import { McpCapabilitiesBadges } from '$lib/components/app/mcp';
+	import { McpCapabilitiesBadges, McpServerIdentity } from '$lib/components/app/mcp';
 	import { MCP_TRANSPORT_LABELS, MCP_TRANSPORT_ICONS } from '$lib/constants';
 	import { MCPTransportType } from '$lib/enums';
 	import type { MCPServerInfo, MCPCapabilitiesInfo } from '$lib/types';
 
 	interface Props {
 		displayName: string;
-		faviconUrl: string | null;
+		faviconUrl?: string | null;
 		enabled: boolean;
 		disabled?: boolean;
 		onToggle: (enabled: boolean) => void;
@@ -32,42 +31,16 @@
 
 <div class="space-y-3">
 	<div class="flex items-start justify-between gap-3">
-		<div class="grid min-w-0 gap-3">
-			<div class="flex items-center gap-2 overflow-hidden">
-				{#if faviconUrl}
-					<img
-						src={faviconUrl}
-						alt=""
-						class="h-5 w-5 shrink-0 rounded"
-						onerror={(e) => {
-							(e.currentTarget as HTMLImageElement).style.display = 'none';
-						}}
-					/>
-				{:else}
-					<div class="flex h-5 w-5 shrink-0 items-center justify-center rounded bg-muted">
-						<Cable class="h-3 w-3 text-muted-foreground" />
-					</div>
-				{/if}
-
-				<p class="min-w-0 shrink-0 truncate leading-none font-medium">{displayName}</p>
-
-				{#if serverInfo?.version}
-					<Badge variant="secondary" class="h-4 min-w-0 truncate px-1 text-[10px]">
-						v{serverInfo.version}
-					</Badge>
-				{/if}
-
-				{#if serverInfo?.websiteUrl}
-					<a
-						href={serverInfo.websiteUrl}
-						target="_blank"
-						rel="noopener noreferrer"
-						class="shrink-0 text-muted-foreground hover:text-foreground"
-						aria-label="Open website"
-					>
-						<ExternalLink class="h-3 w-3" />
-					</a>
-				{/if}
+		<div class="flex min-w-0 flex-col gap-3">
+			<div class="inline-flex items-center gap-2">
+				<McpServerIdentity
+					{displayName}
+					{faviconUrl}
+					{serverInfo}
+					iconClass="h-5 w-5"
+					iconRounded="rounded"
+					nameClass="leading-6 font-medium"
+				/>
 			</div>
 
 			{#if capabilities || transportType}

tools/server/webui/src/lib/components/app/mcp/McpServerIdentity.svelte

@@ -0,0 +1,67 @@
+<script lang="ts">
+	import { ExternalLink } from '@lucide/svelte';
+	import { Badge } from '$lib/components/ui/badge';
+	import { TruncatedText } from '$lib/components/app/misc';
+	import { sanitizeExternalUrl } from '$lib/utils';
+	import type { MCPServerInfo } from '$lib/types';
+
+	interface Props {
+		displayName?: string;
+		faviconUrl?: string | null;
+		serverInfo?: MCPServerInfo;
+		iconClass?: string;
+		iconRounded?: string;
+		showVersion?: boolean;
+		showWebsite?: boolean;
+		nameClass?: string;
+	}
+
+	let {
+		displayName,
+		faviconUrl = null,
+		serverInfo,
+		iconClass = 'h-5 w-5',
+		iconRounded = 'rounded-sm',
+		showVersion = true,
+		showWebsite = true,
+		nameClass
+	}: Props = $props();
+
+	let safeWebsiteUrl = $derived(
+		serverInfo?.websiteUrl ? sanitizeExternalUrl(serverInfo.websiteUrl) : null
+	);
+</script>
+
+<span class="flex min-w-0 items-center gap-1.5">
+	{#if faviconUrl}
+		<img
+			src={faviconUrl}
+			alt=""
+			class={['shrink-0', iconRounded, iconClass]}
+			onerror={(e) => {
+				(e.currentTarget as HTMLImageElement).style.display = 'none';
+			}}
+		/>
+	{/if}
+
+	<TruncatedText text={displayName ?? ''} class={nameClass ?? ''} />
+
+	{#if showVersion && serverInfo?.version}
+		<Badge variant="secondary" class="h-4 min-w-0 shrink px-1 text-[10px]">
+			<TruncatedText text={`v${serverInfo.version}`} />
+		</Badge>
+	{/if}
+
+	{#if showWebsite && safeWebsiteUrl}
+		<a
+			href={safeWebsiteUrl}
+			target="_blank"
+			rel="noopener noreferrer"
+			class="shrink-0 text-muted-foreground hover:text-foreground"
+			aria-label="Open website"
+			onclick={(e) => e.stopPropagation()}
+		>
+			<ExternalLink class="h-3 w-3" />
+		</a>
+	{/if}
+</span>

tools/server/webui/src/lib/components/app/mcp/index.ts

@@ -180,6 +180,25 @@ export { default as McpServerCardDeleteDialog } from './McpServerCard/McpServerC
 /** Skeleton loading state for server card during health checks. */
 export { default as McpServerCardSkeleton } from './McpServerCardSkeleton.svelte';
 
+/**
+ * **McpServerIdentity** - Server identity display (icon, name, version)
+ *
+ * Reusable headless component for displaying server name, favicon/icon, and version badge.
+ * Accepts all data via props with no store dependencies for predictable rendering.
+ *
+ * **Features:**
+ * - Server favicon/icon with fallback
+ * - Truncated display name with max-width
+ * - Optional version badge (v1.2.3)
+ * - Optional external link to server website
+ *
+ * @example
+ * ```svelte
+ * <McpServerIdentity displayName={name} faviconUrl={iconUrl} serverInfo={info} />
+ * ```
+ */
+export { default as McpServerIdentity } from './McpServerIdentity.svelte';
+
 /**
  * **McpServerInfo** - Server instructions display
  *

tools/server/webui/src/lib/components/app/misc/TruncatedText.svelte

@@ -32,7 +32,7 @@
 
 {#if isTruncated && showTooltip}
 	<Tooltip.Root>
-		<Tooltip.Trigger class={className}>
+		<Tooltip.Trigger class="{className} min-w-0">
 			<span bind:this={textElement} class="block truncate">
 				{text}
 			</span>
@@ -43,7 +43,7 @@
 		</Tooltip.Content>
 	</Tooltip.Root>
 {:else}
-	<span bind:this={textElement} class="{className} block truncate">
+	<span bind:this={textElement} class="{className} block min-w-0 truncate">
 		{text}
 	</span>
 {/if}

tools/server/webui/src/lib/components/app/models/ModelsSelectorSheet.svelte

@@ -170,7 +170,7 @@
 			>
 				<Package class="h-3.5 w-3.5" />
 
-				<TruncatedText text={selectedOption?.model || ''} class="min-w-0 font-medium" />
+				<TruncatedText text={selectedOption?.model || ''} class="font-medium" />
 
 				{#if ms.updating}
 					<Loader2 class="h-3 w-3.5 animate-spin" />

tools/server/webui/src/lib/components/app/settings/SettingsChat/SettingsChatToolsTab.svelte

@@ -2,28 +2,15 @@
 	import { ChevronDown, ChevronRight } from '@lucide/svelte';
 	import { Checkbox } from '$lib/components/ui/checkbox';
 	import * as Collapsible from '$lib/components/ui/collapsible';
-	import { TruncatedText } from '$lib/components/app';
+	import { TruncatedText, McpServerIdentity } from '$lib/components/app';
 	import { toolsStore } from '$lib/stores/tools.svelte';
 	import { permissionsStore } from '$lib/stores/permissions.svelte';
 	import { mcpStore } from '$lib/stores/mcp.svelte';
-	import { ToolSource } from '$lib/enums';
 	import { SvelteSet } from 'svelte/reactivity';
 
 	let expandedGroups = new SvelteSet<string>();
 	let groups = $derived(toolsStore.toolGroups);
 
-	function getFavicon(group: { source: ToolSource; label: string }): string | null {
-		if (group.source !== ToolSource.MCP) return null;
-
-		for (const server of mcpStore.getServersSorted()) {
-			if (mcpStore.getServerLabel(server) === group.label) {
-				return mcpStore.getServerFavicon(server.id);
-			}
-		}
-
-		return null;
-	}
-
 	function toggleExpanded(label: string) {
 		if (expandedGroups.has(label)) {
 			expandedGroups.delete(label);
@@ -39,8 +26,6 @@
 	<div class="space-y-2">
 		{#each groups as group (group.label)}
 			{@const isExpanded = expandedGroups.has(group.label)}
-			{@const favicon = getFavicon(group)}
-
 			<Collapsible.Root open={isExpanded} onOpenChange={() => toggleExpanded(group.label)}>
 				<Collapsible.Trigger
 					class="flex w-full items-center gap-2 rounded-lg px-3 py-2 text-sm hover:bg-muted/50"
@@ -51,19 +36,16 @@
 						<ChevronRight class="h-3.5 w-3.5 shrink-0" />
 					{/if}
 
-					<span class="inline-flex min-w-0 items-center gap-1.5 font-medium">
-						{#if favicon}
-							<img
-								src={favicon}
-								alt=""
-								class="h-4 w-4 shrink-0 rounded-sm"
-								onerror={(e) => {
-									(e.currentTarget as HTMLImageElement).style.display = 'none';
-								}}
-							/>
-						{/if}
+					{@const faviconUrl = group.serverId ? mcpStore.getServerFavicon(group.serverId) : null}
 
-						<span class="truncate">{group.label}</span>
+					<span class="inline-flex min-w-0 items-center gap-1.5 font-medium">
+						<McpServerIdentity
+							iconClass="h-4 w-4"
+							iconRounded="rounded-sm"
+							showVersion={false}
+							displayName={group.label}
+							{faviconUrl}
+						/>
 					</span>
 
 					<span class="ml-auto shrink-0 text-xs text-muted-foreground">
@@ -89,7 +71,7 @@
 								: false}
 
 							<div class="flex items-center gap-2 rounded px-2 py-1.5 text-sm hover:bg-muted/50">
-								<TruncatedText text={toolName} class="min-w-0 flex-1 truncate" showTooltip={true} />
+								<TruncatedText text={toolName} class="flex-1" showTooltip={true} />
 
 								<div class="flex w-16 shrink-0 justify-center">
 									<Checkbox

tools/server/webui/src/lib/components/app/settings/SettingsMcpServers.svelte

@@ -54,14 +54,14 @@
 	});
 </script>
 
-<div in:fade={{ duration: 150 }} class="max-h-full overflow-auto">
+<div in:fade={{ duration: 150 }} class="h-full max-h-[100dvh] overflow-y-auto">
 	<div class="flex items-center gap-2 p-4 md:absolute md:top-8 md:left-8 md:px-0 md:py-2">
 		<McpLogo class="h-5 w-5 md:h-6 md:w-6" />
 
 		<h1 class="text-xl font-semibold md:text-2xl">MCP Servers</h1>
 	</div>
 
-	<div class="sticky top-0 z-10 mt-4 flex items-start justify-end gap-4 px-8 py-4">
+	<div class="sticky top-0 z-10 mt-4 flex items-start gap-4 p-4 md:justify-end md:px-8">
 		<Button variant="outline" size="sm" class="shrink-0" onclick={() => (isAddingServer = true)}>
 			<Plus class="h-4 w-4" />
 
@@ -89,7 +89,6 @@
 					{:else}
 						<McpServerCard
 							{server}
-							faviconUrl={mcpStore.getServerFavicon(server.id)}
 							enabled={conversationsStore.isMcpServerEnabledForChat(server.id)}
 							onToggle={async () => {
 								const wasEnabled = conversationsStore.isMcpServerEnabledForChat(server.id);

tools/server/webui/src/lib/constants/favicon.ts

@@ -1,4 +0,0 @@
-export const GOOGLE_FAVICON_BASE_URL = 'https://www.google.com/s2/favicons';
-export const DEFAULT_FAVICON_SIZE = 32;
-export const DOMAIN_SEPARATOR = '.';
-export const ROOT_DOMAIN_MIN_PARTS = 2;

tools/server/webui/src/lib/constants/index.ts

@@ -13,7 +13,6 @@ export * from './code-blocks';
 export * from './code';
 export * from './context-keys';
 export * from './css-classes';
-export * from './favicon';
 export * from './floating-ui-constraints';
 export * from './formatters';
 export * from './key-value-pairs';
@@ -40,4 +39,5 @@ export * from './tools';
 export * from './tooltip-config';
 export * from './ui';
 export * from './uri-template';
+export * from './url';
 export * from './viewport';

tools/server/webui/src/lib/constants/mcp.ts

@@ -13,7 +13,9 @@ export const MCP_ALLOWED_ICON_MIME_TYPES = new Set([
 	MimeTypeImage.JPEG,
 	MimeTypeImage.JPG,
 	MimeTypeImage.SVG,
-	MimeTypeImage.WEBP
+	MimeTypeImage.WEBP,
+	MimeTypeImage.ICO,
+	MimeTypeImage.ICO_MICROSOFT
 ]);
 
 /**

tools/server/webui/src/lib/constants/url.ts

@@ -0,0 +1,186 @@
+const STD = ['com', 'net', 'org', 'gov', 'edu'] as const;
+
+const STD_MIL = [...STD, 'mil'] as const;
+
+const ccTLD_PREFIXES: Record<string, readonly string[]> = {
+	// --- Standard 5 only ---
+	ar: STD,
+	bd: STD,
+	bg: STD,
+	cn: STD_MIL,
+	eg: STD,
+	gr: STD,
+	hk: STD,
+	hr: STD,
+	lk: STD,
+	mx: STD_MIL,
+	my: STD_MIL,
+	ng: STD,
+	ph: STD,
+	pk: STD,
+	pl: STD,
+	ro: STD,
+	ru: STD,
+	sa: STD,
+	si: STD,
+	tr: STD,
+	tw: STD,
+	ua: STD,
+	ve: STD,
+
+	au: [...STD_MIL, 'id', 'asn', 'csiro'],
+	br: [
+		...STD_MIL,
+		'art',
+		'eco',
+		'eng',
+		'inf',
+		'med',
+		'psi',
+		'tmp',
+		'etc',
+		'adm',
+		'adv',
+		'arq',
+		'bio',
+		'bmd',
+		'cim',
+		'cng',
+		'cnt',
+		'coop',
+		'ecn',
+		'esp',
+		'far',
+		'fm',
+		'fnd',
+		'fot',
+		'fst',
+		'g12',
+		'ggf',
+		'imb',
+		'ind',
+		'jor',
+		'jus',
+		'leg',
+		'lel',
+		'mat',
+		'mp',
+		'mus',
+		'not',
+		'ntr',
+		'odo',
+		'ppg',
+		'pro',
+		'psc',
+		'qsl',
+		'rec',
+		'slg',
+		'srv',
+		'trd',
+		'tur',
+		'tv',
+		'vet',
+		'vlog',
+		'wiki',
+		'zlg'
+	],
+	id: [...STD_MIL, 'co', 'go', 'or', 'web', 'sch'],
+	in: [...STD_MIL, 'co', 'gen', 'ind', 'firm', 'ernet', 'nic'],
+	kr: [...STD_MIL, 'co', 'go', 'or', 'ac', 're'],
+	nz: [
+		...STD_MIL,
+		'co',
+		'gen',
+		'geek',
+		'kiwi',
+		'maori',
+		'school',
+		'govt',
+		'health',
+		'iwi',
+		'parliament'
+	],
+	sg: [...STD, 'per'],
+	th: ['co', 'go', 'or', 'in', 'ac', 'mi', 'net'],
+
+	ae: ['co', 'net', 'org', 'gov', 'ac', 'sch'],
+	hu: ['co', 'net', 'org', 'gov', 'edu'],
+	il: ['co', 'net', 'org', 'gov', 'ac', 'muni'],
+	jp: ['ac', 'ad', 'co', 'ed', 'go', 'gr', 'lg', 'ne', 'or'],
+	ke: ['co', 'or', 'ne', 'go', 'ac', 'sc'],
+	rs: ['co', 'net', 'org', 'gov', 'edu'],
+	uk: ['co', 'org', 'net', 'ac', 'gov', 'mil', 'nhs', 'police', 'mod', 'ltd', 'plc', 'me', 'sch'],
+	za: ['co', 'org', 'net', 'web', 'law', 'mil']
+};
+
+const WILDCARD_BASES: Record<string, readonly string[]> = {
+	br: ['nom', 'blog'],
+	jp: [
+		'kobe',
+		'kyoto',
+		'nagoya',
+		'osaka',
+		'sapporo',
+		'sendai',
+		'tokyo',
+		'yokohama',
+		'aichi',
+		'akita',
+		'aomori',
+		'chiba',
+		'ehime',
+		'fukui',
+		'fukuoka',
+		'fukushima',
+		'gifu',
+		'gunma',
+		'hiroshima',
+		'hokkaido',
+		'hyogo',
+		'ibaraki',
+		'ishikawa',
+		'iwate',
+		'kagawa',
+		'kagoshima',
+		'kanagawa',
+		'kochi',
+		'kumamoto',
+		'mie',
+		'miyagi',
+		'miyazaki',
+		'nagano',
+		'nara',
+		'niigata',
+		'oita',
+		'okayama',
+		'okinawa',
+		'saga',
+		'saitama',
+		'shiga',
+		'shimane',
+		'shizuoka',
+		'tochigi',
+		'tokushima',
+		'tottori',
+		'toyama',
+		'wakayama',
+		'yamagata',
+		'yamaguchi',
+		'yamanashi'
+	]
+};
+
+function buildSuffixSet(suffixes: Record<string, readonly string[]>): Set<string> {
+	const set = new Set<string>();
+
+	for (const [tld, parts] of Object.entries(suffixes)) {
+		for (const part of parts) {
+			set.add(`${part}.${tld}`);
+		}
+	}
+
+	return set;
+}
+
+export const TWO_PART_PUBLIC_SUFFIXES = buildSuffixSet(ccTLD_PREFIXES);
+export const WILDCARD_PUBLIC_SUFFIXES = buildSuffixSet(WILDCARD_BASES);

tools/server/webui/src/lib/enums/files.ts

@@ -182,7 +182,9 @@ export enum MimeTypeImage {
 	PNG = 'image/png',
 	GIF = 'image/gif',
 	WEBP = 'image/webp',
-	SVG = 'image/svg+xml'
+	SVG = 'image/svg+xml',
+	ICO = 'image/x-icon',
+	ICO_MICROSOFT = 'image/vnd.microsoft.icon'
 }
 
 export enum MimeTypeText {

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

@@ -24,10 +24,10 @@ export enum McpPromptVariant {
  */
 export enum UrlProtocol {
 	DATA = 'data:',
-	HTTP = 'http://',
-	HTTPS = 'https://',
-	WEBSOCKET = 'ws://',
-	WEBSOCKET_SECURE = 'wss://'
+	HTTP = 'http:',
+	HTTPS = 'https:',
+	WEBSOCKET = 'ws:',
+	WEBSOCKET_SECURE = 'wss:'
 }
 
 export enum HtmlInputType {

tools/server/webui/src/lib/services/mcp.service.ts

@@ -27,6 +27,7 @@ import {
 } from '$lib/enums';
 import type {
 	MCPServerConfig,
+	MCPResourceIcon,
 	ToolCallParams,
 	ToolExecutionResult,
 	Implementation,
@@ -469,10 +470,11 @@ export class MCPService {
 			title: impl.title,
 			description: impl.description,
 			websiteUrl: impl.websiteUrl,
-			icons: impl.icons?.map((icon: { src: string; mimeType?: string; sizes?: string }) => ({
+			icons: impl.icons?.map((icon: MCPResourceIcon) => ({
 				src: icon.src,
 				mimeType: icon.mimeType,
-				sizes: icon.sizes
+				sizes: icon.sizes,
+				theme: icon.theme
 			}))
 		};
 	}
@@ -581,7 +583,6 @@ export class MCPService {
 			this.createLog(MCPConnectionPhase.INITIALIZING, 'Sending initialize request...')
 		);
 
-		console.log(`[MCPService][${serverName}] Connecting to server...`);
 		try {
 			await client.connect(transport);
 			// Transport diagnostics are only for the initial handshake, not long-lived traffic.

tools/server/webui/src/lib/stores/mcp.svelte.ts

@@ -26,11 +26,10 @@ import { config, settingsStore } from '$lib/stores/settings.svelte';
 import { mcpResourceStore } from '$lib/stores/mcp-resources.svelte';
 import { mode } from 'mode-watcher';
 import {
-	getProxiedUrlString,
 	parseMcpServerSettings,
 	detectMcpTransportFromUrl,
-	getFaviconUrl,
-	uuid
+	uuid,
+	extractRootDomain
 } from '$lib/utils';
 import {
 	MCPConnectionPhase,
@@ -413,7 +412,9 @@ class MCPStore {
 	#isValidIconUri(src: string): boolean {
 		try {
 			if (src.startsWith(UrlProtocol.DATA)) return true;
+
 			const url = new URL(src);
+
 			return url.protocol === UrlProtocol.HTTPS;
 		} catch {
 			return false;
@@ -446,40 +447,29 @@ class MCPStore {
 
 		// 1. Prefer icon explicitly matching the current color scheme
 		const themedIcon = validIcons.find((icon) => icon.theme === preferredTheme);
-		if (themedIcon) return this.#proxyIconSrc(themedIcon.src);
+		if (themedIcon) return themedIcon.src;
 
 		// 2. Handle universal icons (no theme specified)
 		const universalIcons = validIcons.filter((icon) => !icon.theme);
 
 		if (universalIcons.length === EXPECTED_THEMED_ICON_PAIR_COUNT) {
 			// Heuristic: two theme-less icons → assume [0] = light, [1] = dark
-			return this.#proxyIconSrc(universalIcons[isDark ? 1 : 0].src);
+			return universalIcons[isDark ? 1 : 0].src;
 		}
 
 		if (universalIcons.length > 0) {
-			return this.#proxyIconSrc(universalIcons[0].src);
+			return universalIcons[0].src;
 		}
 
 		// 3. Last resort: use opposite-theme icon
-		return this.#proxyIconSrc(validIcons[0].src);
-	}
-
-	/**
-	 * Route an icon src through the CORS proxy if it's an HTTPS URL.
-	 * Data URIs are returned as-is.
-	 */
-	#proxyIconSrc(src: string): string {
-		if (src.startsWith('data:')) return src;
-		if (!this._proxyAvailable) return src;
-
-		return getProxiedUrlString(src);
+		return validIcons[0].src;
 	}
 
 	/**
 	 * Get icon URL for an MCP server by its ID.
-	 * Prefers the server's own icons (from MCP spec) and falls back
-	 * to Google's favicon service.
-	 * Returns null if server is not found.
+	 * Returns the best icon from the MCP server's `icons` array
+	 * (see MCP spec: spec.modelcontextprotocol.io).
+	 * Returns null if no icon is available.
 	 */
 	getServerFavicon(serverId: string): string | null {
 		const server = this.getServerById(serverId);
@@ -497,7 +487,39 @@ class MCPStore {
 			}
 		}
 
-		return getFaviconUrl(server.url, this._proxyAvailable);
+		// Fallback: try favicon from root domain
+		const fallbackUrl = this.#getServerFaviconFallback(server.url);
+		if (fallbackUrl) {
+			return fallbackUrl;
+		}
+
+		return null;
+	}
+
+	/**
+	 * Construct a fallback favicon URL from the MCP server URL.
+	 * e.g. https://mcp.exa.ai/mcp -> https://exa.ai/favicon.ico
+	 */
+	#getServerFaviconFallback(serverUrl: string): string | null {
+		try {
+			const url = new URL(serverUrl);
+			const rootDomain = extractRootDomain(url);
+			if (!rootDomain) return null;
+
+			const origin = `${url.protocol}//${rootDomain}`;
+			const candidates = ['favicon.ico', 'favicon.svg', 'favicon.png'];
+
+			for (const path of candidates) {
+				const faviconUrl = `${origin}/${path}`;
+				if (this.#isValidIconUri(faviconUrl)) {
+					return faviconUrl;
+				}
+			}
+		} catch {
+			// Invalid URL, return null
+		}
+
+		return null;
 	}
 
 	isAnyServerLoading(): boolean {

tools/server/webui/src/lib/utils/cors-proxy.ts

@@ -33,12 +33,3 @@ export function buildProxiedHeaders(headers: Record<string, string>): Record<str
 
 	return proxiedHeaders;
 }
-
-/**
- * Get a proxied URL string for use in fetch requests.
- * @param targetUrl - The original URL to proxy
- * @returns Proxied URL as string
- */
-export function getProxiedUrlString(targetUrl: string): string {
-	return buildProxiedUrl(targetUrl).href;
-}

tools/server/webui/src/lib/utils/favicon.ts

@@ -1,34 +0,0 @@
-/**
- * Favicon utility functions for extracting favicons from URLs.
- */
-
-import { getProxiedUrlString } from './cors-proxy';
-import {
-	GOOGLE_FAVICON_BASE_URL,
-	DEFAULT_FAVICON_SIZE,
-	DOMAIN_SEPARATOR,
-	ROOT_DOMAIN_MIN_PARTS
-} from '$lib/constants';
-
-/**
- * Gets a favicon URL for a given URL using Google's favicon service.
- * Returns null if the URL is invalid.
- *
- * @param urlString - The URL to get the favicon for
- * @returns The favicon URL or null if invalid
- */
-export function getFaviconUrl(urlString: string, useProxy = true): string | null {
-	try {
-		const url = new URL(urlString);
-		const hostnameParts = url.hostname.split(DOMAIN_SEPARATOR);
-		const rootDomain =
-			hostnameParts.length >= ROOT_DOMAIN_MIN_PARTS
-				? hostnameParts.slice(-ROOT_DOMAIN_MIN_PARTS).join(DOMAIN_SEPARATOR)
-				: url.hostname;
-
-		const googleFaviconUrl = `${GOOGLE_FAVICON_BASE_URL}?domain=${rootDomain}&sz=${DEFAULT_FAVICON_SIZE}`;
-		return useProxy ? getProxiedUrlString(googleFaviconUrl) : googleFaviconUrl;
-	} catch {
-		return null;
-	}
-}

tools/server/webui/src/lib/utils/index.ts

@@ -39,7 +39,10 @@ export { highlightCode, detectIncompleteCodeBlock, type IncompleteCodeBlock } fr
 export { setConfigValue, getConfigValue, configToParameterRecord } from './config-helpers';
 
 // CORS Proxy
-export { buildProxiedUrl, getProxiedUrlString, buildProxiedHeaders } from './cors-proxy';
+export { buildProxiedUrl, buildProxiedHeaders } from './cors-proxy';
+
+// URL utilities
+export { extractRootDomain, sanitizeExternalUrl } from './url';
 
 // Conversation utilities
 export { createMessageCountMap, getMessageCount } from './conversation-utils';
@@ -146,9 +149,6 @@ export { createBase64DataUrl } from './data-url';
 // Header utilities
 export { parseHeadersToArray, serializeHeaders } from './headers';
 
-// Favicon utilities
-export { getFaviconUrl } from './favicon';
-
 // Agentic content utilities (structured section derivation)
 export {
 	deriveAgenticSections,

tools/server/webui/src/lib/utils/url.ts

@@ -0,0 +1,72 @@
+import { TWO_PART_PUBLIC_SUFFIXES, WILDCARD_PUBLIC_SUFFIXES } from '$lib/constants';
+import { UrlProtocol } from '$lib/enums';
+
+/**
+ * Check whether a hostname looks like an IPv4 or IPv6 address.
+ */
+function isIpAddress(hostname: string): boolean {
+	if (hostname.includes(':')) return true;
+
+	if (/^\d{1,3}(\.\d{1,3}){3}$/.test(hostname)) return true;
+
+	return false;
+}
+
+/**
+ * Extract the registrable root domain from a URL.
+ *
+ * @example
+ *   'mcp.example.com'           -> 'example.com'
+ *   'www.example.co.uk'         -> 'example.co.uk'
+ *   'bar.foo.nom.br'            -> 'bar.foo.nom.br'
+ *   '192.168.1.1'               -> null
+ *   'localhost'                 -> null
+ */
+export function extractRootDomain(url: URL): string | null {
+	const hostname = url.hostname.toLowerCase();
+	if (!hostname || isIpAddress(hostname)) return null;
+
+	const parts = hostname.split('.');
+
+	if (parts.length < 2) return null;
+
+	if (parts.length >= 3) {
+		const suffix2 = `${parts[parts.length - 2]}.${parts[parts.length - 1]}`;
+
+		if (TWO_PART_PUBLIC_SUFFIXES.has(suffix2)) {
+			return parts.slice(-3).join('.');
+		}
+	}
+
+	for (let i = 2; i <= parts.length; i++) {
+		const candidate = parts.slice(-i).join('.');
+
+		if (WILDCARD_PUBLIC_SUFFIXES.has(candidate)) {
+			if (parts.length === i + 1) {
+				return hostname;
+			}
+
+			return parts.slice(-(i + 2)).join('.');
+		}
+	}
+
+	return parts.slice(-2).join('.');
+}
+
+/**
+ * Sanitize an external URL string for safe use in an `<a href>`.
+ * Only allows http: and https: schemes. Returns `null` for anything else.
+ */
+export function sanitizeExternalUrl(raw: string): string | null {
+	try {
+		const url = new URL(raw);
+
+		if (url.protocol !== UrlProtocol.HTTP && url.protocol !== UrlProtocol.HTTPS) {
+			return null;
+		}
+
+		return url.href;
+	} catch {
+		return null;
+	}
+}

vendor/cpp-httplib/httplib.cpp

@@ -2506,6 +2506,10 @@ void get_remote_ip_and_port(socket_t sock, std::string &ip, int &port) {
   }
 }
 
+// Recursive form retained so operator""_t below can compute hashes for
+// switch-case labels at compile time (C++11 constexpr forbids loops). Do not
+// call from runtime paths with arbitrary-length inputs — use str2tag()
+// instead, which is iterative and stack-safe.
 constexpr unsigned int str2tag_core(const char *s, size_t l,
                                            unsigned int h) {
   return (l == 0)
@@ -2519,7 +2523,16 @@ constexpr unsigned int str2tag_core(const char *s, size_t l,
 }
 
 unsigned int str2tag(const std::string &s) {
-  return str2tag_core(s.data(), s.size(), 0);
+  // Iterative form of str2tag_core: the recursive constexpr version is kept
+  // for compile-time UDL evaluation of short string literals, but at runtime
+  // we may receive arbitrarily long inputs (e.g. fuzzed Content-Type) that
+  // would blow the stack with one frame per character.
+  unsigned int h = 0;
+  for (auto c : s) {
+    h = (((std::numeric_limits<unsigned int>::max)() >> 6) & h * 33) ^
+        static_cast<unsigned char>(c);
+  }
+  return h;
 }
 
 namespace udl {
@@ -9777,7 +9790,15 @@ bool ClientImpl::process_request(Stream &strm, Request &req,
           output_error_log(error, &req);
           return false;
         }
-        res.body.reserve(static_cast<size_t>(len));
+        // Cap the reservation by payload_max_length_ to avoid OOM when a
+        // hostile or malformed server sends an enormous Content-Length.
+        // The actual body read below is bounded by payload_max_length_,
+        // so reserving more than that is never useful.
+        auto reserve_len = static_cast<size_t>(len);
+        if (payload_max_length_ > 0 && reserve_len > payload_max_length_) {
+          reserve_len = payload_max_length_;
+        }
+        res.body.reserve(reserve_len);
       }
     }
 

vendor/cpp-httplib/httplib.h

@@ -8,8 +8,8 @@
 #ifndef CPPHTTPLIB_HTTPLIB_H
 #define CPPHTTPLIB_HTTPLIB_H
 
-#define CPPHTTPLIB_VERSION "0.43.2"
-#define CPPHTTPLIB_VERSION_NUM "0x002b02"
+#define CPPHTTPLIB_VERSION "0.43.3"
+#define CPPHTTPLIB_VERSION_NUM "0x002b03"
 
 #ifdef _WIN32
 #if defined(_WIN32_WINNT) && _WIN32_WINNT < 0x0A00