spec : auto-detect mtp draft model type (#27005)

* metal: add TQ2_0 support

Add support for the GGML_TYPE_TQ2_0 (ternary, 2 bits per element) type in
the Metal backend.

Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731

* cont : optimize mul_mv kernel

- float ops over integer ops
- precalculate sums
- hoist coef out of the inner loop
- contiguous y loads

llama.cpp:DeepSeek-v4-Flash-0731

common/arg.cpp

@@ -561,6 +561,15 @@ void common_models_handler_apply(common_models_handler & handler, common_params
         }
     }
 
+    // infer the speculative type from the draft GGUF metadata when none is requested
+    // note: reads only the first split - sharded drafts need an explicit --spec-type
+    if (spec_types_is_default(params) && !params.speculative.draft.mparams.path.empty()) {
+        const auto types_gguf = common_speculative_types_from_gguf(params.speculative.draft.mparams.path);
+        if (!types_gguf.empty()) {
+            params.speculative.types = types_gguf;
+        }
+    }
+
     // when a sidecar type is requested, the draft repo resolves to its sidecar instead of a full model
     const bool spec_sidecar_found = !plan_spec.mtp.local_path.empty() ||
                                     !plan_spec.dflash.local_path.empty() ||

common/speculative.cpp

@@ -2,6 +2,7 @@
 
 #include "common.h"
 #include "ggml.h"
+#include "ggml-cpp.h"
 #include "llama.h"
 #include "log.h"
 #include "ngram-cache.h"
@@ -912,6 +913,9 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
 
     std::vector<common_sampler_ptr> smpls;
 
+    // backend sampler chain per seq, attached to ctx_dft
+    std::vector<llama_sampler *> backend_chains;
+
     int32_t n_embd_dec = 0;  // draft hidden size
     int32_t n_embd_enc = 0;  // target_layer_ids_n * target_hidden_size
     int32_t n_embd_tgt = 0;  // target model hidden size
@@ -985,6 +989,22 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
             s.reset(common_sampler_init(model_dft, sparams));
         }
 
+        // offload draft sampling to the backend
+        backend_chains.assign(n_seq, nullptr);
+        if (this->params.backend_sampling) {
+            for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) {
+                llama_sampler * chain = llama_sampler_chain_init(llama_sampler_chain_default_params());
+                llama_sampler_chain_add(chain, llama_sampler_init_top_k(10));
+
+                if (!llama_set_sampler(ctx_dft, seq_id, chain)) {
+                    SPC_WRN("backend offload failed for seq_id=%d; using CPU sampler\n", (int) seq_id);
+                    llama_sampler_free(chain);
+                    chain = nullptr;
+                }
+                backend_chains[seq_id] = chain;
+            }
+        }
+
         // turn on extraction of the target layers' input embeddings
         for (uint32_t k = 0; k < target_layer_ids_n; ++k) {
             llama_set_embeddings_layer_inp(ctx_tgt, (uint32_t) target_layer_ids[k], true);
@@ -995,6 +1015,18 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
     }
 
     ~common_speculative_impl_draft_dflash() override {
+        auto * ctx_dft = this->params.ctx_dft;
+        for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) backend_chains.size(); ++seq_id) {
+            if (backend_chains[seq_id] == nullptr) {
+                continue;
+            }
+            if (ctx_dft) {
+                llama_set_sampler(ctx_dft, seq_id, nullptr);
+            }
+            llama_sampler_free(backend_chains[seq_id]);
+        }
+        backend_chains.clear();
+
         llama_batch_free(batch);
         llama_batch_free(batch_inject);
     }
@@ -2196,6 +2228,43 @@ common_speculative_type common_speculative_type_from_name(const std::string & na
     return it->second;
 }
 
+std::vector<common_speculative_type> common_speculative_types_from_gguf(const std::string & path) {
+    struct gguf_init_params gguf_params = {
+        /* .no_alloc = */ true,
+        /* .ctx      = */ nullptr,
+    };
+
+    gguf_context_ptr gguf_ctx(gguf_init_from_file(path.c_str(), gguf_params));
+    if (!gguf_ctx) {
+        return {};
+    }
+
+    const int64_t arch_id = gguf_find_key(gguf_ctx.get(), "general.architecture");
+    if (arch_id < 0 || gguf_get_kv_type(gguf_ctx.get(), arch_id) != GGUF_TYPE_STRING) {
+        return {};
+    }
+
+    const std::string arch = gguf_get_val_str(gguf_ctx.get(), arch_id);
+    if (arch != "dflash") {
+        const uint32_t block_count = gguf_get_val_u32(gguf_ctx.get(), gguf_find_key(gguf_ctx.get(), (arch + ".block_count").c_str()));
+
+        if (gguf_find_tensor(gguf_ctx.get(), ("blk." + std::to_string(block_count - 1) + ".nextn.eh_proj.weight").c_str()) >= 0) {
+            return { COMMON_SPECULATIVE_TYPE_DRAFT_MTP };
+        }
+
+        return {};
+    }
+
+    // the Markov head distinguishes draft-dspark from draft-dflash
+    const auto type = gguf_find_tensor(gguf_ctx.get(), "markov_w1.weight") >= 0
+                    ? COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK
+                    : COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH;
+
+    SPC_INF("auto-detected speculative type '%s' from the draft model metadata\n", common_speculative_type_to_str(type).c_str());
+
+    return { type };
+}
+
 static uint32_t common_get_enabled_speculative_configs(const std::vector<common_speculative_type> & configs) {
     uint32_t result = 0;
     for (size_t i = 0; i < configs.size(); i++) {
@@ -2263,6 +2332,23 @@ common_params common_base_params_to_speculative(const common_params & params) {
     result.n_outputs_max = params.n_parallel;
     result.n_outputs_max_per_seq = 1;
 
+    // dflash/dspark decode the whole noise block in a single pass and sample every block position on the backend
+    // TODO: refactor such properties to be announced by the speculative types
+    //       something like `struct common_speculative_type_props common_speculative_type_get_props(...);`
+    const bool has_block_draft = std::any_of(
+        params.speculative.types.begin(), params.speculative.types.end(),
+        [](common_speculative_type t) {
+            return t == COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH || t == COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK;
+        });
+    if (has_block_draft) {
+        // per-seq output positions: DFlash decodes anchor + n_max masks (n_max + 1); DSpark n_max -> +1 covers both
+        const int32_t per_seq = std::max(1, params_spec.n_max + 1);
+        result.n_outputs_max = params.n_parallel * per_seq;
+        if (params_spec.backend_sampling) {
+            result.n_outputs_max_per_seq = per_seq;
+        }
+    }
+
     return result;
 }
 

common/speculative.h

@@ -14,6 +14,9 @@ const char * common_speculative_all_types_str();
 // parse user provided types
 std::vector<enum common_speculative_type> common_speculative_types_from_names(const std::vector<std::string> & names);
 
+// infer the spec types from the GGUF metadata of a draft model; empty if unknown
+std::vector<enum common_speculative_type> common_speculative_types_from_gguf(const std::string & path);
+
 // convert string to type
 enum common_speculative_type common_speculative_type_from_name(const std::string & name);
 

docs/backend/SYCL.md

@@ -804,6 +804,7 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm
 | GGML_SYCL_MKL_FA_DEBUG | 0 (default) or 1 | Enable per-call diagnostic logging for MKL flash attention: GEMM/softmax timings, interleaved-head detection, and buffer memory usage. |
 | GGML_SYCL_MKL_FA_DIAG | 0 (default) or 1 | Enable output fingerprinting for MKL flash attention. Dumps the first 64 float output values for the first 6 FA calls with n_kv ≥ 1024, labeled with kernel type (MKL/TILE/VEC) for cross-kernel comparison. |
 | GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute (currently top-k MoE gating). |
+| GGML_SYCL_ENABLE_ESIMD | 0 or 1 (default)| Enable ESIMD kernels when available. |
 | ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.<br>Recommended to use when --split-mode = layer |
 | UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS | 0 (default) or 1 | Allow SYCL/Unified Runtime Level Zero device allocations larger than 4 GiB. llama.cpp's direct Level Zero allocation path requests the relaxed maximum-size limit itself when GGML_SYCL_ENABLE_LEVEL_ZERO=1. |
 | GGML_SYCL_USM_SYSTEM | 0 (default) or 1 | Enable experimental support for [USM system allocations](https://github.khronos.org/SYCL_Reference/iface/usm_basic_concept.html#system-allocations) for large GPU buffers. This requires enough host memory for model weights and caches, an Intel Xe2+ GPU such as BMG or newer and supported on Linux only, with CONFIG_DRM_XE_GPUSVM enabled. |

ggml/src/ggml-cpu/ops.cpp

@@ -8941,7 +8941,7 @@ static void ggml_compute_forward_flash_attn_ext_tiled(
             for (int tk = 0; tk < kv_tile; tk++) {
                 const char * v_data = (const char *)v->data + (ic + tk)*nbv1 + iv2*nbv2 + iv3*nbv3;
                 if (kv_type == GGML_TYPE_F16) {
-                    ggml_fp16_to_fp32_row((const ggml_fp16_t *)v_data, V32 + tk * DV, DV);
+                    ggml_cpu_fp16_to_fp32((const ggml_fp16_t *)v_data, V32 + tk * DV, DV);
                 } else {
                     memcpy(V32 + tk * DV, v_data, DV * sizeof(float));
                 }

ggml/src/ggml-hip/CMakeLists.txt

@@ -126,9 +126,6 @@ if (GGML_HIP_EXPORT_METRICS)
     set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -Rpass-analysis=kernel-resource-usage --save-temps")
 endif()
 
-# Fast math for HIP, like CUDA's -use_fast_math. Not -ffast-math: that implies -ffinite-math-only, which breaks ggml's INFINITY masking and produces NaNs.
-set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -funsafe-math-optimizations")
-
 if (NOT GGML_CUDA_FA)
     add_compile_definitions(GGML_CUDA_NO_FA)
 endif()

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

@@ -953,6 +953,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
                 nr0 = N_R0_IQ4_XS;
                 smem = 32*sizeof(float);
             } break;
+        case GGML_TYPE_TQ2_0:
+            {
+                nsg = N_SG_TQ2_0;
+                nr0 = N_R0_TQ2_0;
+            } break;
         default:
             {
                 GGML_LOG_ERROR("Asserting on type %d\n", (int) tsrc0);
@@ -1182,6 +1187,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
                 nr0 = N_R0_IQ4_XS;
                 smem = 32*sizeof(float);
             } break;
+        case GGML_TYPE_TQ2_0:
+            {
+                nsg = N_SG_TQ2_0;
+                nr0 = N_R0_TQ2_0;
+            } break;
         default:
             {
                 GGML_LOG_ERROR("Asserting on type %d\n", (int)op->src[2]->type);

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

@@ -1407,6 +1407,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
                            case GGML_TYPE_Q5_0:
                            case GGML_TYPE_Q5_1:
                            case GGML_TYPE_IQ4_NL:
+                           case GGML_TYPE_TQ2_0:
                            case GGML_TYPE_I32:
                                 return true;
                            default:
@@ -1435,6 +1436,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
                     case GGML_TYPE_Q5_0:
                     case GGML_TYPE_Q5_1:
                     case GGML_TYPE_Q8_0:
+                    case GGML_TYPE_TQ2_0:
                         switch (op->type) {
                             case GGML_TYPE_F32:
                             case GGML_TYPE_F16:
@@ -1470,6 +1472,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
                     case GGML_TYPE_Q5_0:
                     case GGML_TYPE_Q5_1:
                     case GGML_TYPE_IQ4_NL:
+                    case GGML_TYPE_TQ2_0:
                         return true;
                     default:
                         return false;

ggml/src/ggml-metal/ggml-metal-impl.h

@@ -87,6 +87,9 @@
 #define N_R0_IQ4_XS 2
 #define N_SG_IQ4_XS 2
 
+#define N_R0_TQ2_0 4
+#define N_SG_TQ2_0 2
+
 // function constants offsets
 #define FC_FLASH_ATTN_EXT_PAD          100
 #define FC_FLASH_ATTN_EXT_BLK          200

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

@@ -468,6 +468,34 @@ void quantize_iq4_nl(device const float * src, device block_iq4_nl & dst) {
     dst.d = sumq2 > 0 ? sumqx/sumq2 : d;
 }
 
+void quantize_tq2_0(device const float * src, device block_tq2_0 & dst) {
+#pragma METAL fp math_mode(safe)
+    float amax = 0.0f; // absolute max
+
+    for (int j = 0; j < QK_K; j++) {
+        const float v = src[j];
+        amax = MAX(amax, fabs(v));
+    }
+
+    const float d = amax;
+    const float id = d ? 1.0f/d : 0.0f;
+
+    dst.d = (half) d;
+
+    for (int j = 0; j < QK_K/4; j += 32) {
+        for (int m = 0; m < 32; ++m) {
+            uint8_t q = 0;
+            for (int n = 0; n < 4; ++n) {
+                // -1, 0, 1 -> 0, 1, 2
+                int xi = (int)round(src[m + n*32] * id) + 1;
+                q += (uint8_t)((xi & 3) << (2*n));
+            }
+            dst.qs[j + m] = q;
+        }
+        src += 4*32;
+    }
+}
+
 template <typename type4x4>
 void dequantize_q4_1(device const block_q4_1 * xb, short il, thread type4x4 & reg) {
     device const uint16_t * qs = ((device const uint16_t *)xb + 2);
@@ -1021,6 +1049,25 @@ void dequantize_iq4_xs(device const block_iq4_xs * xb, short il, thread type4x4
     }
 }
 
+template <typename type4x4>
+void dequantize_tq2_0(device const block_tq2_0 * xb, short il, thread type4x4 & reg) {
+    device const uint8_t * qs = xb->qs;
+    const float d = xb->d;
+
+    float4x4 reg_f;
+
+    // 2 bits per element, 4 elements per byte, 128 elements per 32-byte group
+    const short base = il * 16;
+    for (int k = 0; k < 16; k++) {
+        const int i = base + k;
+        const int byte = ((i >> 7) & 1) * 32 + (i & 31);
+        const int l = (i >> 5) & 3;
+        reg_f[k/4][k%4] = d * (float)(((qs[byte] >> (2*l)) & 3) - 1);
+    }
+
+    reg = (type4x4) reg_f;
+}
+
 enum ggml_sort_order {
     GGML_SORT_ORDER_ASC,
     GGML_SORT_ORDER_DESC,
@@ -8001,6 +8048,7 @@ template [[host_name("kernel_cpy_f32_q4_1")]]   kernel cpy_f_q_t kernel_cpy_f32_
 template [[host_name("kernel_cpy_f32_q5_0")]]   kernel cpy_f_q_t kernel_cpy_f32_q<QK5_0,  block_q5_0,   quantize_q5_0>;
 template [[host_name("kernel_cpy_f32_q5_1")]]   kernel cpy_f_q_t kernel_cpy_f32_q<QK5_1,  block_q5_1,   quantize_q5_1>;
 template [[host_name("kernel_cpy_f32_iq4_nl")]] kernel cpy_f_q_t kernel_cpy_f32_q<QK4_NL, block_iq4_nl, quantize_iq4_nl>;
+template [[host_name("kernel_cpy_f32_tq2_0")]]  kernel cpy_f_q_t kernel_cpy_f32_q<QK_K,   block_tq2_0,  quantize_tq2_0>;
 
 template<typename T4x4, typename block_q, short nl, void (*dequantize_func)(device const block_q *, short, thread T4x4 &)>
 kernel void kernel_cpy_q_f32(
@@ -8048,6 +8096,8 @@ template [[host_name("kernel_cpy_q5_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32<
 template [[host_name("kernel_cpy_q5_1_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32<float4x4, block_q5_1, 2, dequantize_q5_1>;
 template [[host_name("kernel_cpy_q8_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32<float4x4, block_q8_0, 2, dequantize_q8_0>;
 
+template [[host_name("kernel_cpy_tq2_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32<float4x4, block_tq2_0, QK_NL, dequantize_tq2_0>;
+
 template [[host_name("kernel_cpy_q1_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32<half4x4, block_q1_0, 8, dequantize_q1_0>;
 template [[host_name("kernel_cpy_q2_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32<half4x4, block_q2_0, 4, dequantize_q2_0>;
 template [[host_name("kernel_cpy_q4_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32<half4x4, block_q4_0, 2, dequantize_q4_0>;
@@ -8056,6 +8106,8 @@ template [[host_name("kernel_cpy_q5_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32<
 template [[host_name("kernel_cpy_q5_1_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32<half4x4, block_q5_1, 2, dequantize_q5_1>;
 template [[host_name("kernel_cpy_q8_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32<half4x4, block_q8_0, 2, dequantize_q8_0>;
 
+template [[host_name("kernel_cpy_tq2_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32<half4x4, block_tq2_0, QK_NL, dequantize_tq2_0>;
+
 template<typename T>
 kernel void kernel_concat(
         constant ggml_metal_kargs_concat & args,
@@ -9822,6 +9874,121 @@ kernel void kernel_mul_mv_mxfp4_f32(
     kernel_mul_mv_mxfp4_f32_impl<N_R0_MXFP4, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
 }
 
+template<int nr0, typename args_t>
+void kernel_mul_mv_tq2_0_f32_impl(
+        args_t args,
+        device const char * src0,
+        device const char * src1,
+        device       char * dst,
+        threadgroup  char * shmem,
+        uint3  tgpig,
+        ushort tiisg,
+        ushort sgitg) {
+    const short NSG = FC_mul_mv_nsg;
+
+    const int nb = args.ne00/QK_K;
+
+    const int r0 = tgpig.x;
+    const int r1 = tgpig.y;
+    const int im = tgpig.z;
+
+    const int first_row = (r0 * NSG + sgitg) * nr0;
+
+    const uint i12 = im%FC_mul_mv_ne12;
+    const uint i13 = im/FC_mul_mv_ne12;
+
+    const uint64_t offset1 =        r1*args.nb11 + (i12        )*args.nb12 + (i13        )*args.nb13;
+
+    device const float * y = (device const float *) (src1 + offset1);
+
+    device const block_tq2_0 * ax[nr0];
+    for (int row = 0; row < nr0; ++row) {
+        const uint64_t offset0 = (first_row + row)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
+        ax[row] = (device const block_tq2_0 *) ((device char *) src0 + offset0);
+    }
+
+    float sumf[nr0] = {0.f};
+
+    // 8 threads per block, NBLOCK blocks per pass, 2 halves per block per pass
+    constexpr short NBLOCK = 4;
+
+    constexpr short NB = N_SIMDWIDTH/NBLOCK; // threads per block
+
+    const short blk = tiisg / NB;    // 0..NBLOCK-1, block handled by this thread
+    const short htg = tiisg % NB;    // 0..NB-1, thread within block (0..7)
+
+    // byte and y base offsets within the block (32 elements per thread, 4 per byte)
+    device const float4 * yb4 = (device const float4 *)(y + 4*htg + blk*QK_K);
+
+    // hoisted per-byte coefficients (from y) and total y-sum, shared across rows
+    // ref: https://github.com/ggml-org/llama.cpp/pull/26980
+    float4 coef[4];
+
+    for (int ib = blk; ib < nb; ib += NBLOCK) {
+        FOR_UNROLL (short h0 = 0; h0 < 2; ++h0) {
+            const float4 y0 = yb4[ 0 + 32*h0];
+            const float4 y1 = yb4[ 8 + 32*h0];
+            const float4 y2 = yb4[16 + 32*h0];
+            const float4 y3 = yb4[24 + 32*h0];
+
+            float sumy = 0.f;
+            FOR_UNROLL (short j = 0; j < 4; ++j) {
+                coef[j] = float4(
+                        y0[j],
+                        y1[j] - 4.0f*y0[j],
+                        y2[j] - 4.0f*y1[j],
+                        y3[j] - 4.0f*y2[j]);
+
+                sumy += (y0[j] + y1[j]) + (y2[j] + y3[j]);
+            }
+
+            FOR_UNROLL (short row = 0; row < nr0; ++row) {
+                device const block_tq2_0 & xb = ax[row][ib];
+                device const uchar * qs = xb.qs + 4*htg + 32*h0;
+
+                float sum = -sumy;
+                FOR_UNROLL (short j = 0; j < 4; ++j) {
+                    // express the 2-bit field shifts (v>>2, v>>4, v>>6) as float floor ops
+                    const float v = (float)qs[j];
+
+                    const float f0 = v;
+                    const float f1 = floor(v*0.25f);    // v>>2
+                    const float f2 = floor(v*0.0625);   // v>>4
+                    const float f3 = floor(v*0.015625); // v>>6
+
+                    sum += coef[j][0]*f0 + coef[j][1]*f1 + coef[j][2]*f2 + coef[j][3]*f3;
+                }
+
+                sumf[row] += xb.d * sum;
+            }
+        }
+
+        yb4 += QK_K * NBLOCK / 4;
+    }
+
+    device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
+
+    for (int row = 0; row < nr0; ++row) {
+        const float tot = simd_sum(sumf[row]);
+        if (tiisg == 0 && first_row + row < args.ne01) {
+            dst_f32[first_row + row] = tot;
+        }
+    }
+}
+
+[[host_name("kernel_mul_mv_tq2_0_f32")]]
+kernel void kernel_mul_mv_tq2_0_f32(
+        constant ggml_metal_kargs_mul_mv & args,
+        device const char * src0,
+        device const char * src1,
+        device       char * dst,
+        uint3  tgpig[[threadgroup_position_in_grid]],
+        ushort tiisg[[thread_index_in_simdgroup]],
+        ushort sgitg[[simdgroup_index_in_threadgroup]]) {
+
+    kernel_mul_mv_tq2_0_f32_impl<N_R0_TQ2_0, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg);
+}
+
 template<typename block_q, short nl, void (*dequantize_func)(device const block_q *, short, thread float4x4 &)>
 kernel void kernel_get_rows_q(
         constant ggml_metal_kargs_get_rows & args,
@@ -9915,6 +10082,38 @@ template [[host_name("kernel_get_rows_iq1_s")]]   kernel get_rows_q_t kernel_get
 template [[host_name("kernel_get_rows_iq1_m")]]   kernel get_rows_q_t kernel_get_rows_q<block_iq1_m,   QK_NL, dequantize_iq1_m>;
 template [[host_name("kernel_get_rows_iq4_nl")]]  kernel get_rows_q_t kernel_get_rows_q<block_iq4_nl,  2,     dequantize_iq4_nl>;
 template [[host_name("kernel_get_rows_iq4_xs")]]  kernel get_rows_q_t kernel_get_rows_q<block_iq4_xs,  QK_NL, dequantize_iq4_xs>;
+template [[host_name("kernel_get_rows_tq2_0")]]   kernel get_rows_q_t kernel_get_rows_q<block_tq2_0,   QK_NL, dequantize_tq2_0>;
+
+template<typename TS, typename TI, short QK, typename block_q, void (*quantize_func)(device const float *, device block_q &)>
+kernel void kernel_set_rows_q(
+        constant ggml_metal_kargs_set_rows & args,
+        device const  void * src0,
+        device const  void * src1,
+        device       float * dst,
+        uint3                tgpig[[threadgroup_position_in_grid]],
+        uint                 tiitg[[thread_index_in_threadgroup]],
+        uint3                tptg [[threads_per_threadgroup]]) {
+    const int32_t i03 = tgpig.z;
+    const int32_t i02 = tgpig.y;
+
+    const int32_t i12 = i03%args.ne12;
+    const int32_t i11 = i02%args.ne11;
+
+    const int32_t i01 = tgpig.x*tptg.y + tiitg/tptg.x;
+    if (i01 >= args.ne01) {
+        return;
+    }
+
+    const int32_t i10 = i01;
+    const TI      i1  = ((const device TI *) ((const device char *) src1 + i10*args.nb10 + i11*args.nb11 + i12*args.nb12))[0];
+
+          device block_q * dst_row = (      device block_q *) ((      device char *) dst  +  i1*args.nb1  + i02*args.nb2  + i03*args.nb3);
+    const device TS      * src_row = (const device TS      *) ((const device char *) src0 + i01*args.nb01 + i02*args.nb02 + i03*args.nb03);
+
+    for (int ind = tiitg%tptg.x; ind < args.nk0; ind += tptg.x) {
+        quantize_func(src_row + QK*ind, dst_row[ind]);
+    }
+}
 
 template<typename TS, typename TI, typename block_q, void (*quantize_func)(device const float *, device block_q &)>
 kernel void kernel_set_rows_q32(
@@ -10011,6 +10210,11 @@ template [[host_name("kernel_set_rows_f32_i32_q5_1")]]   kernel set_rows_q32_t k
 template [[host_name("kernel_set_rows_f32_i64_iq4_nl")]] kernel set_rows_q32_t kernel_set_rows_q32<float, int64_t, block_iq4_nl, quantize_iq4_nl>;
 template [[host_name("kernel_set_rows_f32_i32_iq4_nl")]] kernel set_rows_q32_t kernel_set_rows_q32<float, int32_t, block_iq4_nl, quantize_iq4_nl>;
 
+typedef decltype(kernel_set_rows_q<float, int64_t, QK_K, block_tq2_0, quantize_tq2_0>) set_rows_qK_t;
+
+template [[host_name("kernel_set_rows_f32_i64_tq2_0")]]  kernel set_rows_qK_t kernel_set_rows_q<float, int64_t, QK_K, block_tq2_0, quantize_tq2_0>;
+template [[host_name("kernel_set_rows_f32_i32_tq2_0")]]  kernel set_rows_qK_t kernel_set_rows_q<float, int32_t, QK_K, block_tq2_0, quantize_tq2_0>;
+
 kernel void kernel_diag_f32(
         constant ggml_metal_kargs_diag & args,
         device   const char * src0,
@@ -10786,6 +10990,7 @@ template [[host_name("kernel_mul_mm_iq1_s_f32")]]   kernel mul_mm_t kernel_mul_m
 template [[host_name("kernel_mul_mm_iq1_m_f32")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  float, float2x4>;
 template [[host_name("kernel_mul_mm_iq4_nl_f32")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  float, float2x4>;
 template [[host_name("kernel_mul_mm_iq4_xs_f32")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  float, float2x4>;
+template [[host_name("kernel_mul_mm_tq2_0_f32")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_tq2_0,   QK_NL, dequantize_tq2_0,   float,  float4x4,  float, float2x4>;
 
 template [[host_name("kernel_mul_mm_f32_f16")]]     kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   float4x4,      1,     dequantize_f32,     float,  float4x4,  half, half2x4>;
 template [[host_name("kernel_mul_mm_f16_f16")]]     kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   half4x4,       1,     dequantize_f16,     half,   half4x4,   half, half2x4>;
@@ -10811,6 +11016,7 @@ template [[host_name("kernel_mul_mm_iq1_s_f16")]]   kernel mul_mm_t kernel_mul_m
 template [[host_name("kernel_mul_mm_iq1_m_f16")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  half, half2x4>;
 template [[host_name("kernel_mul_mm_iq4_nl_f16")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  half, half2x4>;
 template [[host_name("kernel_mul_mm_iq4_xs_f16")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  half, half2x4>;
+template [[host_name("kernel_mul_mm_tq2_0_f16")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_tq2_0,   QK_NL, dequantize_tq2_0,   float,  float4x4,  half, half2x4>;
 
 //
 // indirect matrix-matrix multiplication
@@ -10845,6 +11051,7 @@ template [[host_name("kernel_mul_mm_id_iq1_s_f32")]]   kernel mul_mm_id kernel_m
 template [[host_name("kernel_mul_mm_id_iq1_m_f32")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  float, float2x4>;
 template [[host_name("kernel_mul_mm_id_iq4_nl_f32")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  float, float2x4>;
 template [[host_name("kernel_mul_mm_id_iq4_xs_f32")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  float, float2x4>;
+template [[host_name("kernel_mul_mm_id_tq2_0_f32")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_tq2_0,   QK_NL, dequantize_tq2_0,   float,  float4x4,  float, float2x4>;
 
 template [[host_name("kernel_mul_mm_id_f32_f16")]]     kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   float4x4,      1,     dequantize_f32,     float,  float4x4,  half, half2x4>;
 template [[host_name("kernel_mul_mm_id_f16_f16")]]     kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   half4x4,       1,     dequantize_f16,     half,   half4x4,   half, half2x4>;
@@ -10870,6 +11077,7 @@ template [[host_name("kernel_mul_mm_id_iq1_s_f16")]]   kernel mul_mm_id kernel_m
 template [[host_name("kernel_mul_mm_id_iq1_m_f16")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  half, half2x4>;
 template [[host_name("kernel_mul_mm_id_iq4_nl_f16")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  half, half2x4>;
 template [[host_name("kernel_mul_mm_id_iq4_xs_f16")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  half, half2x4>;
+template [[host_name("kernel_mul_mm_id_tq2_0_f16")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_tq2_0,   QK_NL, dequantize_tq2_0,   float,  float4x4,  half, half2x4>;
 
 //
 // matrix-vector multiplication
@@ -11027,6 +11235,7 @@ template [[host_name("kernel_mul_mv_id_iq3_s_f32")]]   kernel kernel_mul_mv_id_t
 template [[host_name("kernel_mul_mv_id_iq2_s_f32")]]   kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_s_f32_impl  <N_R0_IQ2_S>>>;
 template [[host_name("kernel_mul_mv_id_iq4_nl_f32")]]  kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq4_nl_f32_impl <N_R0_IQ4_NL>>>;
 template [[host_name("kernel_mul_mv_id_iq4_xs_f32")]]  kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq4_xs_f32_impl <N_R0_IQ4_XS>>>;
+template [[host_name("kernel_mul_mv_id_tq2_0_f32")]]   kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_tq2_0_f32_impl  <N_R0_TQ2_0>>>;
 
 kernel void kernel_pool_2d_max_f32(
         constant    ggml_metal_kargs_pool_2d & args,

ggml/src/ggml-sycl/common.hpp

@@ -61,6 +61,7 @@ void ggml_sycl_host_free(void* ptr);
 extern int g_ggml_sycl_debug;
 extern int g_ggml_sycl_enable_optimize;
 extern int g_ggml_sycl_enable_fusion;
+extern int g_ggml_sycl_enable_esimd;
 extern int g_ggml_sycl_prioritize_dmmv;
 extern int g_ggml_sycl_enable_flash_attention;
 extern int g_ggml_sycl_dev2dev_memcpy;

ggml/src/ggml-sycl/concat.cpp

@@ -184,8 +184,8 @@ void concat_impl_sycl(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
             const size_t size0 = ggml_nbytes(src0);
             const size_t size1 = ggml_nbytes(src1);
 
-            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0).wait()));
-            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d + size0 / type_size, src1_d, size1).wait()));
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d + size0 / type_size, src1_d, size1)));
         }
     } else {
         concat_T_sycl_non_cont<T>(stream, (const char *) src0->data, (const char *) src1->data, (char *) dst->data,
@@ -196,6 +196,270 @@ void concat_impl_sycl(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
     }
 }
 
+static void concat_impl_q4_0_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+    const ggml_tensor * src0   = dst->src[0];
+    const ggml_tensor * src1   = dst->src[1];
+    queue_ptr           stream = ctx.stream();
+
+    const int32_t dim = ((int32_t *) dst->op_params)[0];
+
+    GGML_ASSERT(src0->type == GGML_TYPE_Q4_0);
+    GGML_ASSERT(src1->type == GGML_TYPE_Q4_0);
+    GGML_ASSERT(dst->type == GGML_TYPE_Q4_0);
+    GGML_ASSERT(src0->ne[0] % QK4_0 == 0);
+    GGML_ASSERT(src1->ne[0] % QK4_0 == 0);
+    GGML_ASSERT(dst->ne[0] % QK4_0 == 0);
+
+    const int ne00_blk = src0->ne[0] / QK4_0;
+    const int ne0_blk  = dst->ne[0] / QK4_0;
+
+    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
+        const block_q4_0 * src0_d = (const block_q4_0 *) src0->data;
+        const block_q4_0 * src1_d = (const block_q4_0 *) src1->data;
+        block_q4_0 * dst_d        = (block_q4_0 *) dst->data;
+        const size_t type_size    = sizeof(block_q4_0);
+
+        if (dim != 3) {
+            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
+                concat_T_sycl<block_q4_0>(
+                    src0_d + i3 * (src0->nb[3] / type_size),
+                    src1_d + i3 * (src1->nb[3] / type_size),
+                    dst_d + i3 * (dst->nb[3] / type_size),
+                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
+                    dst->ne[1], dst->ne[2], dim, stream);
+            }
+        } else {
+            const size_t size0 = ggml_nbytes(src0);
+            const size_t size1 = ggml_nbytes(src1);
+
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
+        }
+    } else {
+        concat_T_sycl_non_cont<block_q4_0>(
+            stream, (const char *) src0->data, (const char *) src1->data,
+            (char *) dst->data,
+            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
+            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
+            src1->ne[0] / QK4_0, src1->ne[1], src1->ne[2], src1->ne[3],
+            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
+            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
+            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
+    }
+}
+
+static void concat_impl_q4_1_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+    const ggml_tensor * src0   = dst->src[0];
+    const ggml_tensor * src1   = dst->src[1];
+    queue_ptr           stream = ctx.stream();
+
+    const int32_t dim = ((int32_t *) dst->op_params)[0];
+
+    GGML_ASSERT(src0->type == GGML_TYPE_Q4_1);
+    GGML_ASSERT(src1->type == GGML_TYPE_Q4_1);
+    GGML_ASSERT(dst->type == GGML_TYPE_Q4_1);
+    GGML_ASSERT(src0->ne[0] % QK4_1 == 0);
+    GGML_ASSERT(src1->ne[0] % QK4_1 == 0);
+    GGML_ASSERT(dst->ne[0] % QK4_1 == 0);
+
+    const int ne00_blk = src0->ne[0] / QK4_1;
+    const int ne0_blk  = dst->ne[0] / QK4_1;
+
+    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
+        const block_q4_1 * src0_d = (const block_q4_1 *) src0->data;
+        const block_q4_1 * src1_d = (const block_q4_1 *) src1->data;
+        block_q4_1 * dst_d        = (block_q4_1 *) dst->data;
+        const size_t type_size    = sizeof(block_q4_1);
+
+        if (dim != 3) {
+            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
+                concat_T_sycl<block_q4_1>(
+                    src0_d + i3 * (src0->nb[3] / type_size),
+                    src1_d + i3 * (src1->nb[3] / type_size),
+                    dst_d + i3 * (dst->nb[3] / type_size),
+                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
+                    dst->ne[1], dst->ne[2], dim, stream);
+            }
+        } else {
+            const size_t size0 = ggml_nbytes(src0);
+            const size_t size1 = ggml_nbytes(src1);
+
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
+        }
+    } else {
+        concat_T_sycl_non_cont<block_q4_1>(
+            stream, (const char *) src0->data, (const char *) src1->data,
+            (char *) dst->data,
+            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
+            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
+            src1->ne[0] / QK4_1, src1->ne[1], src1->ne[2], src1->ne[3],
+            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
+            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
+            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
+    }
+}
+
+static void concat_impl_q5_0_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+    const ggml_tensor * src0   = dst->src[0];
+    const ggml_tensor * src1   = dst->src[1];
+    queue_ptr           stream = ctx.stream();
+
+    const int32_t dim = ((int32_t *) dst->op_params)[0];
+
+    GGML_ASSERT(src0->type == GGML_TYPE_Q5_0);
+    GGML_ASSERT(src1->type == GGML_TYPE_Q5_0);
+    GGML_ASSERT(dst->type == GGML_TYPE_Q5_0);
+    GGML_ASSERT(src0->ne[0] % QK5_0 == 0);
+    GGML_ASSERT(src1->ne[0] % QK5_0 == 0);
+    GGML_ASSERT(dst->ne[0] % QK5_0 == 0);
+
+    const int ne00_blk = src0->ne[0] / QK5_0;
+    const int ne0_blk  = dst->ne[0] / QK5_0;
+
+    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
+        const block_q5_0 * src0_d = (const block_q5_0 *) src0->data;
+        const block_q5_0 * src1_d = (const block_q5_0 *) src1->data;
+        block_q5_0 * dst_d        = (block_q5_0 *) dst->data;
+        const size_t type_size    = sizeof(block_q5_0);
+
+        if (dim != 3) {
+            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
+                concat_T_sycl<block_q5_0>(
+                    src0_d + i3 * (src0->nb[3] / type_size),
+                    src1_d + i3 * (src1->nb[3] / type_size),
+                    dst_d + i3 * (dst->nb[3] / type_size),
+                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
+                    dst->ne[1], dst->ne[2], dim, stream);
+            }
+        } else {
+            const size_t size0 = ggml_nbytes(src0);
+            const size_t size1 = ggml_nbytes(src1);
+
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
+        }
+    } else {
+        concat_T_sycl_non_cont<block_q5_0>(
+            stream, (const char *) src0->data, (const char *) src1->data,
+            (char *) dst->data,
+            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
+            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
+            src1->ne[0] / QK5_0, src1->ne[1], src1->ne[2], src1->ne[3],
+            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
+            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
+            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
+    }
+}
+
+static void concat_impl_q5_1_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+    const ggml_tensor * src0   = dst->src[0];
+    const ggml_tensor * src1   = dst->src[1];
+    queue_ptr           stream = ctx.stream();
+
+    const int32_t dim = ((int32_t *) dst->op_params)[0];
+
+    GGML_ASSERT(src0->type == GGML_TYPE_Q5_1);
+    GGML_ASSERT(src1->type == GGML_TYPE_Q5_1);
+    GGML_ASSERT(dst->type == GGML_TYPE_Q5_1);
+    GGML_ASSERT(src0->ne[0] % QK5_1 == 0);
+    GGML_ASSERT(src1->ne[0] % QK5_1 == 0);
+    GGML_ASSERT(dst->ne[0] % QK5_1 == 0);
+
+    const int ne00_blk = src0->ne[0] / QK5_1;
+    const int ne0_blk  = dst->ne[0] / QK5_1;
+
+    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
+        const block_q5_1 * src0_d = (const block_q5_1 *) src0->data;
+        const block_q5_1 * src1_d = (const block_q5_1 *) src1->data;
+        block_q5_1 * dst_d        = (block_q5_1 *) dst->data;
+        const size_t type_size    = sizeof(block_q5_1);
+
+        if (dim != 3) {
+            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
+                concat_T_sycl<block_q5_1>(
+                    src0_d + i3 * (src0->nb[3] / type_size),
+                    src1_d + i3 * (src1->nb[3] / type_size),
+                    dst_d + i3 * (dst->nb[3] / type_size),
+                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
+                    dst->ne[1], dst->ne[2], dim, stream);
+            }
+        } else {
+            const size_t size0 = ggml_nbytes(src0);
+            const size_t size1 = ggml_nbytes(src1);
+
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
+        }
+    } else {
+        concat_T_sycl_non_cont<block_q5_1>(
+            stream, (const char *) src0->data, (const char *) src1->data,
+            (char *) dst->data,
+            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
+            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
+            src1->ne[0] / QK5_1, src1->ne[1], src1->ne[2], src1->ne[3],
+            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
+            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
+            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
+    }
+}
+
+static void concat_impl_q8_0_sycl(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
+    const ggml_tensor * src0   = dst->src[0];
+    const ggml_tensor * src1   = dst->src[1];
+    queue_ptr           stream = ctx.stream();
+
+    const int32_t dim = ((int32_t *) dst->op_params)[0];
+
+    GGML_ASSERT(src0->type == GGML_TYPE_Q8_0);
+    GGML_ASSERT(src1->type == GGML_TYPE_Q8_0);
+    GGML_ASSERT(dst->type == GGML_TYPE_Q8_0);
+    GGML_ASSERT(src0->ne[0] % QK8_0 == 0);
+    GGML_ASSERT(src1->ne[0] % QK8_0 == 0);
+    GGML_ASSERT(dst->ne[0] % QK8_0 == 0);
+
+    const int ne00_blk = src0->ne[0] / QK8_0;
+    const int ne0_blk  = dst->ne[0] / QK8_0;
+
+    if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) {
+        const block_q8_0 * src0_d = (const block_q8_0 *) src0->data;
+        const block_q8_0 * src1_d = (const block_q8_0 *) src1->data;
+        block_q8_0 * dst_d        = (block_q8_0 *) dst->data;
+        const size_t type_size    = sizeof(block_q8_0);
+
+        if (dim != 3) {
+            for (int i3 = 0; i3 < dst->ne[3]; i3++) {
+                concat_T_sycl<block_q8_0>(
+                    src0_d + i3 * (src0->nb[3] / type_size),
+                    src1_d + i3 * (src1->nb[3] / type_size),
+                    dst_d + i3 * (dst->nb[3] / type_size),
+                    ne00_blk, src0->ne[1], src0->ne[2], ne0_blk,
+                    dst->ne[1], dst->ne[2], dim, stream);
+            }
+        } else {
+            const size_t size0 = ggml_nbytes(src0);
+            const size_t size1 = ggml_nbytes(src1);
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dst_d, src0_d, size0)));
+            SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy((char *) dst_d + size0, src1_d, size1)));
+        }
+    } else {
+        concat_T_sycl_non_cont<block_q8_0>(
+            stream, (const char *) src0->data, (const char *) src1->data,
+            (char *) dst->data,
+            ne00_blk, src0->ne[1], src0->ne[2], src0->ne[3],
+            src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
+            src1->ne[0] / QK8_0, src1->ne[1], src1->ne[2], src1->ne[3],
+            src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
+            ne0_blk, dst->ne[1], dst->ne[2], dst->ne[3],
+            dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], dim);
+    }
+}
+
 void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
 
     switch (dst->type) {
@@ -222,6 +486,21 @@ void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
     case GGML_TYPE_I8:
         concat_impl_sycl<int8_t>(ctx, dst);
         break;
+    case GGML_TYPE_Q4_0:
+        concat_impl_q4_0_sycl(ctx, dst);
+        break;
+    case GGML_TYPE_Q4_1:
+        concat_impl_q4_1_sycl(ctx, dst);
+        break;
+    case GGML_TYPE_Q5_0:
+        concat_impl_q5_0_sycl(ctx, dst);
+        break;
+    case GGML_TYPE_Q5_1:
+        concat_impl_q5_1_sycl(ctx, dst);
+        break;
+    case GGML_TYPE_Q8_0:
+        concat_impl_q8_0_sycl(ctx, dst);
+        break;
     default:
         fprintf(stderr, "%s: unsupported types: dst: %s\n", __func__, ggml_type_name(dst->type));
         GGML_ASSERT(false);

ggml/src/ggml-sycl/dmmv.cpp

@@ -8,6 +8,9 @@
         #include <sycl/ext/oneapi/bfloat16.hpp>
         #define GGML_SYCL_DMMV_HAS_BF16
     #endif
+    #include <sycl/ext/intel/esimd.hpp>
+    #include "esimd.hpp"
+    #define GGML_SYCL_DMMV_HAS_ESIMD
 #endif
 
 static void convert_f16(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){
@@ -1864,6 +1867,113 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
         });
 }
 
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+using ggml_sycl_esimd::GGML_SYCL_DMMV_ESIMD_WG_SIZE;
+
+// generic reordered dequantize-matvec: each work-group owns a pair of
+// consecutive output rows and updates one 32-wide accumulator per row
+template <ggml_type T>
+ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd(
+        const void * vx, const float * y, float * dst,
+        const int ncols, const int nrows,
+        sycl::local_accessor<float, 1> lmem,
+        const sycl::nd_item<1> & it) {
+    using namespace sycl::ext::intel::esimd;
+    using traits = ggml_sycl_esimd::esimd_reorder_q_traits<T>;
+
+    const int    num_blocks_per_row = ncols / QK_K;
+    const size_t nb = (size_t) nrows * num_blocks_per_row;
+    const auto   ps = traits::make_ptrs(vx, nb);
+
+    const int  tid      = it.get_local_id(0);
+    const int  row_pair = it.get_group(0);
+    const int  row0     = row_pair * 2; // two consecutive output rows
+    const bool has_row1 = row0 + 1 < nrows;
+
+    // one 32-wide accumulator per output row (small footprint, no spill)
+    simd<float, 32> acc0 = 0.0f;
+    simd<float, 32> acc1 = 0.0f;
+
+    for (int ib = tid; ib < num_blocks_per_row; ib += GGML_SYCL_DMMV_ESIMD_WG_SIZE) {
+        simd<float, 256> y_vec = block_load<float, 256>(y + (size_t) ib * QK_K);
+
+        const size_t bi0 = (size_t) (row0 + 0) * num_blocks_per_row + ib;
+        const size_t bi1 = (size_t) (row0 + 1) * num_blocks_per_row + ib;
+
+        traits::mac_pair(ps, bi0, ps, bi1, has_row1, y_vec, acc0, acc1);
+    }
+
+    lmem[tid * 2 + 0] = reduce<float>(acc0, std::plus<>{});
+    lmem[tid * 2 + 1] = reduce<float>(acc1, std::plus<>{});
+    it.barrier(sycl::access::fence_space::local_space);
+
+    if (tid == 0) {
+        float sum0 = 0.0f;
+        float sum1 = 0.0f;
+        for (int p = 0; p < GGML_SYCL_DMMV_ESIMD_WG_SIZE; ++p) {
+            sum0 += lmem[p * 2 + 0];
+            sum1 += lmem[p * 2 + 1];
+        }
+        dst[row0 + 0] = sum0;
+        if (has_row1) {
+            dst[row0 + 1] = sum1;
+        }
+    }
+}
+
+static void dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(const void *vx, const float *y,
+                                                           float *dst, const int ncols,
+                                                           const int nrows,
+                                                           dpct::queue_ptr stream) {
+    GGML_ASSERT(ncols % QK_K == 0);
+    const int workgroups = (nrows + 1) / 2;
+    stream->submit([&](sycl::handler &h) {
+        sycl::local_accessor<float, 1> lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h);
+        h.parallel_for(
+            sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)),
+            [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
+                dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q3_K>(
+                    vx, y, dst, ncols, nrows, lmem, it);
+            });
+    });
+}
+
+static void dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(const void *vx, const float *y,
+                                                           float *dst, const int ncols,
+                                                           const int nrows,
+                                                           dpct::queue_ptr stream) {
+    GGML_ASSERT(ncols % QK_K == 0);
+    const int workgroups = (nrows + 1) / 2;
+    stream->submit([&](sycl::handler &h) {
+        sycl::local_accessor<float, 1> lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h);
+        h.parallel_for(
+            sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)),
+            [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
+                dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q4_K>(
+                    vx, y, dst, ncols, nrows, lmem, it);
+            });
+    });
+}
+
+static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const float *y,
+                                                           float *dst, const int ncols,
+                                                           const int nrows,
+                                                           dpct::queue_ptr stream) {
+    GGML_ASSERT(ncols % QK_K == 0);
+    const int workgroups = (nrows + 1) / 2;
+    stream->submit([&](sycl::handler &h) {
+        sycl::local_accessor<float, 1> lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h);
+        h.parallel_for(
+            sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)),
+            [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
+                dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q6_K>(
+                    vx, y, dst, ncols, nrows, lmem, it);
+            });
+    });
+}
+
+#endif // GGML_SYCL_DMMV_HAS_ESIMD
+
 static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y,
                                                      float *dst, const int ncols,
                                                      const int nrows,
@@ -1992,7 +2102,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
         case GGML_TYPE_Q3_K:
             if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
                 ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
-                dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+                if (g_ggml_sycl_enable_esimd) {
+                    dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
+                else
+#endif
+                {
+                    dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
             } else {
                 dequantize_mul_mat_vec_q3_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }
@@ -2000,7 +2118,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
         case GGML_TYPE_Q4_K:
             if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
                 ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
-                dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+                if (g_ggml_sycl_enable_esimd) {
+                    dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
+                else
+#endif
+                {
+                    dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
             } else {
                 dequantize_mul_mat_vec_q4_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }
@@ -2016,7 +2142,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
         case GGML_TYPE_Q6_K:
             if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
                 ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
-                dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+                if (g_ggml_sycl_enable_esimd) {
+                    dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
+                else
+#endif
+                {
+                    dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
             } else {
                 dequantize_mul_mat_vec_q6_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }

ggml/src/ggml-sycl/element_wise.cpp

@@ -448,6 +448,47 @@ static void unary_gated_op_generic_kernel(
     }
 }
 
+// Fused UNARY + MUL. Unlike the gated ops above, `x` and `g` are separate tensors of the
+// same shape; `o0`/`o1` are their row strides in elements, so a half-view needs no repack.
+// `dst` is contiguous and indexed flat. Math is done in f32, as the CPU and CUDA references do.
+template<typename T, typename F>
+static void unary_mul_flat_kernel(const T * x, const T * g, T * dst, const int64_t k, const sycl::nd_item<1> &item_ct1, F op) {
+    SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
+        dst[i] = (T) (op((float) x[i]) * (float) g[i]);
+    }
+}
+
+template<typename T, typename F>
+static void unary_mul_strided_kernel(const T * x, const T * g, T * dst, const int64_t k, const sycl::uint3 n_fd, const int64_t o0, const int64_t o1, const sycl::nd_item<1> &item_ct1, F op) {
+    SYCL_GLOBAL_ID_LOOP(k, item_ct1) {
+        const sycl::uint2 rc = fast_div_modulo((uint32_t) i, n_fd);
+        const int64_t j0 = rc.x() * o0 + rc.y();
+        const int64_t j1 = o0 == o1 ? j0 : rc.x() * o1 + rc.y();
+        dst[i] = (T) (op((float) x[j0]) * (float) g[j1]);
+    }
+}
+
+template<typename T, typename F>
+static void unary_mul_sycl(const T * x, const T * g, T * dst, const int64_t k, const int64_t n, const int64_t o0, const int64_t o1, queue_ptr main_stream, F op) {
+    const size_t            num_blocks = ceil_div((size_t) k, (size_t) SYCL_GLU_BLOCK_SIZE);
+    const sycl::nd_range<1> range(num_blocks * sycl::range<1>(SYCL_GLU_BLOCK_SIZE), sycl::range<1>(SYCL_GLU_BLOCK_SIZE));
+
+    // o0 == o1 == n makes (i/n)*o0 + (i%n) == i, so the strided kernel degenerates to the flat one
+    if (o0 == n && o1 == n) {
+        main_stream->parallel_for(range, [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+            unary_mul_flat_kernel(x, g, dst, k, item_ct1, op);
+        });
+        return;
+    }
+
+    // 32-bit fastdiv, exact only below 2^31; ggml_sycl_can_fuse() already declined past that
+    GGML_ASSERT(k < ((int64_t) 1 << 31));
+    const sycl::uint3 n_fd = init_fastdiv_values((uint32_t) n);
+    main_stream->parallel_for(range, [=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
+        unary_mul_strided_kernel(x, g, dst, k, n_fd, o0, o1, item_ct1, op);
+    });
+}
+
 namespace ggml_sycl_detail {
 static void acc_f32_sycl(const char *x, const char *y, float *dst,
                          const int64_t n_elements,
@@ -991,6 +1032,52 @@ static inline void ggml_sycl_op_swiglu(ggml_backend_sycl_context & ctx, ggml_ten
     });
 }
 
+// dst = op(unary_node->src[0]) * other, written straight to the MUL output, saving the
+// standalone unary launch. Preconditions come from ggml_sycl_can_fuse(); re-asserted here.
+void ggml_sycl_op_unary_mul_fused(ggml_backend_sycl_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node) {
+    scope_op_debug_print scope_dbg_print(__func__, mul_node, /*num_src=*/2);
+
+    const ggml_tensor * x = unary_node->src[0];
+    const ggml_tensor * g = (mul_node->src[0] == unary_node) ? mul_node->src[1] : mul_node->src[0];
+
+    // g is picked by elimination; ggml_can_fuse()'s single-use rule rules out MUL(unary, unary)
+    GGML_ASSERT(g != unary_node);
+    GGML_ASSERT(x->type == g->type && x->type == mul_node->type);
+    GGML_ASSERT(ggml_are_same_shape(x, g) && ggml_are_same_shape(x, mul_node));
+    GGML_ASSERT(ggml_is_contiguous_1(x) && ggml_is_contiguous_1(g));
+    // dst is indexed flat
+    GGML_ASSERT(ggml_is_contiguous(mul_node));
+
+    queue_ptr main_stream = ctx.stream();
+    SYCL_CHECK(ggml_sycl_set_device(ctx.device));
+
+    const int64_t k = ggml_nelements(mul_node);
+    const int64_t n = mul_node->ne[0];
+
+    const auto dispatch_type = [&](auto op) {
+        switch (mul_node->type) {
+            case GGML_TYPE_F32:
+                unary_mul_sycl((const float *) x->data, (const float *) g->data, (float *) mul_node->data,
+                               k, n, x->nb[1] / sizeof(float), g->nb[1] / sizeof(float), main_stream, op);
+                break;
+            case GGML_TYPE_F16:
+                unary_mul_sycl((const sycl::half *) x->data, (const sycl::half *) g->data, (sycl::half *) mul_node->data,
+                               k, n, x->nb[1] / sizeof(sycl::half), g->nb[1] / sizeof(sycl::half), main_stream, op);
+                break;
+            default:
+                GGML_ABORT("fused unary+mul: unsupported type %s", ggml_type_name(mul_node->type));
+        }
+    };
+
+    switch (ggml_get_unary_op(unary_node)) {
+        case GGML_UNARY_OP_SILU:     dispatch_type([](float v) { return op_silu(v); });     break;
+        case GGML_UNARY_OP_SIGMOID:  dispatch_type([](float v) { return op_sigmoid(v); });  break;
+        case GGML_UNARY_OP_SOFTPLUS: dispatch_type([](float v) { return op_softplus(v); }); break;
+        default:
+            GGML_ABORT("fused unary+mul: unsupported unary op %s", ggml_unary_op_name(ggml_get_unary_op(unary_node)));
+    }
+}
+
 __dpct_inline__ float ggml_sycl_op_swiglu_oai_single(float x, float g, float alpha = 1.702f, float limit = 7.0f) {
     x = sycl::fmin(x, limit);
     g = sycl::fmax(sycl::fmin(g, limit), -limit);

ggml/src/ggml-sycl/element_wise.hpp

@@ -95,4 +95,7 @@ void ggml_sycl_trunc(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
 
 void ggml_sycl_arange(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
 
+// fused UNARY(silu|sigmoid|softplus) + MUL; see ggml_sycl_can_fuse() for the accepted shapes
+void ggml_sycl_op_unary_mul_fused(ggml_backend_sycl_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node);
+
 #endif // GGML_SYCL_ELEMENTWISE_HPP

ggml/src/ggml-sycl/esimd.hpp

@@ -0,0 +1,392 @@
+//
+// MIT license
+// Copyright (C) 2026 Intel Corporation
+// SPDX-License-Identifier: MIT
+//
+
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+
+#ifndef GGML_SYCL_ESIMD_HPP
+#define GGML_SYCL_ESIMD_HPP
+
+#include <sycl/ext/intel/esimd.hpp>
+
+#include "common.hpp"
+
+namespace ggml_sycl_esimd {
+
+constexpr int GGML_SYCL_DMMV_ESIMD_WG_SIZE = 4;
+
+//
+// Shared ESIMD building blocks for the reordered K-quant dequantize-matvec
+// kernels.
+//
+// The reordered K-quant ESIMD matvec kernels share one skeleton: per super-block,
+// load a 256-float activation slice, load one weight block, dequantize it into 8
+// chunks of 32 and MAC each chunk against the matching activation slice, then
+// reduce and run a lane-0 epilogue.
+//
+// Each K-quant kernel emits exactly 8 chunks of 32 mapping to activation slices
+// 0..7, so the per-block work is captured by esimd_reorder_q_traits<T>::mac_pair,
+// which dequantizes two weight blocks and MACs both against a shared activation
+// vector with the two FMA chains interleaved (co-scheduled to hide FMA latency).
+// The "pair" is the (row0,row1) row pair owned by one work-group, so the
+// layout+dequant is written once per quant type here.
+//
+
+template <ggml_type T> struct esimd_reorder_q_traits;
+
+// build a 32-lane vector whose low 16 lanes are `lo` and high 16 are `hi`
+// (a super-chunk splits into two 16-wide halves with distinct scale/min codes).
+static ESIMD_INLINE sycl::ext::intel::esimd::simd<float, 32> splat_lo_hi(float lo, float hi) {
+    using namespace sycl::ext::intel::esimd;
+    simd<float, 32> v;
+    v.select<16, 1>(0)  = lo;
+    v.select<16, 1>(16) = hi;
+    return v;
+}
+
+// unpack one block of Q4_K/Q5_K scale/min codes (get_scale_min_k4 layout) into 8
+// float scales (dall * sc) and 8 float mins (-dmin * m); the min carries the
+// negation so the dequant epilogue adds.
+static ESIMD_INLINE void unpack_scale_min_k4(
+        sycl::ext::intel::esimd::simd<uint8_t, 12> scales, float dall, float dmin,
+        sycl::ext::intel::esimd::simd<float, 8> & scale_f,
+        sycl::ext::intel::esimd::simd<float, 8> & min_f) {
+    using namespace sycl::ext::intel::esimd;
+    simd<uint8_t, 8> sc = 0;
+    simd<uint8_t, 8> m  = 0;
+    simd<uint8_t, 4> scale_lo = scales.select<4, 1>(0);
+    simd<uint8_t, 4> min_lo   = scales.select<4, 1>(4);
+    simd<uint8_t, 4> hi_bits  = scales.select<4, 1>(8);
+    sc.select<4, 1>(0) = scale_lo & simd<uint8_t, 4>(0x3F);
+    sc.select<4, 1>(4) = (hi_bits & simd<uint8_t, 4>(0x0F)) |
+                         ((scale_lo >> simd<uint8_t, 4>(6)) << simd<uint8_t, 4>(4));
+    m.select<4, 1>(0)  = min_lo & simd<uint8_t, 4>(0x3F);
+    m.select<4, 1>(4)  = (hi_bits >> simd<uint8_t, 4>(4)) |
+                         ((min_lo >> simd<uint8_t, 4>(6)) << simd<uint8_t, 4>(4));
+    scale_f = convert<float>(sc) * dall;
+    min_f   = convert<float>(m) * (-dmin);
+}
+
+// ---------------------------------------------------------------------------
+// Q3_K, SOA reorder layout produced by reorder_qw_q3_k:
+//   [qs: nb*(QK_K/4)] [hmask: nb*(QK_K/8)] [scales: nb*12] [d: nb*sizeof(half)]
+// with nb = nrows*num_blocks_per_row. Single super-block scale d, no dmin.
+//
+// 3 bits per weight: 2 low bits in qs, 1 high bit in hmask. The 8 output chunks
+// of 32 (matching dequantize_row_q3_K) map to super-chunk s (0..7): byte base
+// 32*(s/4) into the 64-byte qs array, bit shift 2*(s%4); the low 16 lanes use
+// scale code 2s, the high 16 use 2s+1. hmask is a 32-byte array (like Q5_K's
+// qh) where chunk s uses bit s of the same 32 bytes, but INVERTED: the value is
+// (q & 3) - (hmask_bit_set ? 0 : 4), i.e. (q & 3) + 4*bit - 4.
+//
+// The 16 6-bit scale codes are packed into 12 bytes (get_scale_min layout for
+// Q3_K): low nibbles from bytes 0..7, high 2 bits from bytes 8..11 shifted by
+// 0/2/4/6; the dequant scale is d * (code - 32).
+// ---------------------------------------------------------------------------
+template <> struct esimd_reorder_q_traits<GGML_TYPE_Q3_K> {
+    struct ptrs {
+        const uint8_t *    qs;
+        const uint8_t *    hmask;
+        const uint8_t *    scales;
+        const sycl::half * d;
+    };
+
+    static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
+        const uint8_t * qs     = (const uint8_t *) vx;
+        const uint8_t * hmask  = qs + nb * (QK_K / 4);
+        const uint8_t * scales = hmask + nb * (QK_K / 8);
+        const sycl::half * d   = (const sycl::half *) (scales + nb * 12);
+        return { qs, hmask, scales, d };
+    }
+
+    // unpack the 12 packed bytes into 16 6-bit scale codes (dequantize_row_q3_K
+    // aux layout), returned as float scale = d * (code - 32).
+    // done with wide (8/16-lane) ops rather than four 4-lane groups.
+    static ESIMD_INLINE sycl::ext::intel::esimd::simd<float, 16> unpack_scales(
+            sycl::ext::intel::esimd::simd<uint8_t, 12> in, float d) {
+        using namespace sycl::ext::intel::esimd;
+
+        // low 6-bit part: codes 0..7 = low nibble of bytes 0..7,
+        //                 codes 8..15 = high nibble of bytes 0..7
+        simd<uint8_t, 8>  lo8 = in.select<8, 1>(0);
+        simd<uint8_t, 16> code;
+        code.select<8, 1>(0) = lo8 & simd<uint8_t, 8>(0x0F);
+        code.select<8, 1>(8) = lo8 >> simd<uint8_t, 8>(4);
+
+        // high 2-bit part: bytes 8..11 replicated 4x, group g (0..3) shifted 2*g
+        simd<uint8_t, 16> hib;
+        hib.select<4, 1>(0)  = in.select<4, 1>(8);
+        hib.select<4, 1>(4)  = in.select<4, 1>(8);
+        hib.select<4, 1>(8)  = in.select<4, 1>(8);
+        hib.select<4, 1>(12) = in.select<4, 1>(8);
+        simd<uint8_t, 16> hshift;
+        hshift.select<4, 1>(0)  = 0;
+        hshift.select<4, 1>(4)  = 2;
+        hshift.select<4, 1>(8)  = 4;
+        hshift.select<4, 1>(12) = 6;
+        hib = (hib >> hshift) & simd<uint8_t, 16>(0x03);
+
+        code = code | (hib << simd<uint8_t, 16>(4));
+        return (convert<float>(code) - 32.0f) * d;
+    }
+
+    static ESIMD_INLINE void mac_pair(
+            const ptrs & pa, size_t bia,
+            const ptrs & pb, size_t bib, bool has_b,
+            sycl::ext::intel::esimd::simd<float, 256> & y_vec,
+            sycl::ext::intel::esimd::simd<float, 32> & acc_a,
+            sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
+        using namespace sycl::ext::intel::esimd;
+
+        simd<uint8_t, 64> qs_a     = block_load<uint8_t, 64>(pa.qs + bia * (QK_K / 4));
+        simd<uint8_t, 64> qs_b     = 0;
+        simd<uint8_t, 32> hmask_a  = block_load<uint8_t, 32>(pa.hmask + bia * (QK_K / 8));
+        simd<uint8_t, 32> hmask_b  = 0;
+        simd<uint8_t, 12> scales_a = block_load<uint8_t, 12>(pa.scales + bia * 12);
+        simd<uint8_t, 12> scales_b = 0;
+
+        const float d_a = (float) pa.d[bia];
+        float d_b = 0.0f;
+        if (has_b) {
+            qs_b     = block_load<uint8_t, 64>(pb.qs + bib * (QK_K / 4));
+            hmask_b  = block_load<uint8_t, 32>(pb.hmask + bib * (QK_K / 8));
+            scales_b = block_load<uint8_t, 12>(pb.scales + bib * 12);
+            d_b = (float) pb.d[bib];
+        }
+
+        simd<float, 16> scale_f_a = unpack_scales(scales_a, d_a);
+        simd<float, 16> scale_f_b = unpack_scales(scales_b, d_b);
+
+#pragma unroll
+        for (int s = 0; s < 8; ++s) {
+            const int     byte_base = 32 * (s / 4);
+            const uint8_t shift     = (uint8_t) (2 * (s % 4));
+            simd<float, 32> y_s = y_vec.select<32, 1>(s * 32);
+
+            // 2 low bits from qs, high bit from hmask (bit s of the same 32 bytes);
+            // value = (q & 3) + 4*bit - 4  (inverted hmask: subtract 4 when bit clear).
+            // merge in the integer domain: q3 = (q & 3) | (bit << 2) in {0..7},
+            // then a single convert + subtract yields q3 - 4 (one convert, not two)
+            simd<uint16_t, 32> q3_a = convert<uint16_t>(
+                    (qs_a.select<32, 1>(byte_base) >> shift) & simd<uint8_t, 32>(3));
+            q3_a |= convert<uint16_t>(
+                    ((hmask_a >> simd<uint8_t, 32>((uint8_t) s)) & simd<uint8_t, 32>(1)) << simd<uint8_t, 32>(2));
+            simd<uint16_t, 32> q3_b = convert<uint16_t>(
+                    (qs_b.select<32, 1>(byte_base) >> shift) & simd<uint8_t, 32>(3));
+            q3_b |= convert<uint16_t>(
+                    ((hmask_b >> simd<uint8_t, 32>((uint8_t) s)) & simd<uint8_t, 32>(1)) << simd<uint8_t, 32>(2));
+
+            simd<float, 32> qf_a = convert<float>(q3_a) - 4.0f;
+            simd<float, 32> qf_b = convert<float>(q3_b) - 4.0f;
+
+            const float scale_a_lo = scale_f_a[2 * s + 0];
+            const float scale_a_hi = scale_f_a[2 * s + 1];
+            const float scale_b_lo = scale_f_b[2 * s + 0];
+            const float scale_b_hi = scale_f_b[2 * s + 1];
+
+            simd<float, 32> scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
+            simd<float, 32> scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
+
+            simd<float, 32> deq_a = qf_a * scale_vec_a;
+            simd<float, 32> deq_b = qf_b * scale_vec_b;
+
+            acc_a += y_s * deq_a;
+            acc_b += y_s * deq_b;
+        }
+    }
+};
+
+// ---------------------------------------------------------------------------
+// Q4_K, SOA reorder layout produced by reorder_qw_q4_k:
+//   [qs: nb*(QK_K/2)] [scales: nb*K_SCALE_SIZE] [dm: nb*sizeof(half2)]
+// with nb = nrows*num_blocks_per_row.
+// ---------------------------------------------------------------------------
+template <> struct esimd_reorder_q_traits<GGML_TYPE_Q4_K> {
+    struct ptrs {
+        const uint8_t *    qs;
+        const uint8_t *    scales;
+        const sycl::half * dm;
+    };
+
+    static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
+        const uint8_t * qs     = (const uint8_t *) vx;
+        const uint8_t * scales = qs + nb * (QK_K / 2);
+        const sycl::half * dm  = (const sycl::half *) (scales + nb * K_SCALE_SIZE);
+        return { qs, scales, dm };
+    }
+
+    static ESIMD_INLINE void mac_pair(
+            const ptrs & pa, size_t bia,
+            const ptrs & pb, size_t bib, bool has_b,
+            sycl::ext::intel::esimd::simd<float, 256> & y_vec,
+            sycl::ext::intel::esimd::simd<float, 32> & acc_a,
+            sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
+        using namespace sycl::ext::intel::esimd;
+
+        simd<uint8_t, 128> qs_a     = block_load<uint8_t, 128>(pa.qs + bia * (QK_K / 2));
+        simd<uint8_t, 128> qs_b     = 0;
+        simd<uint8_t, 12>  scales_a = block_load<uint8_t, 12>(pa.scales + bia * K_SCALE_SIZE);
+        simd<uint8_t, 12>  scales_b = 0;
+
+        const float dall_a = (float) pa.dm[bia * 2 + 0];
+        const float dmin_a = (float) pa.dm[bia * 2 + 1];
+        float dall_b = 0.0f;
+        float dmin_b = 0.0f;
+        if (has_b) {
+            qs_b     = block_load<uint8_t, 128>(pb.qs + bib * (QK_K / 2));
+            scales_b = block_load<uint8_t, 12>(pb.scales + bib * K_SCALE_SIZE);
+            dall_b = (float) pb.dm[bib * 2 + 0];
+            dmin_b = (float) pb.dm[bib * 2 + 1];
+        }
+
+        simd<float, 8> scale_f_a, min_f_a, scale_f_b, min_f_b;
+        unpack_scale_min_k4(scales_a, dall_a, dmin_a, scale_f_a, min_f_a);
+        unpack_scale_min_k4(scales_b, dall_b, dmin_b, scale_f_b, min_f_b);
+
+        simd<uint8_t, 128> qs_lo_a = qs_a & simd<uint8_t, 128>(0x0F);
+        simd<uint8_t, 128> qs_hi_a = qs_a >> simd<uint8_t, 128>(4);
+        simd<uint8_t, 128> qs_lo_b = qs_b & simd<uint8_t, 128>(0x0F);
+        simd<uint8_t, 128> qs_hi_b = qs_b >> simd<uint8_t, 128>(4);
+
+#pragma unroll
+        for (int sb = 0; sb < 8; sb += 2) {
+            const int q_offset = sb * 16;
+            simd<float, 32> y_lo = y_vec.select<32, 1>(sb * 32);
+            simd<float, 32> y_hi = y_vec.select<32, 1>((sb + 1) * 32);
+
+            const float scale_a_lo = scale_f_a[sb];
+            const float scale_a_hi = scale_f_a[sb + 1];
+            const float min_a_lo   = min_f_a[sb];
+            const float min_a_hi   = min_f_a[sb + 1];
+            const float scale_b_lo = scale_f_b[sb];
+            const float scale_b_hi = scale_f_b[sb + 1];
+            const float min_b_lo   = min_f_b[sb];
+            const float min_b_hi   = min_f_b[sb + 1];
+
+            simd<uint8_t, 32> qa_lo = qs_lo_a.select<32, 1>(q_offset);
+            simd<uint8_t, 32> qa_hi = qs_hi_a.select<32, 1>(q_offset);
+            simd<uint8_t, 32> qb_lo = qs_lo_b.select<32, 1>(q_offset);
+            simd<uint8_t, 32> qb_hi = qs_hi_b.select<32, 1>(q_offset);
+
+            simd<float, 32> deq_a_lo = convert<float>(qa_lo) * scale_a_lo + min_a_lo;
+            simd<float, 32> deq_a_hi = convert<float>(qa_hi) * scale_a_hi + min_a_hi;
+            simd<float, 32> deq_b_lo = convert<float>(qb_lo) * scale_b_lo + min_b_lo;
+            simd<float, 32> deq_b_hi = convert<float>(qb_hi) * scale_b_hi + min_b_hi;
+
+            acc_a += y_lo * deq_a_lo;
+            acc_b += y_lo * deq_b_lo;
+            acc_a += y_hi * deq_a_hi;
+            acc_b += y_hi * deq_b_hi;
+        }
+    }
+};
+
+// ---------------------------------------------------------------------------
+// Q6_K, SOA reorder layout:
+//   [ql: nb*(QK_K/2)] [qh: nb*(QK_K/4)] [scales(int8): nb*(QK_K/16)] [d: nb*half]
+// ---------------------------------------------------------------------------
+template <> struct esimd_reorder_q_traits<GGML_TYPE_Q6_K> {
+    struct ptrs {
+        const uint8_t *    ql;
+        const uint8_t *    qh;
+        const int8_t *     scales;
+        const sycl::half * d;
+    };
+
+    static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
+        const uint8_t *    ql     = (const uint8_t *) vx;
+        const uint8_t *    qh     = ql + nb * (QK_K / 2);
+        const int8_t *     scales = (const int8_t *) (qh + nb * (QK_K / 4));
+        const sycl::half * d      = (const sycl::half *) (scales + nb * (QK_K / 16));
+        return { ql, qh, scales, d };
+    }
+
+    static ESIMD_INLINE void mac_pair(
+            const ptrs & pa, size_t bia,
+            const ptrs & pb, size_t bib, bool has_b,
+            sycl::ext::intel::esimd::simd<float, 256> & y_vec,
+            sycl::ext::intel::esimd::simd<float, 32> & acc_a,
+            sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
+        using namespace sycl::ext::intel::esimd;
+
+        simd<uint8_t, 128> ql_a     = block_load<uint8_t, 128>(pa.ql + bia * (QK_K / 2));
+        simd<uint8_t, 128> ql_b     = 0;
+        simd<uint8_t, 64>  qh_a     = block_load<uint8_t, 64>(pa.qh + bia * (QK_K / 4));
+        simd<uint8_t, 64>  qh_b     = 0;
+        simd<int8_t, 16>   scales_a = block_load<int8_t, 16>(pa.scales + bia * (QK_K / 16));
+        simd<int8_t, 16>   scales_b = 0;
+
+        const float d_a = (float) pa.d[bia];
+        float d_b = 0.0f;
+        if (has_b) {
+            ql_b     = block_load<uint8_t, 128>(pb.ql + bib * (QK_K / 2));
+            qh_b     = block_load<uint8_t, 64>(pb.qh + bib * (QK_K / 4));
+            scales_b = block_load<int8_t, 16>(pb.scales + bib * (QK_K / 16));
+            d_b = (float) pb.d[bib];
+        }
+
+        simd<float, 16> sc_a = convert<float>(scales_a);
+        simd<float, 16> sc_b = convert<float>(scales_b);
+
+#pragma unroll
+        for (int im = 0; im < 2; ++im) {
+            simd<uint8_t, 32> ql_lo_a   = ql_a.select<32, 1>(64 * im);
+            simd<uint8_t, 32> ql_hi_a   = ql_a.select<32, 1>(64 * im + 32);
+            simd<uint8_t, 32> qh_bits_a = qh_a.select<32, 1>(32 * im);
+            simd<uint8_t, 32> ql_lo_b   = ql_b.select<32, 1>(64 * im);
+            simd<uint8_t, 32> ql_hi_b   = ql_b.select<32, 1>(64 * im + 32);
+            simd<uint8_t, 32> qh_bits_b = qh_b.select<32, 1>(32 * im);
+
+            // reconstruct each 32-wide 6-bit group (matches dequantize_row_q6_K)
+#pragma unroll
+            for (int g = 0; g < 4; ++g) {
+                simd<float, 32> y_g = y_vec.select<32, 1>(32 * (4 * im + g));
+
+                const float scale_a_lo = sc_a[8 * im + 2 * g + 0] * d_a;
+                const float scale_a_hi = sc_a[8 * im + 2 * g + 1] * d_a;
+                const float scale_b_lo = sc_b[8 * im + 2 * g + 0] * d_b;
+                const float scale_b_hi = sc_b[8 * im + 2 * g + 1] * d_b;
+
+                simd<float, 32> scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
+                simd<float, 32> scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
+
+                simd<uint8_t, 32> qa;
+                simd<uint8_t, 32> qb;
+                switch (g) {
+                    case 0:
+                        qa = (ql_lo_a & simd<uint8_t, 32>(0x0F)) | ((qh_bits_a & simd<uint8_t, 32>(0x03)) << simd<uint8_t, 32>(4));
+                        qb = (ql_lo_b & simd<uint8_t, 32>(0x0F)) | ((qh_bits_b & simd<uint8_t, 32>(0x03)) << simd<uint8_t, 32>(4));
+                        break;
+                    case 1:
+                        qa = (ql_hi_a & simd<uint8_t, 32>(0x0F)) | ((qh_bits_a & simd<uint8_t, 32>(0x0C)) << simd<uint8_t, 32>(2));
+                        qb = (ql_hi_b & simd<uint8_t, 32>(0x0F)) | ((qh_bits_b & simd<uint8_t, 32>(0x0C)) << simd<uint8_t, 32>(2));
+                        break;
+                    case 2:
+                        qa = (ql_lo_a >> simd<uint8_t, 32>(4)) | (qh_bits_a & simd<uint8_t, 32>(0x30));
+                        qb = (ql_lo_b >> simd<uint8_t, 32>(4)) | (qh_bits_b & simd<uint8_t, 32>(0x30));
+                        break;
+                    default:
+                        qa = (ql_hi_a >> simd<uint8_t, 32>(4)) | ((qh_bits_a & simd<uint8_t, 32>(0xC0)) >> simd<uint8_t, 32>(2));
+                        qb = (ql_hi_b >> simd<uint8_t, 32>(4)) | ((qh_bits_b & simd<uint8_t, 32>(0xC0)) >> simd<uint8_t, 32>(2));
+                        break;
+                }
+
+                simd<float, 32> deq_a = (convert<float>(qa) - 32.0f) * scale_vec_a;
+                simd<float, 32> deq_b = (convert<float>(qb) - 32.0f) * scale_vec_b;
+
+                acc_a += y_g * deq_a;
+                acc_b += y_g * deq_b;
+            }
+        }
+    }
+};
+
+} // namespace ggml_sycl_esimd
+
+#endif // GGML_SYCL_ESIMD_HPP

ggml/src/ggml-sycl/fusion.cpp

@@ -1,6 +1,14 @@
 #include "fusion.hpp"
 
-bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops) {
+#include <algorithm>
+
+bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops,
+                        std::initializer_list<enum ggml_unary_op> unary_ops) {
+#ifndef NDEBUG
+    const size_t num_unary = std::count(ops.begin(), ops.end(), GGML_OP_UNARY);
+    GGML_ASSERT(unary_ops.size() == num_unary);
+#endif
+
     if (!g_ggml_sycl_enable_fusion) {
         return false;
     }
@@ -40,5 +48,45 @@ bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializ
         return true;
     }
 
+    if (ops.size() == 2 && ops.begin()[0] == GGML_OP_UNARY && ops.begin()[1] == GGML_OP_MUL &&
+        unary_ops.size() == 1) {
+        const ggml_tensor * unary = cgraph->nodes[node_idx];
+        const ggml_tensor * mul   = cgraph->nodes[node_idx + 1];
+
+        const ggml_unary_op unary_op = ggml_get_unary_op(unary);
+        if (unary_op != unary_ops.begin()[0]) {
+            return false;
+        }
+
+        // the ops ggml_sycl_op_unary_mul_fused() has a kernel for
+        if (unary_op != GGML_UNARY_OP_SILU && unary_op != GGML_UNARY_OP_SIGMOID &&
+            unary_op != GGML_UNARY_OP_SOFTPLUS) {
+            return false;
+        }
+
+        if (unary->type != GGML_TYPE_F32 && unary->type != GGML_TYPE_F16) {
+            return false;
+        }
+
+        const ggml_tensor * other = (mul->src[0] == unary) ? mul->src[1] : mul->src[0];
+        if (other->type != unary->type) {
+            return false;
+        }
+
+        // one row stride per source comes from nb[1], so rows must be contiguous and equally
+        // shaped; the destination is written flat, so it must be fully contiguous
+        if (!ggml_is_contiguous_1(unary->src[0]) || !ggml_is_contiguous_1(other) ||
+            !ggml_are_same_shape(other, unary) || !ggml_is_contiguous(mul)) {
+            return false;
+        }
+
+        // the 32-bit fastdiv is inexact past 2^31; decline, the unfused path handles it
+        if (ggml_nelements(mul) >= ((int64_t) 1 << 31)) {
+            return false;
+        }
+
+        return true;
+    }
+
     return false;
 }

ggml/src/ggml-sycl/fusion.hpp

@@ -6,10 +6,12 @@
 #include "common.hpp"
 
 // Backend-side fusability test. `ops` names a candidate op sequence starting at cgraph node
-// `node_idx`; the result is true only if ggml considers that subgraph fusable *and* the SYCL
+// `node_idx`, and `unary_ops` the GGML_UNARY_OP each GGML_OP_UNARY in `ops` must carry, in
+// order; the result is true only if ggml considers that subgraph fusable *and* the SYCL
 // kernel which would service it accepts the tensors involved (types, shapes, contiguity).
 //
 // Lives in its own translation unit because it grows a branch per supported op sequence.
-bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops);
+bool ggml_sycl_can_fuse(const ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops,
+                        std::initializer_list<enum ggml_unary_op> unary_ops);
 
 #endif  // GGML_SYCL_FUSION_HPP

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -43,6 +43,9 @@
 #    include <sycl/ext/oneapi/virtual_mem/virtual_mem.hpp>
 #    define GGML_SYCL_SUPPORT_VMM
 #endif
+#if defined(__INTEL_LLVM_COMPILER)
+    #define GGML_SYCL_DMMV_HAS_ESIMD
+#endif
 #include <sycl/half_type.hpp>
 
 #include "ggml.h"
@@ -90,6 +93,7 @@ int g_ggml_sycl_fa_onednn = 1;
 int g_ggml_sycl_fa_onednn_max_kv = 0;
 int g_ggml_sycl_enable_vmm = 1;
 int g_ggml_sycl_enable_fusion = 1;
+int g_ggml_sycl_enable_esimd = 1;
 int g_ggml_sycl_prioritize_dmmv = 0;
 int g_ggml_sycl_use_async_mem_op = 0;
 int g_ggml_sycl_use_async_mem_op_requested = 1;
@@ -298,6 +302,7 @@ static void ggml_check_sycl() try {
         g_ggml_sycl_fa_onednn_max_kv = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN_MAX_KV", 0);
         g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1);
         g_ggml_sycl_enable_fusion = ggml_sycl_get_env("GGML_SYCL_ENABLE_FUSION", 1);
+        g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1);
         g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
 
         g_ggml_sycl_dev2dev_memcpy = ggml_sycl_get_env("GGML_SYCL_DEV2DEV_MEMCPY", DEV2DEV_MEMCPY_SYCL);
@@ -392,6 +397,12 @@ static void ggml_check_sycl() try {
 
         GGML_LOG_INFO("  GGML_SYCL_ENABLE_FUSION: %d\n", g_ggml_sycl_enable_fusion);
 
+#if defined(__INTEL_LLVM_COMPILER)
+        GGML_LOG_INFO("  GGML_SYCL_ENABLE_ESIMD: %d\n", g_ggml_sycl_enable_esimd);
+#else
+        GGML_LOG_INFO("  GGML_SYCL_ENABLE_ESIMD: %d disabled by compile flag\n", g_ggml_sycl_enable_esimd);
+#endif
+
         GGML_LOG_INFO("  GGML_SYCL_PRIORITIZE_DMMV: %d\n", g_ggml_sycl_prioritize_dmmv);
 
         g_ggml_sycl_use_async_mem_op_requested = ggml_sycl_get_env("GGML_SYCL_USE_ASYNC_MEM_OP", 1);
@@ -2676,21 +2687,15 @@ inline void ggml_sycl_op_mul_mat_sycl(
         else
 #endif
         {
-            ggml_sycl_pool_alloc<sycl::half> dst_f16(ctx.pool(), row_diff * src1_ncols);
-
-            const sycl::half alpha_f16 = 1.0f;
-            const sycl::half beta_f16  = 0.0f;
+            const float alpha = 1.0f;
+            const float beta  = 0.0f;
             SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm(
                 *stream, oneapi::mkl::transpose::trans,
                 oneapi::mkl::transpose::nontrans, row_diff, src1_ncols, ne10,
-                &alpha_f16, src0_ptr, dpct::library_data_t::real_half, ne00,
-                src1_ptr, dpct::library_data_t::real_half, ne10, &beta_f16,
-                dst_f16.get(), dpct::library_data_t::real_half, ldc,
-                dpct::library_data_t::real_half)));
-            scope_op_debug_print scope_dbg_print(__func__, "/to_fp32_sycl", dst, /*num_src=*/2,
-                                                 " : converting dst to fp32");
-            const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16, dst);
-            to_fp32_sycl(dst_f16.get(), dst_dd_i, row_diff*src1_ncols, stream);
+                &alpha, src0_ptr, dpct::library_data_t::real_half, ne00,
+                src1_ptr, dpct::library_data_t::real_half, ne10, &beta,
+                dst_dd_i, dpct::library_data_t::real_float, ldc,
+                dpct::library_data_t::real_float)));
         }
     } else {
         ggml_sycl_pool_alloc<float> src0_ddq_as_f32(ctx.pool());
@@ -3740,6 +3745,22 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) {
     }
 }
 
+static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) {
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+    switch (type) {
+        case GGML_TYPE_Q3_K:
+        case GGML_TYPE_Q4_K:
+        case GGML_TYPE_Q6_K:
+            return true;
+        default:
+            return false;
+    }
+#else
+    GGML_UNUSED(type);
+    return false;
+#endif
+}
+
 static bool ggml_sycl_supports_dmmv(enum ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q1_0:
@@ -4443,19 +4464,22 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor
     use_mul_mat_q = use_mul_mat_q && (src1->ne[1] <= MMQ_MAX_BATCH_SIZE);
 #endif // SYCL_USE_XMX
 
-    // Dispatch becomes obscure with the reorder, MMVQ when the reorder optimization
-    // is enabled takes precedence over DMMV, the current if-else implementation
-    // requires disabling DMMV if both conditions are met
+    // When reorder is enabled, both ESIMD, MMVQ and DMMV kernels may be used. For
+    // best performance use ESIMD when supported, followed by MMVQ, and finally DMMV.
+    // But the reordered ESIMD path cannot be used without reordered MMVQ. A later
+    // multi-token call (ne[1] in 2..8) will take the MMVQ path and it would read the
+    // reordered bytes as if they were still the unreordered layout.
 
     if (!g_ggml_sycl_prioritize_dmmv && ((should_reorder_tensor(ctx, dst) &&
                                           ggml_sycl_supports_reorder_mmvq(src0->type)))) {
-      // Arc770 get benefit with Q4_0 by skipping it.
-      if (!(ggml_sycl_info().devices[ctx.device].hw_info.arch ==
-                gpu_arch::intel_gpu_acm_g10 &&
-            src0->type == GGML_TYPE_Q4_0)) {
-        use_dequantize_mul_mat_vec =
-            use_dequantize_mul_mat_vec && !use_mul_mat_vec_q;
-      }
+        bool use = g_ggml_sycl_enable_esimd && ggml_sycl_supports_reorder_esimd(src0->type);
+        // Arc770 get benefit with Q4_0 by skipping MMVQ path
+        if (!(ggml_sycl_info().devices[ctx.device].hw_info.arch ==
+                    gpu_arch::intel_gpu_acm_g10 &&
+                src0->type == GGML_TYPE_Q4_0)) {
+            use = use || !use_mul_mat_vec_q;
+        }
+        use_dequantize_mul_mat_vec = use_dequantize_mul_mat_vec && use;
     }
 
     if (!split && src0->type == GGML_TYPE_F16 && ggml_is_permuted(src0) && ggml_is_permuted(src1) && src1->ne[1] == 1) {
@@ -5422,11 +5446,17 @@ static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * syc
         }
 #endif
         if (node->op == GGML_OP_RMS_NORM &&
-            ggml_sycl_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL })) {
+            ggml_sycl_can_fuse(cgraph, i, { GGML_OP_RMS_NORM, GGML_OP_MUL }, {})) {
             ggml_sycl_op_rms_norm_fused(*sycl_ctx, node, cgraph->nodes[i + 1]);
             i++;
             continue;
         }
+        if (node->op == GGML_OP_UNARY &&
+            ggml_sycl_can_fuse(cgraph, i, { GGML_OP_UNARY, GGML_OP_MUL }, { ggml_get_unary_op(node) })) {
+            ggml_sycl_op_unary_mul_fused(*sycl_ctx, node, cgraph->nodes[i + 1]);
+            i++;
+            continue;
+        }
 
         bool ok = ggml_sycl_compute_forward(*sycl_ctx, node);
         if (!ok) {

tests/test-backend-ops.cpp

@@ -3695,6 +3695,117 @@ struct test_relu_sqr : public test_case {
     }
 };
 
+// GGML_OP_UNARY(SILU|SIGMOID|SOFTPLUS) + GGML_OP_MUL (fused operation).
+// `layout` and `tail` are used for fallback cases where fusion must be skipped
+struct test_unary_mul : public test_case {
+    const ggml_unary_op op;
+    const ggml_type type;
+    const std::array<int64_t, 4> ne;
+    const bool swap;          // unary result is the second MUL operand
+    const std::string layout; // operand layout, see build_graph()
+    const std::string tail;   // extra consumer past the MUL, see build_graph()
+
+    std::string op_desc(ggml_tensor * t) override {
+        GGML_UNUSED(t);
+        return std::string(ggml_unary_op_name(op)) + "_MUL";
+    }
+
+    bool run_whole_graph() override { return true; }
+
+    double max_nmse_err() override {
+        // the fused kernel elides the rounding of the unary result that the CPU chain
+        // performs; relax the tolerance to match that drift
+        switch (type) {
+            case GGML_TYPE_F16: return 5e-5;
+            default:            return 1e-7;
+        }
+    }
+
+    std::string vars() override {
+        return VARS_TO_STR5(type, ne, swap, layout, tail);
+    }
+
+    test_unary_mul(ggml_unary_op op,
+            ggml_type type = GGML_TYPE_F32,
+            std::array<int64_t, 4> ne = {128, 2, 2, 2},
+            bool swap = false,
+            std::string layout = "packed",
+            std::string tail = "")
+        : op(op), type(type), ne(ne), swap(swap), layout(std::move(layout)), tail(std::move(tail)) {}
+
+    // `ne` viewed out of a wider tensor: rows stay contiguous, but the stride exceeds the width
+    ggml_tensor * padded(ggml_context * ctx, const char * name, int64_t mul0, int64_t off0) {
+        std::array<int64_t, 4> ne_w = ne;
+        ne_w[0] *= mul0;
+        ggml_tensor * base = ggml_new_tensor(ctx, type, 4, ne_w.data());
+        ggml_set_name(base, name);
+        return ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3],
+                            base->nb[1], base->nb[2], base->nb[3], off0 * base->nb[0]);
+    }
+
+    ggml_tensor * build_graph(ggml_context * ctx) override {
+        ggml_tensor * a = nullptr; // unary source
+        ggml_tensor * b = nullptr; // other MUL operand
+
+        if (layout == "packed") {
+            a = ggml_new_tensor(ctx, type, 4, ne.data());
+            b = ggml_new_tensor(ctx, type, 4, ne.data());
+        } else if (layout == "pad_unary") {
+            a = padded(ctx, "a", 3, 0);
+            b = ggml_new_tensor(ctx, type, 4, ne.data());
+        } else if (layout == "pad_other") {
+            a = ggml_new_tensor(ctx, type, 4, ne.data());
+            b = padded(ctx, "b", 3, 0);
+        } else if (layout == "halves") {
+            // the shape the Conformer audio encoders build: one tensor split in two
+            std::array<int64_t, 4> ne_w = ne;
+            ne_w[0] *= 2;
+            ggml_tensor * base = ggml_new_tensor(ctx, type, 4, ne_w.data());
+            ggml_set_name(base, "base");
+            b = ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3], base->nb[1], base->nb[2], base->nb[3], 0);
+            a = ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3], base->nb[1], base->nb[2], base->nb[3],
+                             ne[0] * base->nb[0]);
+        } else if (layout == "strided_dim1") {
+            // contiguous rows but a strided dim 1: not ggml_is_contiguous_1, must not fuse
+            std::array<int64_t, 4> ne_w = ne;
+            ne_w[1] *= 3;
+            ggml_tensor * base = ggml_new_tensor(ctx, type, 4, ne_w.data());
+            ggml_set_name(base, "a");
+            a = ggml_view_4d(ctx, base, ne[0], ne[1], ne[2], ne[3], base->nb[1], base->nb[2], base->nb[3], 0);
+            b = ggml_new_tensor(ctx, type, 4, ne.data());
+        } else if (layout == "bcast") {
+            a = ggml_new_tensor(ctx, type, 4, ne.data());
+            b = ggml_new_tensor_4d(ctx, type, ne[0], 1, 1, 1);
+        } else {
+            GGML_ABORT("unknown layout %s", layout.c_str());
+        }
+        ggml_set_name(a, "a");
+        ggml_set_name(b, "b");
+
+        ggml_tensor * u = ggml_unary(ctx, a, op);
+        ggml_set_name(u, "unary");
+
+        // a broadcasting operand can only be the second one
+        const bool second = swap && layout != "bcast";
+        ggml_tensor * out = second ? ggml_mul(ctx, b, u) : ggml_mul(ctx, u, b);
+
+        if (tail == "reuse") {
+            // a second read of the unary result must block the fusion
+            ggml_set_name(out, "mul");
+            out = ggml_add(ctx, out, u);
+        } else if (tail == "consumer") {
+            // fusion still applies; catches a dispatcher that skips one node too many
+            ggml_set_name(out, "mul");
+            out = ggml_add(ctx, out, b);
+        } else if (!tail.empty()) {
+            GGML_ABORT("unknown tail %s", tail.c_str());
+        }
+        ggml_set_name(out, "out");
+
+        return out;
+    }
+};
+
 // SNAKE activation fusion: y = x + sin(a*x)^2 * inv_b
 // CUDA backend matches the naive 5-op chain (mul, sin, sqr, mul, add)
 // and dispatches a single fused kernel.
@@ -8065,6 +8176,25 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
         test_cases.emplace_back(new test_relu_sqr(type, { 5, 7, 11, 13 }));
     }
 
+    // fused unary + mul (gated activations that are not expressed as GGML_OP_GLU)
+    for (ggml_unary_op op : { GGML_UNARY_OP_SILU, GGML_UNARY_OP_SIGMOID, GGML_UNARY_OP_SOFTPLUS }) {
+        for (ggml_type type : { GGML_TYPE_F16, GGML_TYPE_F32 }) {
+            for (bool swap : { false, true }) {
+                test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, swap));
+            }
+            test_cases.emplace_back(new test_unary_mul(op, type, { 5, 7, 11, 13 }));
+            test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "pad_unary"));
+            // a view only stays out from between the two ops when the unary result is second
+            test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, true, "pad_other"));
+            test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, true, "halves"));
+            test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "packed", "consumer"));
+            // must not fuse
+            test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "strided_dim1"));
+            test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "bcast"));
+            test_cases.emplace_back(new test_unary_mul(op, type, { 128, 2, 2, 2 }, false, "packed", "reuse"));
+        }
+    }
+
     // SNAKE activation fusion: x + sin(a*x)^2 * inv_b
     for (ggml_type type : { GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_BF16 }) {
         test_cases.emplace_back(new test_snake_fuse(type, {   5,   7, 1, 1}));   // primes sub-block
@@ -8865,7 +8995,13 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
 
     for (ggml_type type_a : all_types) {
         for (int i = 1; i < 10; ++i) {
-            test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 16,  i, 256, { 1,  1}, {1, 1}));
+            test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 16,  i, 1*256, { 1,  1}, {1, 1}));
+            //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 12,  i, 2*256, { 2,  1}, {1, 1}));
+            //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 11,  i, 3*256, { 1,  3}, {5, 1}));
+            //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 13,  i, 4*256, { 2,  3}, {1, 1}));
+            //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 17,  i, 31*256, { 4,  1}, {1, 1}));
+            //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 18,  i, 32*256, { 1,  1}, {8, 1}));
+            //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 19,  i, 33*256, { 1,  1}, {1, 1}));
         }
     }
 

tools/server/server-common.cpp

@@ -60,6 +60,33 @@ json format_error_response(const std::string & message, const enum error_type ty
     };
 }
 
+//
+// server_slot_stats
+//
+
+json server_slot_stats::to_json() const {
+    json base = {
+        {"cache_n",                n_prompt_cached},
+
+        {"prompt_n",               n_prompt_processed},
+        {"prompt_ms",              t_prompt_ms()},
+        {"prompt_per_token_ms",    t_prompt_per_token_ms()},
+        {"prompt_per_second",      n_prompt_tps()},
+
+        {"predicted_n",            n_gen},
+        {"predicted_ms",           t_gen_ms()},
+        {"predicted_per_token_ms", t_gen_per_token_ms()},
+        {"predicted_per_second",   n_gen_tps()},
+    };
+
+    if (n_draft_tokens > 0) {
+        base["draft_n"]          = n_draft_tokens;
+        base["draft_n_accepted"] = n_draft_accepted;
+    }
+
+    return base;
+}
+
 //
 // random string / id
 //

tools/server/server-common.h

@@ -334,6 +334,160 @@ json format_response_rerank(
         std::vector<std::string> & texts,
         int top_n);
 
+//
+// stats and metrics
+//
+
+// shared between server_slot and server_task_result_*
+struct server_slot_stats {
+    uint64_t n_prompt_cached    = 0;
+    uint64_t n_prompt_processed = 0;
+    uint64_t n_gen              = 0;
+
+    // speculative decoding stats
+    // note: the per-position breakdown lives in server_slot, it is not needed in a task result
+    uint64_t n_draft_tokens      = 0;
+    uint64_t n_draft_accepted    = 0;
+    uint64_t n_draft_verif_steps = 0;
+
+    // these are absolute timestamps (in us)
+    // note: must be signed - they are subtracted before the later ones are set
+    int64_t t_start       = 0;
+    int64_t t_prompt_last = 0;
+    int64_t t_gen_last    = 0;
+
+    // can only move one direction: start -> prompt -> gen
+    void update_prompt_start() {
+        GGML_ASSERT(t_start == 0);
+        t_start = ggml_time_us();
+    }
+    void set_prompt_last(int64_t t_us) {
+        GGML_ASSERT(t_start > 0);
+        t_prompt_last = t_us;
+    }
+    void update_prompt_last() {
+        set_prompt_last(ggml_time_us());
+    }
+    void update_gen_last() {
+        GGML_ASSERT(t_prompt_last > 0);
+        t_gen_last = ggml_time_us();
+    }
+
+    // these are time durations
+    int64_t t_elapsed_us() const {
+        return ggml_time_us() - t_start;
+    }
+    double t_prompt_ms() const {
+        if (t_prompt_last == 0) {
+            return 0.0; // the prompt is not processed yet
+        }
+        return (t_prompt_last - t_start) / 1000.0;
+    }
+    int64_t t_gen_us() const {
+        if (t_gen_last == 0) {
+            return 0; // the generation is not started yet
+        }
+        // clamp to 1 us, the first token can land in the same us as t_prompt_last
+        return std::max<int64_t>(1, t_gen_last - t_prompt_last);
+    }
+    double t_gen_ms() const {
+        return t_gen_us() / 1000.0;
+    }
+
+    // number of decode steps spent on generation
+    // the first token is free, it comes from the logits of the last prompt batch
+    uint64_t n_gen_steps() const {
+        return n_gen > 0 ? n_gen - 1 : 0;
+    }
+
+    // other derived metrics
+    // note: all of them return 0.0 if the divisor is not known yet
+    double t_prompt_per_token_ms() const {
+        return n_prompt_processed > 0 ? t_prompt_ms() / n_prompt_processed : 0.0;
+    }
+    double t_gen_per_token_ms() const {
+        return n_gen_steps() > 0 ? t_gen_ms() / n_gen_steps() : 0.0;
+    }
+    double n_prompt_tps() const {
+        const double t_ms = t_prompt_ms();
+        return t_ms > 0.0 ? 1e3 / t_ms * n_prompt_processed : 0.0;
+    }
+    double n_gen_tps() const {
+        const double t_ms = t_gen_ms();
+        return t_ms > 0.0 ? 1e3 / t_ms * n_gen_steps() : 0.0;
+    }
+
+    // false if the slot never started, i.e. the task result carries no stats
+    bool is_set() const {
+        return t_start > 0;
+    }
+
+    json to_json() const;
+};
+
+// shared between server_context_impl and server_task_result_*
+// unlike server_slot_stats, server_metrics is server-global and cumulative, not tied to a slot
+struct server_metrics {
+    int64_t t_start = 0;
+
+    struct bucket {
+        uint64_t count = 0; // number of tokens
+        uint64_t steps = 0; // number of decode steps,
+                            // this excludes first generated token (logits from prompt batch)
+        uint64_t time  = 0; // in microseconds
+
+        // the rate uses the decode steps, so that "free" tokens do not inflate it
+        double n_per_second() const {
+            return time > 0 ? (double) steps / (double) time * 1e6 : 0.0;
+        }
+
+        void add(uint64_t n, uint64_t n_steps, uint64_t t_us) {
+            count += n;
+            steps += n_steps;
+            time  += t_us;
+        }
+    };
+
+    // these are reset by reset_bucket(), only the rate is read from them
+    bucket prompt_bucket;
+    bucket predict_bucket;
+
+    // metrics below are cumulative since the server started
+    bucket prompt; // only processed tokens, cached ones are counted separately below
+    bucket predict;
+
+    // tokens reused from the cache need no decode, so they only have a count
+    uint64_t n_prompt_cached = 0;
+
+    uint64_t n_tokens_max = 0;
+
+    uint64_t n_decode     = 0;
+    uint64_t n_busy_slots = 0;
+
+    uint64_t n_draft_tokens      = 0; // Total draft tokens generated
+    uint64_t n_draft_accepted    = 0; // Draft tokens actually accepted
+    uint64_t n_draft_verif_steps = 0; // Total draft token verification steps by the target model
+    std::vector<uint64_t> n_accepted_per_pos; // Accepted tokens per draft position
+
+    void init() {
+        t_start = ggml_time_us();
+    }
+
+    void reset_bucket() {
+        prompt_bucket  = {};
+        predict_bucket = {};
+    }
+
+    void add_prompt(uint64_t n_tokens, uint64_t t_us) {
+        prompt       .add(n_tokens, n_tokens, t_us);
+        prompt_bucket.add(n_tokens, n_tokens, t_us);
+    }
+
+    void add_prompt_cached(uint64_t n_tokens) {
+        n_prompt_cached += n_tokens;
+    }
+};
+
 //
 // other utils
 //

tools/server/server-task.cpp

@@ -10,6 +10,8 @@
 #include "speculative.h"
 #include "server-common.h"
 
+#include <sstream>
+
 using json = nlohmann::ordered_json;
 
 //
@@ -236,34 +238,6 @@ common_chat_msg task_result_state::update_chat_msg(
     return chat_msg;
 }
 
-//
-
-// result_timings
-//
-
-json result_timings::to_json() const {
-    json base = {
-        {"cache_n",                cache_n},
-
-        {"prompt_n",               prompt_n},
-        {"prompt_ms",              prompt_ms},
-        {"prompt_per_token_ms",    prompt_per_token_ms},
-        {"prompt_per_second",      prompt_per_second},
-
-        {"predicted_n",            predicted_n},
-        {"predicted_ms",           predicted_ms},
-        {"predicted_per_token_ms", predicted_per_token_ms},
-        {"predicted_per_second",   predicted_per_second},
-    };
-
-    if (draft_n > 0) {
-        base["draft_n"] = draft_n;
-        base["draft_n_accepted"] = draft_n_accepted;
-    }
-
-    return base;
-}
-
 //
 // result_prompt_progress
 //
@@ -382,7 +356,7 @@ json server_task_result_cmpl_final::to_json_non_oaicompat() {
         {"stop_type",           stop_type_to_str(stop)},
         {"stopping_word",       stopping_word},
         {"tokens_cached",       n_tokens_cached},
-        {"timings",             timings.to_json()},
+        {"timings",             stats.to_json()},
     };
     if (!stream && !probs_output.empty()) {
         res["completion_probabilities"] = completion_token_output::probs_vector_to_json(probs_output, post_sampling_probs);
@@ -432,8 +406,8 @@ json server_task_result_cmpl_final::to_json_oaicompat() {
     if (verbose) {
         res["__verbose"] = to_json_non_oaicompat();
     }
-    if (timings.prompt_n >= 0) {
-        res.push_back({"timings", timings.to_json()});
+    if (stats.is_set()) {
+        res.push_back({"timings", stats.to_json()});
     }
 
     return res;
@@ -480,8 +454,8 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat() {
     if (verbose) {
         res["__verbose"] = to_json_non_oaicompat();
     }
-    if (timings.prompt_n >= 0) {
-        res.push_back({"timings", timings.to_json()});
+    if (stats.is_set()) {
+        res.push_back({"timings", stats.to_json()});
     }
 
     return res;
@@ -541,8 +515,8 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat_stream() {
         });
     }
 
-    if (timings.prompt_n >= 0) {
-        deltas.back().push_back({"timings", timings.to_json()});
+    if (stats.is_set()) {
+        deltas.back().push_back({"timings", stats.to_json()});
     }
 
     // extra fields for debugging purposes
@@ -734,8 +708,8 @@ json server_task_result_cmpl_final::to_json_oaicompat_resp_stream() {
         }}
     });
 
-    if (timings.prompt_n >= 0) {
-        server_sent_events.back().at("data").push_back({"timings", timings.to_json()});
+    if (stats.is_set()) {
+        server_sent_events.back().at("data").push_back({"timings", stats.to_json()});
     }
 
     return server_sent_events;
@@ -1086,8 +1060,8 @@ json server_task_result_cmpl_partial::to_json_non_oaicompat() {
         {"tokens_evaluated", n_prompt_tokens},
     };
     // populate the timings object when needed (usually for the last response or with timings_per_token enabled)
-    if (timings.prompt_n > 0) {
-        res.push_back({"timings", timings.to_json()});
+    if (stats.is_set()) {
+        res.push_back({"timings", stats.to_json()});
     }
     if (is_progress) {
         res.push_back({"prompt_progress", progress.to_json()});
@@ -1126,8 +1100,8 @@ json server_task_result_cmpl_partial::to_json_oaicompat() {
     if (verbose) {
         res["__verbose"] = to_json_non_oaicompat();
     }
-    if (timings.prompt_n >= 0) {
-        res.push_back({"timings", timings.to_json()});
+    if (stats.is_set()) {
+        res.push_back({"timings", stats.to_json()});
     }
     if (is_progress) {
         res.push_back({"prompt_progress", progress.to_json()});
@@ -1180,8 +1154,8 @@ json server_task_result_cmpl_partial::to_json_oaicompat_chat() {
             };
         }
 
-        if (timings.prompt_n >= 0) {
-            last_json.push_back({"timings", timings.to_json()});
+        if (stats.is_set()) {
+            last_json.push_back({"timings", stats.to_json()});
         }
         if (is_progress) {
             last_json.push_back({"prompt_progress", progress.to_json()});
@@ -1330,8 +1304,8 @@ json server_task_result_cmpl_partial::to_json_oaicompat_resp() {
 
     if (!events.empty()) {
         json & data = events.back().at("data");
-        if (timings.prompt_n >= 0) {
-            data.push_back({"timings", timings.to_json()});
+        if (stats.is_set()) {
+            data.push_back({"timings", stats.to_json()});
         }
         if (is_progress) {
             data.push_back({"prompt_progress", progress.to_json()});
@@ -1539,34 +1513,104 @@ json server_task_result_error::to_json() {
 // server_task_result_metrics
 //
 json server_task_result_metrics::to_json() {
-    return json {
-        { "idle",                            n_idle_slots },
-        { "processing",                      n_processing_slots },
-        { "deferred",                        n_tasks_deferred },
-        { "t_start",                         t_start },
+    return slots_data;
+}
 
-        { "n_prompt_tokens_processed_total", n_prompt_tokens_processed_total },
-        { "t_tokens_generation_total",       t_tokens_generation_total },
-        { "n_tokens_predicted_total",        n_tokens_predicted_total },
-        { "t_prompt_processing_total",       t_prompt_processing_total },
-
-        { "n_tokens_max",                    n_tokens_max },
-
-        { "n_prompt_tokens_processed",       n_prompt_tokens_processed },
-        { "t_prompt_processing",             t_prompt_processing },
-        { "n_tokens_predicted",              n_tokens_predicted },
-        { "t_tokens_generation",             t_tokens_generation },
-
-        { "n_decode_total",                  n_decode_total },
-        { "n_busy_slots_total",              n_busy_slots_total },
-
-        { "n_draft_tokens_total",            n_draft_tokens_total },
-        { "n_draft_accepted_total",          n_draft_accepted_total },
-        { "n_draft_verif_steps_total",       n_draft_verif_steps_total },
-        { "n_accepted_per_pos_total",        n_accepted_per_pos_total },
-
-        { "slots",                           slots_data },
+// metrics definition: https://prometheus.io/docs/practices/naming/#metric-names
+std::string server_task_result_metrics::to_metrics() {
+    const std::vector<metric_item> counters = {
+        {
+            "prompt_tokens_total",
+            "Number of prompt tokens processed, excluding cached tokens",
+            (double) metrics.prompt.count
+        }, {
+            "prompt_tokens_cached_total",
+            "Number of prompt tokens reused from the cache",
+            (double) metrics.n_prompt_cached
+        }, {
+            "prompt_seconds_total",
+            "Total time spent processing prompts",
+            metrics.prompt.time / 1.e6
+        }, {
+            "tokens_predicted_total",
+            "Number of generation tokens processed",
+            (double) metrics.predict.count
+        }, {
+            "tokens_predicted_seconds_total",
+            "Total time spent generating tokens",
+            metrics.predict.time / 1.e6
+        }, {
+            "n_decode_total",
+            "Total number of llama_decode() calls, excluding speculative decoding and multimodal decoding",
+            (double) metrics.n_decode
+        }, {
+            "n_tokens_max",
+            "Largest observed sequence length (prompt + generation)",
+            (double) metrics.n_tokens_max
+        }, {
+            "spec_decode_num_draft_tokens_total",
+            "Speculative: Total draft tokens generated",
+            (double) metrics.n_draft_tokens
+        }, {
+            "spec_decode_num_accepted_tokens_total",
+            "Speculative: Total draft tokens accepted by the target model",
+            (double) metrics.n_draft_accepted
+        }, {
+            "spec_decode_num_drafts_total",
+            "Speculative: Total speculative decoding verification steps",
+            (double) metrics.n_draft_verif_steps
+        },
     };
+
+    const std::vector<metric_item> gauges = {
+        {
+            "prompt_tokens_seconds",
+            "Average prompt throughput in tokens/s",
+            metrics.prompt_bucket.n_per_second()
+        }, {
+            "predicted_tokens_seconds",
+            "Average generation throughput in tokens/s",
+            metrics.predict_bucket.n_per_second()
+        }, {
+            "requests_processing",
+            "Number of requests processing",
+            (double) n_processing_slots
+        }, {
+            "requests_deferred",
+            "Number of requests deferred",
+            (double) n_tasks_deferred
+        }, {
+            "n_busy_slots_per_decode",
+            "Average number of busy slots per llama_decode() call",
+            (double) metrics.n_busy_slots / std::max((double) metrics.n_decode, 1.0)
+        },
+    };
+
+    std::stringstream prometheus;
+
+    auto add_items = [&prometheus](const char * type, const std::vector<metric_item> & items) {
+        for (const auto & item : items) {
+            prometheus << "# HELP llamacpp:" << item.name << " " << item.description << "\n"
+                       << "# TYPE llamacpp:" << item.name << " " << type             << "\n"
+                       << "llamacpp:"        << item.name << " " << item.value       << "\n";
+        }
+    };
+
+    add_items("counter", counters);
+    add_items("gauge",   gauges);
+
+    // labeled counter: one time series per draft position
+    if (!metrics.n_accepted_per_pos.empty()) {
+        prometheus << "# HELP llamacpp:spec_decode_num_accepted_tokens_per_pos_total"
+                      " Accepted tokens per draft position\n"
+                   << "# TYPE llamacpp:spec_decode_num_accepted_tokens_per_pos_total counter\n";
+        for (size_t i = 0; i < metrics.n_accepted_per_pos.size(); i++) {
+            prometheus << "llamacpp:spec_decode_num_accepted_tokens_per_pos_total{position=\""
+                       << i << "\"} " << metrics.n_accepted_per_pos[i] << "\n";
+        }
+    }
+
+    return prometheus.str();
 }
 
 //

tools/server/server-task.h

@@ -259,26 +259,6 @@ struct server_task {
     }
 };
 
-struct result_timings {
-    int32_t cache_n = -1;
-
-    int32_t prompt_n = -1;
-    double prompt_ms = 0.0;
-    double prompt_per_token_ms = 0.0;
-    double prompt_per_second = 0.0;
-
-    int32_t predicted_n = -1;
-    double predicted_ms = 0.0;
-    double predicted_per_token_ms = 0.0;
-    double predicted_per_second = 0.0;
-
-    // Optional speculative metrics - only included when > 0
-    int32_t draft_n = 0;
-    int32_t draft_n_accepted = 0;
-
-    json to_json() const;
-};
-
 struct result_prompt_progress {
     int32_t total = 0;
     int32_t cache = 0;
@@ -343,7 +323,7 @@ struct server_task_result_cmpl_final : server_task_result {
 
     bool stream;
     bool include_usage;
-    result_timings timings;
+    server_slot_stats stats;
     std::string prompt;
 
     bool truncated;
@@ -425,7 +405,7 @@ struct server_task_result_cmpl_partial : server_task_result {
     bool is_begin = false; // whether to send 200 status to HTTP client (begin of SSE stream)
                            // ref: https://github.com/ggml-org/llama.cpp/pull/23884
     completion_token_output prob_output;
-    result_timings timings;
+    server_slot_stats stats;
     result_prompt_progress progress;
 
     // response formatting
@@ -510,38 +490,27 @@ struct server_task_result_error : server_task_result {
 };
 
 struct server_task_result_metrics : server_task_result {
+    // these are immediate stats, not accumulated (server_metrics is cumulative)
     int n_idle_slots;
     int n_processing_slots;
     int n_tasks_deferred;
-    int64_t t_start;
 
-    // TODO: somehow reuse server_metrics in the future, instead of duplicating the fields
-    uint64_t n_prompt_tokens_processed_total = 0;
-    uint64_t t_prompt_processing_total       = 0;
-    uint64_t n_tokens_predicted_total        = 0;
-    uint64_t t_tokens_generation_total       = 0;
-
-    uint64_t n_tokens_max = 0;
-
-    uint64_t n_prompt_tokens_processed = 0;
-    uint64_t t_prompt_processing       = 0;
-
-    uint64_t n_tokens_predicted  = 0;
-    uint64_t t_tokens_generation = 0;
-
-    uint64_t n_decode_total     = 0;
-    uint64_t n_busy_slots_total = 0;
-
-    uint64_t n_draft_tokens_total      = 0;
-    uint64_t n_draft_accepted_total    = 0;
-    uint64_t n_draft_verif_steps_total = 0;
-    std::vector<uint64_t> n_accepted_per_pos_total;
+    server_metrics metrics;
 
     // while we can also use std::vector<server_slot> this requires copying the slot object which can be quite messy
     // therefore, we use json to temporarily store the slot.to_json() result
     json slots_data = json::array();
 
+    // used by /slots API
     virtual json to_json() override;
+
+    // used by /metrics API
+    struct metric_item {
+        std::string name;
+        std::string description;
+        double value; // prometheus values are always float64
+    };
+    std::string to_metrics();
 };
 
 struct server_task_result_slot_save_load : server_task_result {

tools/server/tests/unit/test_metrics.py

@@ -0,0 +1,227 @@
+import pytest
+from utils import *
+
+server = ServerPreset.tinyllama2()
+
+
+@pytest.fixture(autouse=True)
+def create_server():
+    global server
+    server = ServerPreset.tinyllama2()
+    server.server_metrics = True
+
+
+def fetch_metrics(server: ServerProcess) -> str:
+    """get /metrics as raw prometheus text"""
+    res = server.make_request("GET", "/metrics")
+    assert res.status_code == 200
+    assert "Process-Start-Time-Unix" in res.headers
+    assert isinstance(res.body, str)
+    return res.body
+
+
+def parse_metrics(text: str) -> dict:
+    """parse the prometheus text format into {name: (type, value)}"""
+    out = {}
+    types = {}
+    for line in text.splitlines():
+        if line.startswith("# TYPE "):
+            _, _, name, kind = line.split(" ", 3)
+            types[name] = kind
+        elif line.startswith("llamacpp:") and "{" not in line:
+            name, value = line.split(" ", 1)
+            assert name in types, f"{name} has no # TYPE line"
+            out[name] = (types[name], float(value))
+    return out
+
+
+def test_metrics_disabled():
+    global server
+    server.server_metrics = False
+    server.start()
+    res = server.make_request("GET", "/metrics")
+    assert res.status_code == 501  # ERROR_TYPE_NOT_SUPPORTED
+
+
+def test_metrics_prometheus_format():
+    global server
+    server.start()
+    server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 8})
+
+    text = fetch_metrics(server)
+    metrics = parse_metrics(text)
+
+    expected_counters = [
+        "llamacpp:prompt_tokens_total",
+        "llamacpp:prompt_tokens_cached_total",
+        "llamacpp:prompt_seconds_total",
+        "llamacpp:tokens_predicted_total",
+        "llamacpp:tokens_predicted_seconds_total",
+        "llamacpp:n_decode_total",
+        "llamacpp:n_tokens_max",
+        "llamacpp:spec_decode_num_draft_tokens_total",
+        "llamacpp:spec_decode_num_accepted_tokens_total",
+        "llamacpp:spec_decode_num_drafts_total",
+    ]
+    expected_gauges = [
+        "llamacpp:prompt_tokens_seconds",
+        "llamacpp:predicted_tokens_seconds",
+        "llamacpp:requests_processing",
+        "llamacpp:requests_deferred",
+        "llamacpp:n_busy_slots_per_decode",
+    ]
+
+    for name in expected_counters:
+        assert metrics[name][0] == "counter"
+    for name in expected_gauges:
+        assert metrics[name][0] == "gauge"
+
+    # every metric must carry a help line
+    for name in expected_counters + expected_gauges:
+        assert f"# HELP {name} " in text
+
+    assert metrics["llamacpp:n_decode_total"][1] > 0
+    assert metrics["llamacpp:requests_processing"][1] == 0
+
+
+def test_metrics_prompt_processed_and_cached():
+    global server
+    server.n_slots = 1  # keep the prompt cache on a single slot
+    server.start()
+
+    prompt = "the quick brown fox jumps over the lazy dog"
+
+    n_processed = 0
+    n_cached = 0
+    for _ in range(2):
+        res = server.make_request("POST", "/completion", data={"prompt": prompt, "n_predict": 4})
+        assert res.status_code == 200
+        n_processed += res.body["timings"]["prompt_n"]
+        n_cached += res.body["timings"]["cache_n"]
+
+    # the second request must reuse the prompt of the first one
+    assert n_cached > 0
+
+    metrics = parse_metrics(fetch_metrics(server))
+
+    # cached tokens are counted apart, they cost no decode
+    assert metrics["llamacpp:prompt_tokens_total"][1] == n_processed
+    assert metrics["llamacpp:prompt_tokens_cached_total"][1] == n_cached
+
+
+def test_metrics_predicted_total_matches_requests():
+    global server
+    server.start()
+
+    n_predicted = 0
+    for n_predict in [1, 4, 16]:
+        res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": n_predict})
+        assert res.status_code == 200
+        n_predicted += res.body["timings"]["predicted_n"]
+
+    metrics = parse_metrics(fetch_metrics(server))
+    assert metrics["llamacpp:tokens_predicted_total"][1] == n_predicted
+
+
+def test_metrics_generation_rate_excludes_first_token():
+    global server
+    server.start()
+
+    # the first token comes from the logits of the last prompt batch, so it costs no decode step
+    res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 1})
+    timings = res.body["timings"]
+    assert timings["predicted_n"] == 1
+    assert timings["predicted_per_second"] == 0.0
+    assert timings["predicted_per_token_ms"] == 0.0
+
+    res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 16})
+    timings = res.body["timings"]
+    assert timings["predicted_n"] == 16
+    # the rate is over 15 decode steps, not 16 tokens
+    expected = 1e3 / timings["predicted_ms"] * 15
+    assert abs(timings["predicted_per_second"] - expected) < 1e-6
+
+
+@pytest.mark.parametrize("n_predict", [1, 8])
+def test_metrics_timings_are_finite(n_predict: int):
+    global server
+    server.start()
+    res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": n_predict})
+    timings = res.body["timings"]
+
+    # a null here means the server produced inf or nan
+    for key, value in timings.items():
+        assert value is not None, f"{key} is null"
+        assert value >= 0, f"{key} is negative"
+
+    assert timings["prompt_ms"] > 0
+    assert timings["prompt_per_token_ms"] > 0
+
+
+def test_metrics_timings_on_prompt_progress():
+    global server
+    server.start()
+
+    # a long prompt so that it is split over several batches (n_batch = 32)
+    prompt = "the quick brown fox jumps over the lazy dog " * 8
+    chunks = list(server.make_stream_request("POST", "/completion", data={
+        "prompt": prompt,
+        "n_predict": 4,
+        "stream": True,
+        "timings_per_token": True,
+        "return_progress": True,
+    }))
+
+    progress = [c for c in chunks if "prompt_progress" in c]
+    assert len(progress) > 1  # the prompt did not fit in a single batch
+
+    # the very first update is sent before any prompt token is decoded
+    first = progress[0]["timings"]
+    assert first["prompt_n"] == 0
+    assert first["prompt_ms"] == 0.0
+    assert first["predicted_n"] == 0
+    assert first["predicted_ms"] == 0.0
+
+    # timings must never go backwards, nor report bogus values
+    prompt_ms = 0.0
+    for chunk in progress:
+        timings = chunk["timings"]
+        for key, value in timings.items():
+            assert value is not None, f"{key} is null"
+            assert value >= 0, f"{key} is negative"
+        assert timings["prompt_ms"] >= prompt_ms
+        prompt_ms = timings["prompt_ms"]
+
+    assert prompt_ms > 0
+
+
+def test_metrics_slots_idle_after_completion():
+    global server
+    server.server_slots = True
+    server.start()
+    server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 8})
+
+    res = server.make_request("GET", "/slots")
+    assert res.status_code == 200
+    for slot in res.body:
+        assert slot["is_processing"] is False
+        if "next_token" in slot:
+            # the budget of the finished task must not leak into the idle slot
+            assert slot["next_token"][0]["n_remain"] == -1
+            assert slot["next_token"][0]["n_decoded"] == 0
+
+
+def test_metrics_embedding_prompt_is_counted():
+    global server
+    server = ServerPreset.bert_bge_small()
+    server.server_metrics = True
+    server.start()
+
+    res = server.make_request("POST", "/v1/embeddings", data={"input": ["hello world", "goodbye world"]})
+    assert res.status_code == 200
+
+    # embedding tasks never sample a token, but their prompt still costs a decode
+    metrics = parse_metrics(fetch_metrics(server))
+    assert metrics["llamacpp:prompt_tokens_total"][1] > 0
+    assert metrics["llamacpp:n_decode_total"][1] > 0
+    assert metrics["llamacpp:tokens_predicted_total"][1] == 0

tools/ui/src/app.d.ts

@@ -32,8 +32,8 @@ import type {
 	ApiRouterModelsStatusResponse,
 	ApiRouterModelsUnloadRequest,
 	ApiRouterModelsUnloadResponse,
-	// Chat types
 	ChatAttachmentDisplayItem,
+	// Chat types
 	ChatMessagePromptProgress,
 	ChatMessageSiblingInfo,
 	ChatMessageTimings,

tools/ui/src/app.html

@@ -2,6 +2,7 @@
 <html lang="en">
 	<head>
 		<meta charset="utf-8" />
+
 		<link rel="icon" href="favicon.ico" sizes="48x48" />
 		<link rel="icon" href="favicon.svg" sizes="any" type="image/svg+xml" />
 

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

@@ -3,7 +3,7 @@
 	import { X } from '@lucide/svelte';
 	import { ActionIcon } from '$lib/components/app';
 	import * as Tooltip from '$lib/components/ui/tooltip';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { mcpStore } from '$lib/stores';
 	import type { MCPResourceAttachment } from '$lib/types';
 	import { getResourceDisplayName, getResourceIcon } from '$lib/utils';
 

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

@@ -5,7 +5,7 @@
 		ChatAttachmentsPreviewNavButtons,
 		ChatAttachmentsPreviewThumbnailStrip
 	} from '$lib/components/app';
-	import { modelsStore } from '$lib/stores/models.svelte';
+	import { modelsStore } from '$lib/stores';
 	import {
 		createBase64DataUrl,
 		formatFileSize,

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

@@ -3,7 +3,7 @@
 	import {
 		ChatAttachmentsList,
 		ChatFormActions,
-		ChatFormContenteditable,
+		ChatFormContentEditable,
 		ChatFormFileInputInvisible,
 		ChatFormMcpResourcesList,
 		ChatFormPickers,
@@ -26,19 +26,16 @@
 		SpecialFileType
 	} from '$lib/enums';
 	import { useChatFormPickers } from '$lib/hooks/use-chat-form-pickers.svelte';
-	import { chatStore } from '$lib/stores/chat.svelte';
 	import {
-		activeConversation,
-		activeMessages,
+		chatStore,
 		conversationsStore,
-		pendingCwd
-	} from '$lib/stores/conversations.svelte';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
-	import { mcpHasResourceAttachments } from '$lib/stores/mcp-resources.svelte';
-	import { modelOptions, selectedModelId } from '$lib/stores/models.svelte';
-	import { isRouterMode } from '$lib/stores/server.svelte';
-	import { config } from '$lib/stores/settings.svelte';
-	import { toolsStore } from '$lib/stores/tools.svelte';
+		mcpResourceStore,
+		mcpStore,
+		modelsStore,
+		serverStore,
+		settingsStore,
+		toolsStore
+	} from '$lib/stores';
 	import type {
 		FileMentionEntry,
 		GetPromptResult,
@@ -143,7 +140,7 @@
 	// float above the box.
 	let mentionAnchor: HTMLDivElement | null = $state(null);
 
-	let cwd = $derived(activeConversation()?.cwd ?? pendingCwd());
+	let cwd = $derived(conversationsStore.activeConversation?.cwd ?? conversationsStore.pendingCwd);
 
 	const pickers = useChatFormPickers({
 		focusInput: refocusInput,
@@ -184,7 +181,7 @@
 	let isResourceDialogOpen = $state(false);
 	let preSelectedResourceUri = $state<string | undefined>(undefined);
 
-	let currentConfig = $derived(config());
+	let currentConfig = $derived(settingsStore.config);
 
 	let pasteLongTextToFileLength = $derived.by(() => {
 		const n = Number(currentConfig.pasteLongTextToFileLen);
@@ -192,18 +189,18 @@
 		return Number.isNaN(n) ? Number(SETTING_CONFIG_DEFAULT.pasteLongTextToFileLen) : n;
 	});
 
-	let isRouter = $derived(isRouterMode());
+	let isRouter = $derived(serverStore.isRouterMode);
 	let conversationModel = $derived(
-		chatStore.getConversationModel(activeMessages() as DatabaseMessage[])
+		chatStore.getConversationModel(conversationsStore.activeMessages as DatabaseMessage[])
 	);
 	let activeModelId = $derived.by(() => {
-		const options = modelOptions();
+		const options = modelsStore.models;
 
 		if (!isRouter) {
 			return options.length > 0 ? options[0].model : null;
 		}
 
-		const selectedId = selectedModelId();
+		const selectedId = modelsStore.selectedModelId;
 
 		if (selectedId) {
 			const model = options.find((m) => m.id === selectedId);
@@ -220,7 +217,9 @@
 		return null;
 	});
 
-	let hasModelSelected = $derived(!isRouter || !!conversationModel || !!selectedModelId());
+	let hasModelSelected = $derived(
+		!isRouter || !!conversationModel || !!modelsStore.selectedModelId
+	);
 	let hasLoadingAttachments = $derived(uploadedFiles.some((f) => f.isLoading));
 	let hasAttachments = $derived(
 		(attachments && attachments.length > 0) || (uploadedFiles && uploadedFiles.length > 0)
@@ -605,7 +604,7 @@
 			class="flex-column relative min-h-12 items-center rounded-4xl md:rounded-3xl py-2 pb-2.25 shadow-sm transition-all focus-within:shadow-md md:py-3!"
 		>
 			{#if useContenteditable}
-				<ChatFormContenteditable
+				<ChatFormContentEditable
 					class="px-5 py-1.5 md:pt-0 mb-0.5"
 					bind:this={inputRef}
 					bind:value
@@ -634,7 +633,7 @@
 				/>
 			{/if}
 
-			{#if mcpHasResourceAttachments()}
+			{#if mcpResourceStore.hasAttachments}
 				<ChatFormMcpResourcesList
 					class="mb-3"
 					onResourceClick={(uri) => {

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte

@@ -15,37 +15,16 @@
 		ICON_CLASS_DEFAULT,
 		TOOLTIP_DELAY_DURATION
 	} from '$lib/constants';
+	import { getChatFormActionsContext } from '$lib/contexts';
 	import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte';
 
 	interface Props {
 		class?: string;
-		disabled?: boolean;
-		hasAudioModality?: boolean;
-		hasVideoModality?: boolean;
-		hasVisionModality?: boolean;
-		hasMcpPromptsSupport?: boolean;
-		hasMcpResourcesSupport?: boolean;
-		onFileUpload?: () => void;
-		onSystemPromptClick?: () => void;
-		onMcpPromptClick?: () => void;
-		onMcpSettingsClick?: () => void;
-		onMcpResourcesClick?: () => void;
 	}
 
-	let {
-		class: className = '',
-		disabled = false,
-		hasAudioModality = false,
-		hasMcpPromptsSupport = false,
-		hasMcpResourcesSupport = false,
-		hasVideoModality = false,
-		hasVisionModality = false,
-		onFileUpload,
-		onMcpPromptClick,
-		onMcpResourcesClick,
-		onMcpSettingsClick,
-		onSystemPromptClick
-	}: Props = $props();
+	let { class: className = '' }: Props = $props();
+
+	const chatFormActions = getChatFormActionsContext();
 
 	let dropdownOpen = $state(false);
 	// The system message action moves focus to the message editor, so the menu
@@ -54,18 +33,23 @@
 
 	function handleMcpSettingsClick() {
 		dropdownOpen = false;
-		onMcpSettingsClick?.();
+		chatFormActions.onMcpSettingsClick?.();
 	}
 
 	const attachmentMenu = useAttachmentMenu(
 		() => ({
-			hasAudioModality,
-			hasMcpPromptsSupport,
-			hasMcpResourcesSupport,
-			hasVideoModality,
-			hasVisionModality
+			hasAudioModality: chatFormActions.hasAudioModality,
+			hasMcpPromptsSupport: chatFormActions.hasMcpPromptsSupport,
+			hasMcpResourcesSupport: chatFormActions.hasMcpResourcesSupport,
+			hasVideoModality: chatFormActions.hasVideoModality,
+			hasVisionModality: chatFormActions.hasVisionModality
+		}),
+		() => ({
+			onFileUpload: chatFormActions.onFileUpload,
+			onMcpPromptClick: chatFormActions.onMcpPromptClick,
+			onMcpResourcesClick: chatFormActions.onMcpResourcesClick,
+			onSystemPromptClick: chatFormActions.onSystemPromptClick
 		}),
-		() => ({ onFileUpload, onMcpPromptClick, onMcpResourcesClick, onSystemPromptClick }),
 		() => {
 			dropdownOpen = false;
 		}
@@ -85,7 +69,7 @@
 							buttonVariants({ variant: 'secondary' }),
 							'file-upload-button h-8 w-8 cursor-pointer rounded-full p-0'
 						)}
-						{disabled}
+						disabled={chatFormActions.disabled}
 					>
 						<span class="sr-only">{ATTACHMENT_TOOLTIP_TEXT}</span>
 
@@ -162,7 +146,7 @@
 				class="flex cursor-pointer items-center gap-2"
 				onclick={() => {
 					suppressCloseAutoFocus = true;
-					onSystemPromptClick?.();
+					chatFormActions.onSystemPromptClick?.();
 				}}
 			>
 				<MessageSquare class={ICON_CLASS_DEFAULT} />
@@ -174,12 +158,12 @@
 
 			<ChatFormActionAddMcpServersSubmenu onMcpSettingsClick={handleMcpSettingsClick} />
 
-			{#if hasMcpPromptsSupport}
+			{#if chatFormActions.hasMcpPromptsSupport}
 				<DropdownMenu.Separator />
 
 				<DropdownMenu.Item
 					class="flex cursor-pointer items-center gap-2"
-					onclick={onMcpPromptClick}
+					onclick={chatFormActions.onMcpPromptClick}
 				>
 					<Zap class={ICON_CLASS_DEFAULT} />
 
@@ -187,10 +171,10 @@
 				</DropdownMenu.Item>
 			{/if}
 
-			{#if hasMcpResourcesSupport}
+			{#if chatFormActions.hasMcpResourcesSupport}
 				<DropdownMenu.Item
 					class="flex cursor-pointer items-center gap-2"
-					onclick={onMcpResourcesClick}
+					onclick={chatFormActions.onMcpResourcesClick}
 				>
 					<FolderOpen class={ICON_CLASS_DEFAULT} />
 

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

@@ -6,8 +6,7 @@
 	import { Switch } from '$lib/components/ui/switch';
 	import { ICON_CLASS_DEFAULT, ROUTES } from '$lib/constants';
 	import { HealthCheckStatus } from '$lib/enums';
-	import { conversationsStore } from '$lib/stores/conversations.svelte';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { conversationsStore, mcpStore } from '$lib/stores';
 	import type { MCPServerSettingsEntry } from '$lib/types';
 
 	interface Props {

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte

@@ -19,44 +19,23 @@
 		ICON_CLASS_DEFAULT,
 		TOOLTIP_DELAY_DURATION
 	} from '$lib/constants';
+	import { getChatFormActionsContext } from '$lib/contexts';
 	import { HealthCheckStatus } from '$lib/enums';
 	import { AttachmentAction } from '$lib/enums/attachment.enums';
 	import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte';
 	import { useReasoningMenu } from '$lib/hooks/use-reasoning-menu.svelte';
 	import { useToolsPanel } from '$lib/hooks/use-tools-panel.svelte';
-	import { conversationsStore } from '$lib/stores/conversations.svelte';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { conversationsStore, mcpStore } from '$lib/stores';
 	import type { Snippet } from 'svelte';
 
 	interface Props {
 		class?: string;
-		disabled?: boolean;
-		hasAudioModality?: boolean;
-		hasVideoModality?: boolean;
-		hasVisionModality?: boolean;
-		hasMcpPromptsSupport?: boolean;
-		hasMcpResourcesSupport?: boolean;
-		onFileUpload?: () => void;
-		onSystemPromptClick?: () => void;
-		onMcpPromptClick?: () => void;
-		onMcpResourcesClick?: () => void;
 		trigger: Snippet<[{ disabled: boolean; onclick?: () => void }]>;
 	}
 
-	let {
-		class: className = '',
-		disabled = false,
-		hasAudioModality = false,
-		hasMcpPromptsSupport = false,
-		hasMcpResourcesSupport = false,
-		hasVideoModality = false,
-		hasVisionModality = false,
-		onFileUpload,
-		onMcpPromptClick,
-		onMcpResourcesClick,
-		onSystemPromptClick,
-		trigger
-	}: Props = $props();
+	let { class: className = '', trigger }: Props = $props();
+
+	const chatFormActions = getChatFormActionsContext();
 
 	let sheetOpen = $state(false);
 	let reasoningExpanded = $state(false);
@@ -66,13 +45,18 @@
 
 	const attachmentMenu = useAttachmentMenu(
 		() => ({
-			hasAudioModality,
-			hasMcpPromptsSupport,
-			hasMcpResourcesSupport,
-			hasVideoModality,
-			hasVisionModality
+			hasAudioModality: chatFormActions.hasAudioModality,
+			hasMcpPromptsSupport: chatFormActions.hasMcpPromptsSupport,
+			hasMcpResourcesSupport: chatFormActions.hasMcpResourcesSupport,
+			hasVideoModality: chatFormActions.hasVideoModality,
+			hasVisionModality: chatFormActions.hasVisionModality
+		}),
+		() => ({
+			onFileUpload: chatFormActions.onFileUpload,
+			onMcpPromptClick: chatFormActions.onMcpPromptClick,
+			onMcpResourcesClick: chatFormActions.onMcpResourcesClick,
+			onSystemPromptClick: chatFormActions.onSystemPromptClick
 		}),
-		() => ({ onFileUpload, onMcpPromptClick, onMcpResourcesClick, onSystemPromptClick }),
 		() => {
 			sheetOpen = false;
 		}
@@ -92,7 +76,7 @@
 
 <div class="flex items-center gap-1 {className}">
 	<Sheet.Root bind:open={sheetOpen}>
-		{@render trigger({ disabled, onclick: () => (sheetOpen = true) })}
+		{@render trigger({ disabled: chatFormActions.disabled, onclick: () => (sheetOpen = true) })}
 
 		<Sheet.Content side="bottom" class="max-h-[85vh] gap-0 overflow-y-auto">
 			<Sheet.Header>
@@ -351,7 +335,7 @@
 					<span>System Message</span>
 				</button>
 
-				{#if hasMcpPromptsSupport}
+				{#if chatFormActions.hasMcpPromptsSupport}
 					<button
 						type="button"
 						class={sheetItemClass}
@@ -363,7 +347,7 @@
 					</button>
 				{/if}
 
-				{#if hasMcpResourcesSupport}
+				{#if chatFormActions.hasMcpResourcesSupport}
 					<button
 						type="button"
 						class={sheetItemClass}

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte

@@ -6,8 +6,7 @@
 	import * as Tooltip from '$lib/components/ui/tooltip';
 	import { CLI_FLAGS, ICON_CLASS_DEFAULT } from '$lib/constants';
 	import { useToolsPanel } from '$lib/hooks/use-tools-panel.svelte';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
-	import { toolsStore } from '$lib/stores/tools.svelte';
+	import { mcpStore, toolsStore } from '$lib/stores';
 
 	const toolsPanel = useToolsPanel();
 	const hasMcpServersAvailable = $derived(mcpStore.getServers().length > 0);

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionsAdd.svelte

@@ -2,66 +2,15 @@
 	import ChatFormActionAddButton from './ChatFormActionAddButton.svelte';
 	import ChatFormActionAddDropdown from './ChatFormActionAddDropdown.svelte';
 	import ChatFormActionAddSheet from './ChatFormActionAddSheet.svelte';
-	import { isMobile } from '$lib/stores/viewport.svelte';
-
-	interface Props {
-		disabled?: boolean;
-		hasAudioModality?: boolean;
-		hasVideoModality?: boolean;
-		hasMcpPromptsSupport?: boolean;
-		hasMcpResourcesSupport?: boolean;
-		hasVisionModality?: boolean;
-		onFileUpload?: () => void;
-		onMcpPromptClick?: () => void;
-		onMcpResourcesClick?: () => void;
-		onMcpSettingsClick?: () => void;
-		onSystemPromptClick?: () => void;
-	}
-
-	let {
-		disabled = false,
-		hasAudioModality = false,
-		hasMcpPromptsSupport = false,
-		hasMcpResourcesSupport = false,
-		hasVideoModality = false,
-		hasVisionModality = false,
-		onFileUpload,
-		onMcpPromptClick,
-		onMcpResourcesClick,
-		onMcpSettingsClick,
-		onSystemPromptClick
-	}: Props = $props();
+	import { isMobile } from '$lib/stores';
 </script>
 
 {#if isMobile.current}
-	<ChatFormActionAddSheet
-		{disabled}
-		{hasAudioModality}
-		{hasVideoModality}
-		{hasVisionModality}
-		{hasMcpPromptsSupport}
-		{hasMcpResourcesSupport}
-		{onFileUpload}
-		{onSystemPromptClick}
-		{onMcpPromptClick}
-		{onMcpResourcesClick}
-	>
+	<ChatFormActionAddSheet>
 		{#snippet trigger({ disabled, onclick })}
 			<ChatFormActionAddButton {disabled} {onclick} />
 		{/snippet}
 	</ChatFormActionAddSheet>
 {:else}
-	<ChatFormActionAddDropdown
-		{disabled}
-		{hasAudioModality}
-		{hasVideoModality}
-		{hasVisionModality}
-		{hasMcpPromptsSupport}
-		{hasMcpResourcesSupport}
-		{onFileUpload}
-		{onMcpPromptClick}
-		{onMcpResourcesClick}
-		{onMcpSettingsClick}
-		{onSystemPromptClick}
-	/>
+	<ChatFormActionAddDropdown />
 {/if}

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte

@@ -1,15 +1,6 @@
 <script lang="ts">
 	import { ModelsSelectorDropdown, ModelsSelectorSheet } from '$lib/components/app';
-	import { chatStore } from '$lib/stores/chat.svelte';
-	import { activeMessages } from '$lib/stores/conversations.svelte';
-	import {
-		modelOptions,
-		modelsStore,
-		selectedModelId,
-		selectedModelName
-	} from '$lib/stores/models.svelte';
-	import { isRouterMode, serverError } from '$lib/stores/server.svelte';
-	import { isMobile } from '$lib/stores/viewport.svelte';
+	import { chatStore, conversationsStore, isMobile, modelsStore, serverStore } from '$lib/stores';
 
 	interface Props {
 		disabled?: boolean;
@@ -35,17 +26,17 @@
 		useGlobalSelection = false
 	}: Props = $props();
 
-	let isRouter = $derived(isRouterMode());
-	let isOffline = $derived(!!serverError());
+	let isRouter = $derived(serverStore.isRouterMode);
+	let isOffline = $derived(!!serverStore.error);
 
 	let conversationModel = $derived(
-		chatStore.getConversationModel(activeMessages() as DatabaseMessage[])
+		chatStore.getConversationModel(conversationsStore.activeMessages as DatabaseMessage[])
 	);
 
 	let lastSyncedConversationModel: string | null = null;
 
 	let selectorModel = $derived.by(() => {
-		const storeModel = selectedModelName();
+		const storeModel = modelsStore.selectedModelName;
 
 		if (storeModel && storeModel !== conversationModel) {
 			return storeModel;
@@ -60,7 +51,7 @@
 
 	$effect(() => {
 		if (conversationModel && conversationModel !== lastSyncedConversationModel) {
-			if (modelOptions().some((m) => m.model === conversationModel)) {
+			if (modelsStore.models.some((m) => m.model === conversationModel)) {
 				modelsStore.selectedModelName = conversationModel;
 				modelsStore.selectModelByName(conversationModel);
 			} else {
@@ -73,24 +64,24 @@
 			isRouter &&
 			!modelsStore.selectedModelId &&
 			modelsStore.loadedModelIds.length > 0 &&
-			activeMessages().length > 0 &&
+			conversationsStore.activeMessages.length > 0 &&
 			!conversationModel
 		) {
 			lastSyncedConversationModel = null;
-			const first = modelOptions().find((m) => modelsStore.loadedModelIds.includes(m.model));
+			const first = modelsStore.models.find((m) => modelsStore.loadedModelIds.includes(m.model));
 
 			if (first) modelsStore.selectModelById(first.id);
 		}
 	});
 
 	let activeModelId = $derived.by(() => {
-		const options = modelOptions();
+		const options = modelsStore.models;
 
 		if (!isRouter) {
 			return options.length > 0 ? options[0].model : null;
 		}
 
-		const selectedId = selectedModelId();
+		const selectedId = modelsStore.selectedModelId;
 
 		if (selectedId) {
 			const model = options.find((m) => m.id === selectedId);
@@ -140,21 +131,23 @@
 	});
 
 	$effect(() => {
-		hasModelSelected = !isRouter || !!conversationModel || !!selectedModelId();
+		hasModelSelected = !isRouter || !!conversationModel || !!modelsStore.selectedModelId;
 	});
 
 	$effect(() => {
 		if (!isRouter) {
 			isSelectedModelInCache = true;
 		} else if (conversationModel) {
-			isSelectedModelInCache = modelOptions().some((option) => option.model === conversationModel);
+			isSelectedModelInCache = modelsStore.models.some(
+				(option) => option.model === conversationModel
+			);
 		} else {
-			const currentModelId = selectedModelId();
+			const currentModelId = modelsStore.selectedModelId;
 
 			if (!currentModelId) {
 				isSelectedModelInCache = false;
 			} else {
-				isSelectedModelInCache = modelOptions().some((option) => option.id === currentModelId);
+				isSelectedModelInCache = modelsStore.models.some((option) => option.id === currentModelId);
 			}
 		}
 	});

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

@@ -11,16 +11,10 @@
 	} from '$lib/components/app';
 	import { Button } from '$lib/components/ui/button';
 	import { ICON_CLASS_DEFAULT, ROUTES } from '$lib/constants';
+	import { setChatFormActionsContext } from '$lib/contexts';
 	import { FileTypeCategory, MessageRole } from '$lib/enums';
 	import { ChatService } from '$lib/services';
-	import {
-		activeProcessingState,
-		isChatStreaming,
-		isLoading as chatIsLoading
-	} from '$lib/stores/chat.svelte';
-	import { activeMessages, conversationsStore } from '$lib/stores/conversations.svelte';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
-	import { config } from '$lib/stores/settings.svelte';
+	import { chatStore, conversationsStore, mcpStore, settingsStore } from '$lib/stores';
 	import { getFileTypeCategory } from '$lib/utils';
 
 	interface Props {
@@ -61,7 +55,7 @@
 		uploadedFiles = []
 	}: Props = $props();
 
-	let currentConfig = $derived(config());
+	let currentConfig = $derived(settingsStore.config);
 
 	let hasMcpPromptsSupport = $derived.by(() => {
 		const perChatOverrides = conversationsStore.getAllMcpServerOverrides();
@@ -103,7 +97,7 @@
 	let hasProcessedTokens = $derived.by(() => {
 		if (!page.params.id) return false;
 
-		const messages = activeMessages() as DatabaseMessage[];
+		const messages = conversationsStore.activeMessages as DatabaseMessage[];
 
 		let totalHistoricalTokens = 0;
 
@@ -125,9 +119,9 @@
 
 		if (totalHistoricalTokens > 0) return true;
 
-		if (!chatIsLoading() && !isChatStreaming()) return false;
+		if (!chatStore.isLoading && !chatStore.isStreaming()) return false;
 
-		const processingState = activeProcessingState();
+		const processingState = chatStore.activeProcessingState;
 
 		if (!processingState) return false;
 
@@ -139,6 +133,42 @@
 
 		return livePromptTokens > 0 || liveOutputTokens > 0;
 	});
+
+	setChatFormActionsContext({
+		get disabled() {
+			return disabled;
+		},
+		get hasAudioModality() {
+			return hasAudioModality;
+		},
+		get hasMcpPromptsSupport() {
+			return hasMcpPromptsSupport;
+		},
+		get hasMcpResourcesSupport() {
+			return hasMcpResourcesSupport;
+		},
+		get hasVideoModality() {
+			return hasVideoModality;
+		},
+		get hasVisionModality() {
+			return hasVisionModality;
+		},
+		get onFileUpload() {
+			return onFileUpload;
+		},
+		get onMcpPromptClick() {
+			return onMcpPromptClick;
+		},
+		get onMcpResourcesClick() {
+			return onMcpResourcesClick;
+		},
+		get onMcpSettingsClick() {
+			return () => goto(ROUTES.MCP_SERVERS);
+		},
+		get onSystemPromptClick() {
+			return onSystemPromptClick;
+		}
+	});
 </script>
 
 <div
@@ -147,19 +177,7 @@
 >
 	{#if showAddButton}
 		<div class="mr-auto flex items-center gap-2">
-			<ChatFormActionsAdd
-				{disabled}
-				{hasAudioModality}
-				{hasVideoModality}
-				{hasVisionModality}
-				{hasMcpPromptsSupport}
-				{hasMcpResourcesSupport}
-				{onFileUpload}
-				{onSystemPromptClick}
-				{onMcpPromptClick}
-				{onMcpResourcesClick}
-				onMcpSettingsClick={() => goto(ROUTES.MCP_SERVERS)}
-			/>
+			<ChatFormActionsAdd />
 		</div>
 	{/if}
 

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

@@ -1,8 +1,9 @@
 <script lang="ts">
 	import { CODE_BLOCK } from '$lib/constants';
 	import { ColorMode } from '$lib/enums';
-	import { isMobile } from '$lib/stores/viewport.svelte';
-	import type { ContentToken, SourceHistoryEntry } from '$lib/utils';
+	import { isMobile } from '$lib/stores';
+	import type { ContentEditableToken } from '$lib/types';
+	import type { SourceHistoryEntry } from '$lib/utils';
 	import {
 		badgeAwareWordJump,
 		buildFragment,
@@ -63,7 +64,7 @@
 		rootElement.dataset.empty = source.length === 0 ? 'true' : 'false';
 	}
 
-	function renderTokens(tokens: ContentToken[]) {
+	function renderTokens(tokens: ContentEditableToken[]) {
 		if (!rootElement) return;
 
 		const caret = rangeToTextOffset(rootElement, safeRange());

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

@@ -1,32 +1,32 @@
 <script lang="ts">
 	import ContextGaugeDial from './ContextGaugeDial.svelte';
 	import { useContextGauge } from '$lib/hooks/use-context-gauge.svelte';
-	import { chatStore, isChatStreaming, isLoading } from '$lib/stores/chat.svelte';
 	import {
+		chatStore,
+		conversationsStore,
 		gaugeTriggerClick,
 		gaugeTriggerEnter,
 		gaugeTriggerKeydown,
 		gaugeTriggerLeave,
 		gaugeTriggerPointerDown
-	} from '$lib/stores/context-gauge-popup.svelte';
-	import { activeConversation, activeMessages } from '$lib/stores/conversations.svelte';
+	} from '$lib/stores';
 	import { untrack } from 'svelte';
 
 	const gauge = useContextGauge();
 
 	$effect(() => {
-		const conv = activeConversation();
+		const conv = conversationsStore.activeConversation;
 
 		untrack(() => chatStore.setActiveProcessingConversation(conv?.id ?? null));
 	});
 
 	$effect(() => {
-		const conv = activeConversation();
-		const messages = activeMessages() as DatabaseMessage[];
+		const conv = conversationsStore.activeConversation;
+		const messages = conversationsStore.activeMessages as DatabaseMessage[];
 
 		if (!conv) return;
 
-		if (isLoading() || isChatStreaming()) return;
+		if (chatStore.isLoading || chatStore.isStreaming()) return;
 
 		if (messages.length === 0) {
 			untrack(() => chatStore.clearProcessingState(conv.id));

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormContextGauge/ContextGaugeDial.svelte

@@ -1,6 +1,6 @@
 <script lang="ts">
-	import type { ColorLevel } from './context-gauge';
 	import { colorLevelTextClass } from './context-gauge';
+	import type { ColorLevel } from '$lib/enums';
 
 	interface Props {
 		percent: number | null;

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormContextGauge/ContextGaugePopup.svelte

@@ -3,12 +3,7 @@
 	import ContextGaugeDetails from './ContextGaugeDetails.svelte';
 	import ContextGaugeLoadModel from './ContextGaugeLoadModel.svelte';
 	import { useContextGauge } from '$lib/hooks/use-context-gauge.svelte';
-	import {
-		gaugeCardEnter,
-		gaugeCardLeave,
-		gaugePopup,
-		gaugePopupClose
-	} from '$lib/stores/context-gauge-popup.svelte';
+	import { gaugeCardEnter, gaugeCardLeave, gaugePopup, gaugePopupClose } from '$lib/stores';
 	import { formatParameters } from '$lib/utils/formatters';
 
 	const gauge = useContextGauge();
@@ -92,7 +87,7 @@
 						<span class={colorLevelTextClass(gauge.colorLevel)}>{gauge.contextPercent}%</span> used
 					</span>
 					<span>
-						{formatParameters((gauge.contextTotal ?? 0) - gauge.contextUsed)} remaining
+						{formatParameters(gauge.contextAvailable ?? 0)} remaining
 					</span>
 				</div>
 			{:else}

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormContextGauge/context-gauge.ts

@@ -1,25 +1,25 @@
-export type ColorLevel = 'ok' | 'warning' | 'critical' | 'neutral';
+import { ColorLevel } from '$lib/enums';
 
 const WARNING_THRESHOLD = 80;
 const CRITICAL_THRESHOLD = 95;
 
 export function colorLevelFromPercent(percent: number | null): ColorLevel {
-	if (percent === null) return 'neutral';
+	if (percent === null) return ColorLevel.NEUTRAL;
 
-	if (percent >= CRITICAL_THRESHOLD) return 'critical';
+	if (percent >= CRITICAL_THRESHOLD) return ColorLevel.CRITICAL;
 
-	if (percent >= WARNING_THRESHOLD) return 'warning';
+	if (percent >= WARNING_THRESHOLD) return ColorLevel.WARNING;
 
-	return 'ok';
+	return ColorLevel.OK;
 }
 
 export function colorLevelTextClass(level: ColorLevel): string {
 	switch (level) {
-		case 'critical':
+		case ColorLevel.CRITICAL:
 			return 'text-red-400';
-		case 'warning':
+		case ColorLevel.WARNING:
 			return 'text-amber-400';
-		case 'ok':
+		case ColorLevel.OK:
 			return 'text-muted-foreground';
 		default:
 			return 'text-muted-foreground';
@@ -28,11 +28,11 @@ export function colorLevelTextClass(level: ColorLevel): string {
 
 export function colorLevelBgClass(level: ColorLevel): string {
 	switch (level) {
-		case 'critical':
+		case ColorLevel.CRITICAL:
 			return 'bg-red-500';
-		case 'warning':
+		case ColorLevel.WARNING:
 			return 'bg-amber-500';
-		case 'ok':
+		case ColorLevel.OK:
 			return 'bg-green-500';
 		default:
 			return 'bg-muted';

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

@@ -3,11 +3,7 @@
 		ChatAttachmentsListItemMcpResource,
 		HorizontalScrollCarousel
 	} from '$lib/components/app';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
-	import {
-		mcpHasResourceAttachments,
-		mcpResourceAttachments
-	} from '$lib/stores/mcp-resources.svelte';
+	import { mcpResourceStore, mcpStore } from '$lib/stores';
 
 	interface Props {
 		class?: string;
@@ -16,8 +12,8 @@
 
 	let { class: className, onResourceClick }: Props = $props();
 
-	const attachments = $derived(mcpResourceAttachments());
-	const hasAttachments = $derived(mcpHasResourceAttachments());
+	const attachments = $derived(mcpResourceStore.attachments);
+	const hasAttachments = $derived(mcpResourceStore.hasAttachments);
 
 	function handleRemove(attachmentId: string) {
 		mcpStore.removeResourceAttachment(attachmentId);

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormPickers/ChatFormMentionPicker.svelte

@@ -8,11 +8,9 @@
 	import { BuiltInTool, FileMentionEntryType, GlobSearchType, KeyboardKey } from '$lib/enums';
 	import { useDebouncedSearch } from '$lib/hooks/use-debounced-search.svelte';
 	import { usePickerNavigation } from '$lib/hooks/use-picker-navigation.svelte';
-	import { config } from '$lib/stores/settings.svelte';
-	import { toolsStore } from '$lib/stores/tools.svelte';
-	import { isMobile } from '$lib/stores/viewport.svelte';
-	import type { FileMentionEntry } from '$lib/types';
-	import { abbreviateHome, type GlobEntryResult, runGlobSearchWithChildren } from '$lib/utils';
+	import { isMobile, settingsStore, toolsStore } from '$lib/stores';
+	import type { FileMentionEntry, GlobEntryResult } from '$lib/types';
+	import { abbreviateHome, runGlobSearchWithChildren } from '$lib/utils';
 
 	/**
 	 * Floating file/folder mention picker. The chat input is the search
@@ -64,7 +62,7 @@
 	// Coerce the depth setting to a positive integer; an invalid value
 	// would otherwise reach the server as max_depth 0 = unlimited.
 	const searchDepth = $derived.by(() => {
-		const n = Number(config().mentionSearchMaxDepth);
+		const n = Number(settingsStore.config.mentionSearchMaxDepth);
 
 		return Number.isInteger(n) && n > 0 ? n : FILE_GLOB_SEARCH_PICKERS.DEFAULT_SEARCH_DEPTH;
 	});

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormPickers/ChatFormPicker/ChatFormPickerItemHeader.svelte

@@ -1,5 +1,5 @@
 <script lang="ts">
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { mcpStore } from '$lib/stores';
 	import type { MCPServerSettingsEntry } from '$lib/types';
 	import type { Snippet } from 'svelte';
 

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormPickers/ChatFormPickerMcpPrompts/ChatFormPickerMcpPrompts.svelte

@@ -9,8 +9,7 @@
 	} from '$lib/components/app/chat';
 	import Badge from '$lib/components/ui/badge/badge.svelte';
 	import { KeyboardKey } from '$lib/enums';
-	import { conversationsStore } from '$lib/stores/conversations.svelte';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { conversationsStore, mcpStore } from '$lib/stores';
 	import type { GetPromptResult, MCPPromptInfo, MCPServerSettingsEntry } from '$lib/types';
 	import { debounce, uuid } from '$lib/utils';
 	import { SvelteMap } from 'svelte/reactivity';

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

@@ -1,5 +1,5 @@
 <script lang="ts">
-	import { isMobile } from '$lib/stores/viewport.svelte';
+	import { isMobile } from '$lib/stores';
 	import { autoResizeTextarea } from '$lib/utils';
 	import { onMount } from 'svelte';
 

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

@@ -10,11 +10,11 @@
 	import { usePickerNavigation } from '$lib/hooks/use-picker-navigation.svelte';
 	import { useScrollActiveRow } from '$lib/hooks/use-scroll-active-row.svelte';
 	import { ToolsService } from '$lib/services/tools.service';
-	import { toolsStore } from '$lib/stores/tools.svelte';
+	import { toolsStore } from '$lib/stores';
+	import type { GlobEntry } from '$lib/types';
 	import {
 		abbreviateHome,
 		buildCaseInsensitiveGlob,
-		type GlobEntry,
 		joinPath,
 		lastPathSegment,
 		runGlobSearchWithChildren

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

@@ -8,18 +8,21 @@
 		ChatMessageUser
 	} from '$lib/components/app/chat';
 	import { REASONING_TAGS, ROUTES, SYSTEM_MESSAGE_PLACEHOLDER } from '$lib/constants';
-	import { getChatActionsContext, setMessageEditContext } from '$lib/contexts';
+	import { setChatMessageActionsContext, setChatMessageEditContext } from '$lib/contexts';
 	import { AgenticSectionType, AttachmentType, MessageRole } from '$lib/enums';
 	import { DatabaseService } from '$lib/services/database.service';
-	import { chatStore, pendingEditMessageId } from '$lib/stores/chat.svelte';
-	import { conversationsStore } from '$lib/stores/conversations.svelte';
-	import { isMobile } from '$lib/stores/viewport.svelte';
-	import type { DatabaseMessageExtraMcpPrompt } from '$lib/types';
+	import { chatStore, conversationsStore, isMobile } from '$lib/stores';
+	import type {
+		ChatMessageActions,
+		ChatMessageDeletionInfo,
+		DatabaseMessageExtraMcpPrompt
+	} from '$lib/types';
 	import { deriveAgenticSections } from '$lib/utils';
 	import { parseFilesToMessageExtras } from '$lib/utils/browser-only';
 
 	interface Props {
 		class?: string;
+		chatActions: ChatMessageActions;
 		message: DatabaseMessage;
 		toolMessages?: DatabaseMessage[];
 		isLastAssistantMessage?: boolean;
@@ -29,6 +32,7 @@
 	}
 
 	let {
+		chatActions,
 		class: className = '',
 		isLastAssistantMessage = false,
 		isLastUserMessage = false,
@@ -38,14 +42,7 @@
 		toolMessages = []
 	}: Props = $props();
 
-	const chatActions = getChatActionsContext();
-
-	let deletionInfo = $state<{
-		totalCount: number;
-		userMessages: number;
-		assistantMessages: number;
-		messageTypes: string[];
-	} | null>(null);
+	let deletionInfo = $state<ChatMessageDeletionInfo | null>(null);
 	// The system message placeholder must never surface as editable content; keeping
 	// it in the derived (not just in handleEdit) guards against prop invalidation
 	// reverting the override while editing
@@ -114,7 +111,7 @@
 	let showSaveOnlyOption = $derived(message.role === MessageRole.USER);
 	let showBranchAfterEditOption = $derived(message.role === MessageRole.ASSISTANT);
 
-	setMessageEditContext({
+	setChatMessageEditContext({
 		cancel: handleCancelEdit,
 		get editedContent() {
 			return editedContent;
@@ -168,6 +165,30 @@
 		startEdit: handleEdit
 	});
 
+	setChatMessageActionsContext({
+		confirmDelete: handleConfirmDelete,
+		copy: handleCopy,
+		get deletionInfo() {
+			return deletionInfo;
+		},
+		get forkConversation() {
+			const isForkableUser = message.role === MessageRole.USER && !mcpPromptExtra;
+
+			return isForkableUser || message.role === MessageRole.ASSISTANT
+				? handleForkConversation
+				: undefined;
+		},
+		navigateToSibling: handleNavigateToSibling,
+		requestDelete: handleDelete,
+		setShowDeleteDialog: handleShowDeleteDialogChange,
+		get showDeleteDialog() {
+			return showDeleteDialog;
+		},
+		get siblingInfo() {
+			return siblingInfo;
+		}
+	});
+
 	let mcpPromptExtra = $derived.by(() => {
 		if (message.role !== MessageRole.USER) return null;
 
@@ -185,7 +206,7 @@
 	});
 
 	$effect(() => {
-		const pendingId = pendingEditMessageId();
+		const pendingId = chatStore.pendingEditMessageId;
 
 		if (pendingId && pendingId === message.id && !isEditing) {
 			handleEdit();
@@ -362,73 +383,22 @@
 
 <div class="chat-message" class:chat-message--synthetic={isSynthetic}>
 	{#if message.role === MessageRole.SYSTEM}
-		<ChatMessageSystem
-			bind:textareaElement
-			class={className}
-			{deletionInfo}
-			{message}
-			onConfirmDelete={handleConfirmDelete}
-			onCopy={handleCopy}
-			onDelete={handleDelete}
-			onEdit={handleEdit}
-			onNavigateToSibling={handleNavigateToSibling}
-			onShowDeleteDialogChange={handleShowDeleteDialogChange}
-			{showDeleteDialog}
-			{siblingInfo}
-		/>
+		<ChatMessageSystem bind:textareaElement class={className} {message} />
 	{:else if mcpPromptExtra}
-		<ChatMessageMcpPrompt
-			class={className}
-			{deletionInfo}
-			{message}
-			mcpPrompt={mcpPromptExtra}
-			onConfirmDelete={handleConfirmDelete}
-			onCopy={handleCopy}
-			onDelete={handleDelete}
-			onEdit={handleEdit}
-			onNavigateToSibling={handleNavigateToSibling}
-			onShowDeleteDialogChange={handleShowDeleteDialogChange}
-			{showDeleteDialog}
-			{siblingInfo}
-		/>
+		<ChatMessageMcpPrompt class={className} {message} mcpPrompt={mcpPromptExtra} />
 	{:else if isSynthetic}
 		<ChatMessageSynthetic {message} class={className} />
 	{:else if message.role === MessageRole.USER}
-		<ChatMessageUser
-			class={className}
-			{deletionInfo}
-			{isLastUserMessage}
-			{message}
-			{nextAssistantMessage}
-			onConfirmDelete={handleConfirmDelete}
-			onCopy={handleCopy}
-			onDelete={handleDelete}
-			onEdit={handleEdit}
-			onForkConversation={handleForkConversation}
-			onNavigateToSibling={handleNavigateToSibling}
-			onShowDeleteDialogChange={handleShowDeleteDialogChange}
-			{showDeleteDialog}
-			{siblingInfo}
-		/>
+		<ChatMessageUser class={className} {isLastUserMessage} {message} {nextAssistantMessage} />
 	{:else}
 		<ChatMessageAssistant
 			bind:textareaElement
 			class={className}
-			{deletionInfo}
 			{isLastAssistantMessage}
 			{message}
 			{toolMessages}
-			onConfirmDelete={handleConfirmDelete}
 			onContinue={handleContinue}
-			onCopy={handleCopy}
-			onDelete={handleDelete}
-			onEdit={handleEdit}
-			onForkConversation={handleForkConversation}
-			onNavigateToSibling={handleNavigateToSibling}
 			onRegenerate={handleRegenerate}
-			onShowDeleteDialogChange={handleShowDeleteDialogChange}
-			{showDeleteDialog}
-			{siblingInfo}
 		/>
 	{/if}
 </div>

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

@@ -8,76 +8,48 @@
 		ChatMessageAssistantStatistics,
 		ChatMessageEditForm
 	} from '$lib/components/app';
-	import { getMessageEditContext } from '$lib/contexts';
+	import { getChatMessageEditContext } from '$lib/contexts';
 	import { MessageRole } from '$lib/enums';
 	import { useProcessingState } from '$lib/hooks/use-processing-state.svelte';
-	import { chatStore, isChatStreaming, isLoading } from '$lib/stores/chat.svelte';
-	import { modelsStore } from '$lib/stores/models.svelte';
-	import { isRouterMode } from '$lib/stores/server.svelte';
-	import { config } from '$lib/stores/settings.svelte';
+	import { chatStore, modelsStore, serverStore, settingsStore } from '$lib/stores';
 	import { modelLoadProgressText } from '$lib/utils';
 	import { hasAgenticContent } from '$lib/utils';
 
 	interface Props {
 		class?: string;
-		deletionInfo: {
-			totalCount: number;
-			userMessages: number;
-			assistantMessages: number;
-			messageTypes: string[];
-		} | null;
 		isLastAssistantMessage?: boolean;
 		message: DatabaseMessage;
 		toolMessages?: DatabaseMessage[];
-		onCopy: () => void;
-		onConfirmDelete: () => void;
 		onContinue?: () => void;
-		onDelete: () => void;
-		onEdit?: () => void;
-		onForkConversation?: (options: { name: string; includeAttachments: boolean }) => void;
-		onNavigateToSibling?: (siblingId: string) => void;
 		onRegenerate: (modelOverride?: string) => void;
-		onShowDeleteDialogChange: (show: boolean) => void;
-		showDeleteDialog: boolean;
-		siblingInfo?: ChatMessageSiblingInfo | null;
 		textareaElement?: HTMLTextAreaElement;
 	}
 
 	let {
 		class: className = '',
-		deletionInfo,
 		isLastAssistantMessage = false,
 		message,
-		onConfirmDelete,
 		onContinue,
-		onCopy,
-		onDelete,
-		onEdit,
-		onForkConversation,
-		onNavigateToSibling,
 		onRegenerate,
-		onShowDeleteDialogChange,
-		showDeleteDialog,
-		siblingInfo = null,
 		textareaElement = $bindable(),
 		toolMessages = []
 	}: Props = $props();
 
 	// Get edit context
-	const editCtx = getMessageEditContext();
+	const editCtx = getChatMessageEditContext();
 
 	const isAgentic = $derived(hasAgenticContent(message, toolMessages));
 	const processingState = useProcessingState();
 
-	let currentConfig = $derived(config());
-	let isRouter = $derived(isRouterMode());
+	let currentConfig = $derived(settingsStore.config);
+	let isRouter = $derived(serverStore.isRouterMode);
 
 	let showRawOutput = $state(false);
 
 	let displayedModel = $derived(message.model ?? null);
 
-	let isCurrentlyLoading = $derived(isLoading());
-	let isStreaming = $derived(isChatStreaming());
+	let isCurrentlyLoading = $derived(chatStore.isLoading);
+	let isStreaming = $derived(chatStore.isStreaming());
 	let hasNoContent = $derived(!message?.content?.trim());
 	let isActivelyProcessing = $derived(isCurrentlyLoading || isStreaming);
 
@@ -175,7 +147,7 @@
 			<ChatMessageAgenticContent
 				{message}
 				{toolMessages}
-				isStreaming={isChatStreaming()}
+				isStreaming={chatStore.isStreaming()}
 				{isLastAssistantMessage}
 			/>
 		{/if}
@@ -190,14 +162,14 @@
 			<div class="inline-flex flex-wrap items-start gap-2 text-xs text-muted-foreground">
 				<ChatMessageAssistantModel
 					{displayedModel}
-					isLoading={isLoading()}
+					isLoading={chatStore.isLoading}
 					{isRouter}
 					{onRegenerate}
 				/>
 
 				<ChatMessageAssistantStatistics
 					{message}
-					isLoading={isLoading()}
+					isLoading={chatStore.isLoading}
 					{processingState}
 					showMessageStats={currentConfig.showMessageStats}
 				/>
@@ -210,18 +182,8 @@
 			role={MessageRole.ASSISTANT}
 			justify="start"
 			actionsPosition="left"
-			{siblingInfo}
-			{showDeleteDialog}
-			{deletionInfo}
-			{onCopy}
-			{onEdit}
 			{onRegenerate}
 			onContinue={currentConfig.enableContinueGeneration ? onContinue : undefined}
-			{onForkConversation}
-			{onDelete}
-			{onConfirmDelete}
-			{onNavigateToSibling}
-			{onShowDeleteDialogChange}
 			showRawOutputSwitch={currentConfig.showRawOutputSwitch}
 			rawOutputEnabled={showRawOutput}
 			onRawOutputToggle={(enabled) => (showRawOutput = enabled)}

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

@@ -1,7 +1,7 @@
 <script lang="ts">
 	import { ModelBadge, ModelsSelectorDropdown } from '$lib/components/app';
 	import { ServerModelStatus } from '$lib/enums';
-	import { modelsStore } from '$lib/stores/models.svelte';
+	import { modelsStore } from '$lib/stores';
 	import { copyToClipboard } from '$lib/utils';
 
 	interface Props {

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

@@ -10,7 +10,7 @@
 
 	let { modelLoadingText, position, processingState }: Props = $props();
 
-	const marginClass = position === 'top' ? 'mt-6' : 'mt-4';
+	const marginClass = $derived(position === 'top' ? 'mt-6' : 'mt-4');
 </script>
 
 <div class="{marginClass} w-full max-w-3xl" in:fade>

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

@@ -2,6 +2,7 @@
 	import { ChatMessageStatistics } from '$lib/components/app';
 	import { ChatMessageStatisticsMode } from '$lib/enums';
 	import type { UseProcessingStateReturn } from '$lib/hooks/use-processing-state.svelte';
+	import { agenticStore } from '$lib/stores';
 
 	interface Props {
 		message: DatabaseMessage;
@@ -11,9 +12,26 @@
 	}
 
 	let { isLoading, message, processingState, showMessageStats }: Props = $props();
+
+	// A running agentic flow stamps per-turn timings on its root message at each
+	// turn boundary and the cumulative agentic totals only on exit; while it runs,
+	// show the session's live totals on the root message instead.
+	const liveLlm = $derived(agenticStore.getLiveLlmTotals(message.convId));
+	const isLiveFlowRoot = $derived(
+		liveLlm !== null && agenticStore.getFlowRootMessageId(message.convId) === message.id
+	);
 </script>
 
-{#if showMessageStats && message.timings && message.timings.predicted_n && message.timings.predicted_ms}
+{#if showMessageStats && isLiveFlowRoot && liveLlm}
+	<ChatMessageStatistics
+		mode={ChatMessageStatisticsMode.GENERATION}
+		isLive
+		promptTokens={liveLlm.prompt_n}
+		promptMs={liveLlm.prompt_ms}
+		predictedTokens={liveLlm.predicted_n}
+		predictedMs={liveLlm.predicted_ms}
+	/>
+{:else if showMessageStats && message.timings && message.timings.predicted_n && message.timings.predicted_ms}
 	{@const agentic = message.timings.agentic}
 	<ChatMessageStatistics
 		mode={ChatMessageStatisticsMode.GENERATION}

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

@@ -4,7 +4,7 @@
 		ChatMessageEditForm,
 		ChatMessageMcpPromptContent
 	} from '$lib/components/app';
-	import { getMessageEditContext } from '$lib/contexts';
+	import { getChatMessageEditContext } from '$lib/contexts';
 	import { McpPromptVariant, MessageRole } from '$lib/enums';
 	import type { DatabaseMessageExtraMcpPrompt } from '$lib/types';
 
@@ -12,39 +12,12 @@
 		class?: string;
 		message: DatabaseMessage;
 		mcpPrompt: DatabaseMessageExtraMcpPrompt;
-		siblingInfo?: ChatMessageSiblingInfo | null;
-		showDeleteDialog: boolean;
-		deletionInfo: {
-			totalCount: number;
-			userMessages: number;
-			assistantMessages: number;
-			messageTypes: string[];
-		} | null;
-		onCopy: () => void;
-		onEdit: () => void;
-		onDelete: () => void;
-		onConfirmDelete: () => void;
-		onNavigateToSibling?: (siblingId: string) => void;
-		onShowDeleteDialogChange: (show: boolean) => void;
 	}
 
-	let {
-		class: className = '',
-		deletionInfo,
-		mcpPrompt,
-		message,
-		onConfirmDelete,
-		onCopy,
-		onDelete,
-		onEdit,
-		onNavigateToSibling,
-		onShowDeleteDialogChange,
-		showDeleteDialog,
-		siblingInfo = null
-	}: Props = $props();
+	let { class: className = '', mcpPrompt, message }: Props = $props();
 
 	// Get edit context
-	const editCtx = getMessageEditContext();
+	const editCtx = getChatMessageEditContext();
 </script>
 
 <div
@@ -63,20 +36,7 @@
 
 		{#if message.timestamp}
 			<div class="max-w-[80%]">
-				<ChatMessageActionIcons
-					actionsPosition="right"
-					{deletionInfo}
-					justify="end"
-					{onConfirmDelete}
-					{onCopy}
-					{onDelete}
-					{onEdit}
-					{onNavigateToSibling}
-					{onShowDeleteDialogChange}
-					{siblingInfo}
-					{showDeleteDialog}
-					role={MessageRole.USER}
-				/>
+				<ChatMessageActionIcons actionsPosition="right" justify="end" role={MessageRole.USER} />
 			</div>
 		{/if}
 	{/if}

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

@@ -3,7 +3,7 @@
 	import { Card } from '$lib/components/ui/card';
 	import * as Tooltip from '$lib/components/ui/tooltip';
 	import { McpPromptVariant } from '$lib/enums';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { mcpStore } from '$lib/stores';
 	import type { DatabaseMessageExtraMcpPrompt } from '$lib/types';
 	import { SvelteMap } from 'svelte/reactivity';
 

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

@@ -4,47 +4,20 @@
 	import { Button } from '$lib/components/ui/button';
 	import { Card } from '$lib/components/ui/card';
 	import { INPUT_CLASSES } from '$lib/constants';
-	import { getMessageEditContext } from '$lib/contexts';
+	import { getChatMessageEditContext } from '$lib/contexts';
 	import { KeyboardKey, MessageRole } from '$lib/enums';
-	import { config } from '$lib/stores/settings.svelte';
+	import { settingsStore } from '$lib/stores';
 	import { autoResizeTextarea, isIMEComposing } from '$lib/utils';
 
 	interface Props {
 		class?: string;
 		message: DatabaseMessage;
-		siblingInfo?: ChatMessageSiblingInfo | null;
-		showDeleteDialog: boolean;
-		deletionInfo: {
-			totalCount: number;
-			userMessages: number;
-			assistantMessages: number;
-			messageTypes: string[];
-		} | null;
-		onCopy: () => void;
-		onEdit: () => void;
-		onDelete: () => void;
-		onConfirmDelete: () => void;
-		onNavigateToSibling?: (siblingId: string) => void;
-		onShowDeleteDialogChange: (show: boolean) => void;
 		textareaElement?: HTMLTextAreaElement;
 	}
 
-	let {
-		class: className = '',
-		deletionInfo,
-		message,
-		onConfirmDelete,
-		onCopy,
-		onDelete,
-		onEdit,
-		onNavigateToSibling,
-		onShowDeleteDialogChange,
-		showDeleteDialog,
-		siblingInfo = null,
-		textareaElement = $bindable()
-	}: Props = $props();
+	let { class: className = '', message, textareaElement = $bindable() }: Props = $props();
 
-	const editCtx = getMessageEditContext();
+	const editCtx = getChatMessageEditContext();
 
 	function handleEditKeydown(event: KeyboardEvent) {
 		if (event.key === KeyboardKey.ENTER && !event.shiftKey && !isIMEComposing(event)) {
@@ -64,7 +37,7 @@
 	let contentHeight = $state(0);
 
 	const MAX_HEIGHT = 200; // pixels
-	const currentConfig = config();
+	const currentConfig = settingsStore.config;
 
 	let showExpandButton = $derived(contentHeight > MAX_HEIGHT);
 
@@ -218,20 +191,7 @@
 
 		{#if message.timestamp}
 			<div class="max-w-[80%]">
-				<ChatMessageActionIcons
-					actionsPosition="right"
-					{deletionInfo}
-					justify="end"
-					{onConfirmDelete}
-					{onCopy}
-					{onDelete}
-					{onEdit}
-					{onNavigateToSibling}
-					{onShowDeleteDialogChange}
-					{siblingInfo}
-					{showDeleteDialog}
-					role={MessageRole.USER}
-				/>
+				<ChatMessageActionIcons actionsPosition="right" justify="end" role={MessageRole.USER} />
 			</div>
 		{/if}
 	{/if}

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

@@ -12,13 +12,8 @@
 	import ChatMessageToolCallBlockSearchResults from './ChatMessageToolCallBlockSearchResults.svelte';
 	import ChatMessageToolCallBlockWriteFile from './ChatMessageToolCallBlockWriteFile.svelte';
 	import { BuiltInTool } from '$lib/enums';
-	import type { DatabaseMessageExtra } from '$lib/types';
-	import {
-		type AgenticSection,
-		extractSearchQuery,
-		extractSearchResults,
-		isWebSearchToolName
-	} from '$lib/utils';
+	import type { AgenticSection, DatabaseMessageExtra } from '$lib/types';
+	import { extractSearchQuery, extractSearchResults, isWebSearchToolName } from '$lib/utils';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -8,14 +8,12 @@
 	import { MarkdownContent, SyntaxHighlightedCode } from '$lib/components/app';
 	import { MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
 	import { AttachmentType, FileTypeText, MimeTypeAudio, ToolResultKind } from '$lib/enums';
-	import type { DatabaseMessageExtra } from '$lib/types';
+	import type { AgenticSection, DatabaseMessageExtra, ToolResultLine } from '$lib/types';
 	import {
-		type AgenticSection,
 		classifyToolResult,
 		formatJsonPretty,
 		getBuiltinToolUi,
-		parseToolResultWithMedia,
-		type ToolResultLine
+		parseToolResultWithMedia
 	} from '$lib/utils';
 	import { createBase64DataUrl } from '$lib/utils/data-url';
 

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

@@ -3,8 +3,9 @@
 	import ToolCallBlock from './ToolCallBlock.svelte';
 	import { XCircle } from '@lucide/svelte';
 	import { MAX_HEIGHT_CODE_BLOCK, RESULT_STAT_SEPARATOR } from '$lib/constants';
-	import { toolsStore } from '$lib/stores/tools.svelte';
-	import { abbreviateHome, type AgenticSection, computeLineDiff, prefixFor } from '$lib/utils';
+	import { toolsStore } from '$lib/stores';
+	import type { AgenticSection } from '$lib/types';
+	import { abbreviateHome, computeLineDiff, prefixFor } from '$lib/utils';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -11,19 +11,17 @@
 	import { AlertTriangle, Check, Loader2, XCircle } from '@lucide/svelte';
 	import { CollapsibleTerminalBlock } from '$lib/components/app';
 	import { SETTINGS_KEYS, TOOL_RUNTIME_SCROLL_AT_BOTTOM_THRESHOLD_PX } from '$lib/constants';
-	import { config } from '$lib/stores/settings.svelte';
-	import { toolsStore } from '$lib/stores/tools.svelte';
-	import type { DatabaseMessageExtra } from '$lib/types';
+	import { AttachmentType } from '$lib/enums';
+	import { settingsStore, toolsStore } from '$lib/stores';
+	import type { AgenticSection, DatabaseMessageExtra, ToolResultLine } from '$lib/types';
 	import {
 		abbreviateHome,
-		type AgenticSection,
 		type ExecShellExitStatus,
 		highlightCode,
 		isExitCodeSummaryLine,
 		parseExecShellCommandError,
 		parseExecShellCommandExitStatus,
-		parseToolResultWithMedia,
-		type ToolResultLine
+		parseToolResultWithMedia
 	} from '$lib/utils';
 
 	interface Props {
@@ -93,7 +91,7 @@
 	);
 
 	const useFullHeightCodeBlocks = $derived(
-		Boolean(config()[SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS])
+		Boolean(settingsStore.config[SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS])
 	);
 
 	const autoScroll = $derived(isLive && !useFullHeightCodeBlocks);
@@ -222,7 +220,7 @@
 			>
 				{#each outputLines as line, i (i)}
 					<div class="font-mono text-[11px] leading-relaxed whitespace-pre-wrap">{line.text}</div>
-					{#if line.media}
+					{#if line.media?.type === AttachmentType.IMAGE}
 						<img
 							src={line.media.base64Url}
 							alt={line.media.name}

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

@@ -2,8 +2,9 @@
 	import { parseFileGlobSearchMeta } from './parsers/file-glob-search';
 	import ToolCallBlock from './ToolCallBlock.svelte';
 	import { XCircle } from '@lucide/svelte';
-	import { toolsStore } from '$lib/stores/tools.svelte';
-	import { abbreviateHome, type AgenticSection } from '$lib/utils';
+	import { toolsStore } from '$lib/stores';
+	import type { AgenticSection } from '$lib/types';
+	import { abbreviateHome } from '$lib/utils';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -1,7 +1,7 @@
 <script lang="ts">
 	import { Clock, Loader2 } from '@lucide/svelte';
 	import { AgenticSectionType } from '$lib/enums';
-	import type { AgenticSection } from '$lib/utils';
+	import type { AgenticSection } from '$lib/types';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -1,8 +1,9 @@
 <script lang="ts">
 	import { Info, Loader2 } from '@lucide/svelte';
 	import { AgenticSectionType } from '$lib/enums';
-	import { toolsStore } from '$lib/stores/tools.svelte';
-	import { abbreviateHome, type AgenticSection } from '$lib/utils';
+	import { toolsStore } from '$lib/stores';
+	import type { AgenticSection } from '$lib/types';
+	import { abbreviateHome } from '$lib/utils';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -2,8 +2,9 @@
 	import { parseGrepSearchMeta } from './parsers/grep-search';
 	import ToolCallBlock from './ToolCallBlock.svelte';
 	import { XCircle } from '@lucide/svelte';
-	import { toolsStore } from '$lib/stores/tools.svelte';
-	import { abbreviateHome, type AgenticSection } from '$lib/utils';
+	import { toolsStore } from '$lib/stores';
+	import type { AgenticSection } from '$lib/types';
+	import { abbreviateHome } from '$lib/utils';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -3,7 +3,7 @@
 	import ToolCallBlock from './ToolCallBlock.svelte';
 	import { SyntaxHighlightedCode } from '$lib/components/app';
 	import { CODE_BLOCK, MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
-	import { type AgenticSection } from '$lib/utils';
+	import type { AgenticSection } from '$lib/types';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -4,7 +4,7 @@
 	import { ATTACHMENT_SAVED_REGEX } from '$lib/constants/agentic.constants';
 	import { AttachmentType, MimeTypeAudio } from '$lib/enums';
 	import type { DatabaseMessageExtraAudioFile, DatabaseMessageExtraImageFile } from '$lib/types';
-	import { type AgenticSection } from '$lib/utils';
+	import type { AgenticSection } from '$lib/types';
 	import { createBase64DataUrl } from '$lib/utils/data-url';
 
 	interface Props {

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

@@ -5,7 +5,8 @@
 	import { SyntaxHighlightedCode } from '$lib/components/app';
 	import { MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
 	import { FileTypeText } from '$lib/enums';
-	import { type AgenticSection, getBuiltinToolUi } from '$lib/utils';
+	import type { AgenticSection } from '$lib/types';
+	import { getBuiltinToolUi } from '$lib/utils';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -4,14 +4,13 @@
 	import * as HoverCard from '$lib/components/ui/hover-card';
 	import { ICON_CLASS_DEFAULT, ICON_CLASS_SPIN } from '$lib/constants';
 	import { AgenticSectionType } from '$lib/enums';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { mcpStore } from '$lib/stores';
+	import type { AgenticSection, SearchResult } from '$lib/types';
 	import {
-		type AgenticSection,
 		extractSearchQuery,
 		extractSearchResults,
 		faviconForUrl,
-		sanitizeExternalUrl,
-		type SearchResult
+		sanitizeExternalUrl
 	} from '$lib/utils';
 
 	interface Props {

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

@@ -4,8 +4,9 @@
 	import { XCircle } from '@lucide/svelte';
 	import { SyntaxHighlightedCode } from '$lib/components/app';
 	import { MAX_HEIGHT_CODE_BLOCK, RESULT_STAT_SEPARATOR } from '$lib/constants';
-	import { toolsStore } from '$lib/stores/tools.svelte';
-	import { abbreviateHome, type AgenticSection } from '$lib/utils';
+	import { toolsStore } from '$lib/stores';
+	import type { AgenticSection } from '$lib/types';
+	import { abbreviateHome } from '$lib/utils';
 
 	interface Props {
 		section: AgenticSection;

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

@@ -13,8 +13,9 @@
 	import { CollapsibleContentBlock } from '$lib/components/app';
 	import { ICON_CLASS_DEFAULT, ICON_CLASS_SPIN } from '$lib/constants';
 	import { AgenticSectionType } from '$lib/enums';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
-	import { type AgenticSection, type BuiltinToolUiEntry, getBuiltinToolUi } from '$lib/utils';
+	import { mcpStore } from '$lib/stores';
+	import type { AgenticSection, BuiltinToolUiEntry } from '$lib/types';
+	import { getBuiltinToolUi } from '$lib/utils';
 	import type { Component, Snippet } from 'svelte';
 
 	type ToolCallBlockMetaWithError = TMeta & { errorMessage?: string };

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts

@@ -5,7 +5,7 @@
 // stay focused on its own format quirks.
 
 import { BuiltInTool } from '$lib/enums';
-import type { AgenticSection } from '$lib/utils/agentic';
+import type { AgenticSection } from '$lib/types/agentic';
 import { parsePartialJsonArgs } from '$lib/utils/parse-partial-json-args';
 
 /**

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts

@@ -6,7 +6,8 @@
 import { parseToolArgs } from './_shared';
 import { FILE_PATH_SEPARATOR_REGEX } from '$lib/constants';
 import { BuiltInTool } from '$lib/enums';
-import { type AgenticSection, tryParseToolResultObject } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
+import { tryParseToolResultObject } from '$lib/utils';
 
 export type EditFileEdit = {
 	oldText: string;

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts

@@ -7,7 +7,7 @@
 
 import { parseToolArgs } from './_shared';
 import { BuiltInTool } from '$lib/enums';
-import type { AgenticSection } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
 
 export type ExecShellCommandMeta = {
 	command: string;

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts

@@ -6,7 +6,8 @@
 
 import { parseToolArgs } from './_shared';
 import { BuiltInTool } from '$lib/enums';
-import { type AgenticSection, splitSearchSummaryList } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
+import { splitSearchSummaryList } from '$lib/utils';
 
 export type FileGlobSearchMeta = {
 	path: string;

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts

@@ -7,7 +7,8 @@
 
 import { parseToolArgs } from './_shared';
 import { BuiltInTool } from '$lib/enums';
-import { type AgenticSection, splitSearchSummaryList } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
+import { splitSearchSummaryList } from '$lib/utils';
 
 export type GrepSearchMatch = {
 	file: string;

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts

@@ -6,7 +6,8 @@
 import { parseToolArgs } from './_shared';
 import { CODE_BLOCK, FILE_PATH_SEPARATOR_REGEX } from '$lib/constants';
 import { BuiltInTool } from '$lib/enums';
-import { type AgenticSection, getFileTypeByExtension } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
+import { getFileTypeByExtension } from '$lib/utils';
 
 export type ReadFileMeta = {
 	fileName: string;

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-media.ts

@@ -5,7 +5,7 @@ import {
 	PREFIX_SIZE,
 	READ_MEDIA_SIZE_REGEX
 } from '$lib/constants/read-media';
-import type { AgenticSection } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
 
 export interface ReadMediaMeta {
 	fileName: string;

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts

@@ -7,7 +7,7 @@
 
 import { parseToolArgs } from './_shared';
 import { BuiltInTool } from '$lib/enums';
-import type { AgenticSection } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
 
 export type RunJavascriptMeta = {
 	code: string;

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts

@@ -6,7 +6,8 @@
 import { parseToolArgs } from './_shared';
 import { CODE_BLOCK, FILE_PATH_SEPARATOR_REGEX } from '$lib/constants';
 import { BuiltInTool } from '$lib/enums';
-import { type AgenticSection, getFileTypeByExtension, tryParseToolResultObject } from '$lib/utils';
+import type { AgenticSection } from '$lib/types';
+import { getFileTypeByExtension, tryParseToolResultObject } from '$lib/utils';
 
 export type WriteFileMeta = {
 	fileName: string;

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

@@ -5,57 +5,31 @@
 		ChatMessageStatistics,
 		ChatMessageUserBubble
 	} from '$lib/components/app/chat';
-	import { getMessageEditContext } from '$lib/contexts';
+	import { getChatMessageEditContext } from '$lib/contexts';
 	import { ChatMessageStatisticsMode, MessageRole } from '$lib/enums';
 	import { useProcessingState } from '$lib/hooks/use-processing-state.svelte';
-	import { isLoading } from '$lib/stores/chat.svelte';
-	import { config } from '$lib/stores/settings.svelte';
+	import { chatStore, settingsStore } from '$lib/stores';
 
 	interface Props {
 		class?: string;
 		message: DatabaseMessage;
-		siblingInfo?: ChatMessageSiblingInfo | null;
-		deletionInfo: {
-			totalCount: number;
-			userMessages: number;
-			assistantMessages: number;
-			messageTypes: string[];
-		} | null;
 		isLastUserMessage?: boolean;
 		nextAssistantMessage?: DatabaseMessage | null;
-		showDeleteDialog: boolean;
-		onEdit: () => void;
-		onDelete: () => void;
-		onConfirmDelete: () => void;
-		onForkConversation?: (options: { name: string; includeAttachments: boolean }) => void;
-		onShowDeleteDialogChange: (show: boolean) => void;
-		onNavigateToSibling?: (siblingId: string) => void;
-		onCopy: () => void;
 	}
 
 	let {
 		class: className = '',
-		deletionInfo,
 		isLastUserMessage = false,
 		message,
-		nextAssistantMessage = null,
-		onConfirmDelete,
-		onCopy,
-		onDelete,
-		onEdit,
-		onForkConversation,
-		onNavigateToSibling,
-		onShowDeleteDialogChange,
-		showDeleteDialog,
-		siblingInfo = null
+		nextAssistantMessage = null
 	}: Props = $props();
 
 	// Get contexts
-	const editCtx = getMessageEditContext();
+	const editCtx = getChatMessageEditContext();
 	const processingState = useProcessingState();
 
-	const currentConfig = $derived(config());
-	const isActivelyProcessing = $derived(isLastUserMessage && isLoading());
+	const currentConfig = $derived(settingsStore.config);
+	const isActivelyProcessing = $derived(isLastUserMessage && chatStore.isLoading);
 
 	// For agentic turns, prefer the cumulative agentic.llm totals over per-call timings.
 	let storedReadingStats = $derived.by(() => {
@@ -133,21 +107,7 @@
 
 		{#if message.timestamp}
 			<div class="max-w-[80%]">
-				<ChatMessageActionIcons
-					actionsPosition="right"
-					{deletionInfo}
-					justify="end"
-					{onConfirmDelete}
-					{onCopy}
-					{onDelete}
-					{onEdit}
-					{onForkConversation}
-					{onNavigateToSibling}
-					{onShowDeleteDialogChange}
-					{siblingInfo}
-					{showDeleteDialog}
-					role={MessageRole.USER}
-				/>
+				<ChatMessageActionIcons actionsPosition="right" justify="end" role={MessageRole.USER} />
 			</div>
 		{/if}
 	{/if}

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

@@ -1,7 +1,7 @@
 <script lang="ts">
 	import { ChatAttachmentsList, MarkdownContent, MentionText } from '$lib/components/app';
 	import { Card } from '$lib/components/ui/card';
-	import { config } from '$lib/stores/settings.svelte';
+	import { settingsStore } from '$lib/stores';
 	import type { DatabaseMessageExtra } from '$lib/types/database';
 
 	interface Props {
@@ -24,7 +24,7 @@
 
 	let isMultiline = $state(false);
 	let messageElement: HTMLElement | undefined = $state();
-	const currentConfig = config();
+	const currentConfig = settingsStore.config;
 
 	$effect(() => {
 		if (!messageElement || !content.trim()) return;

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

@@ -1,7 +1,7 @@
 <script lang="ts">
 	import { ArrowUp, Edit, Trash2 } from '@lucide/svelte';
 	import { ActionIcon, ChatMessageEditForm, ChatMessageUserBubble } from '$lib/components/app';
-	import { useMessageEditContext } from '$lib/hooks/use-message-edit-context.svelte';
+	import { useChatMessageEditContext } from '$lib/hooks/use-chat-message-edit-context.svelte';
 
 	interface Props {
 		class?: string;
@@ -21,7 +21,7 @@
 		onSendImmediately
 	}: Props = $props();
 
-	const editCtx = useMessageEditContext({
+	const editCtx = useChatMessageEditContext({
 		getContent: () => content,
 		getExtras: () => extras,
 		onSave: (content, extras) => onEdit(content, extras)

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCardPermissionRequest.svelte

@@ -7,7 +7,7 @@
 	import { cn } from '$lib/components/ui/utils';
 	import { TOOL_SERVER_LABELS } from '$lib/constants';
 	import { ToolPermissionDecision, ToolSource } from '$lib/enums';
-	import { toolsStore } from '$lib/stores/tools.svelte';
+	import { toolsStore } from '$lib/stores';
 
 	interface Props {
 		toolName: string;

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

@@ -9,30 +9,16 @@
 	import Input from '$lib/components/ui/input/input.svelte';
 	import Label from '$lib/components/ui/label/label.svelte';
 	import { Switch } from '$lib/components/ui/switch';
+	import { getChatMessageActionsContext, getChatMessageEditContext } from '$lib/contexts';
 	import { MessageRole } from '$lib/enums';
-	import { activeConversation } from '$lib/stores/conversations.svelte';
+	import { conversationsStore } from '$lib/stores';
 
 	interface Props {
 		role: MessageRole.USER | MessageRole.ASSISTANT;
 		justify: 'start' | 'end';
 		actionsPosition: 'left' | 'right';
-		siblingInfo?: ChatMessageSiblingInfo | null;
-		showDeleteDialog: boolean;
-		deletionInfo: {
-			totalCount: number;
-			userMessages: number;
-			assistantMessages: number;
-			messageTypes: string[];
-		} | null;
-		onCopy: () => void;
-		onEdit?: () => void;
 		onRegenerate?: () => void;
 		onContinue?: () => void;
-		onForkConversation?: (options: { name: string; includeAttachments: boolean }) => void;
-		onDelete: () => void;
-		onConfirmDelete: () => void;
-		onNavigateToSibling?: (siblingId: string) => void;
-		onShowDeleteDialogChange: (show: boolean) => void;
 		showRawOutputSwitch?: boolean;
 		rawOutputEnabled?: boolean;
 		onRawOutputToggle?: (enabled: boolean) => void;
@@ -40,36 +26,29 @@
 
 	let {
 		actionsPosition,
-		deletionInfo,
 		justify,
-		onConfirmDelete,
 		onContinue,
-		onCopy,
-		onDelete,
-		onEdit,
-		onForkConversation,
-		onNavigateToSibling,
 		onRawOutputToggle,
 		onRegenerate,
-		onShowDeleteDialogChange,
 		rawOutputEnabled = false,
 		role,
-		showDeleteDialog,
-		showRawOutputSwitch = false,
-		siblingInfo = null
+		showRawOutputSwitch = false
 	}: Props = $props();
 
+	const messageActions = getChatMessageActionsContext();
+	const editCtx = getChatMessageEditContext();
+
 	let showForkDialog = $state(false);
 	let forkName = $state('');
 	let forkIncludeAttachments = $state(true);
 
 	function handleConfirmDelete() {
-		onConfirmDelete();
-		onShowDeleteDialogChange(false);
+		messageActions.confirmDelete();
+		messageActions.setShowDeleteDialog(false);
 	}
 
 	function handleOpenForkDialog() {
-		const conv = activeConversation();
+		const conv = conversationsStore.activeConversation;
 
 		forkName = `Fork of ${conv?.name ?? 'Conversation'}`;
 		forkIncludeAttachments = true;
@@ -77,7 +56,10 @@
 	}
 
 	function handleConfirmFork() {
-		onForkConversation?.({ includeAttachments: forkIncludeAttachments, name: forkName.trim() });
+		messageActions.forkConversation?.({
+			includeAttachments: forkIncludeAttachments,
+			name: forkName.trim()
+		});
 		showForkDialog = false;
 	}
 </script>
@@ -88,18 +70,16 @@
 			? 'left-0'
 			: 'right-0'} flex items-center gap-2 opacity-100 transition-opacity"
 	>
-		{#if siblingInfo && siblingInfo.totalSiblings > 1}
-			<ChatMessageActionIconsBranchingControls {siblingInfo} {onNavigateToSibling} />
+		{#if messageActions.siblingInfo && messageActions.siblingInfo.totalSiblings > 1}
+			<ChatMessageActionIconsBranchingControls />
 		{/if}
 
 		<div
 			class="pointer-events-auto inset-0 flex items-center gap-1 opacity-100 transition-all duration-150"
 		>
-			<ActionIcon icon={Copy} tooltip="Copy" onclick={onCopy} />
+			<ActionIcon icon={Copy} tooltip="Copy" onclick={messageActions.copy} />
 
-			{#if onEdit}
-				<ActionIcon icon={Edit} tooltip="Edit" onclick={onEdit} />
-			{/if}
+			<ActionIcon icon={Edit} tooltip="Edit" onclick={editCtx.startEdit} />
 
 			{#if role === MessageRole.ASSISTANT && onRegenerate}
 				<ActionIcon icon={RefreshCw} tooltip="Regenerate" onclick={() => onRegenerate()} />
@@ -109,11 +89,11 @@
 				<ActionIcon icon={ArrowRight} tooltip="Continue" onclick={onContinue} />
 			{/if}
 
-			{#if onForkConversation}
+			{#if messageActions.forkConversation}
 				<ActionIcon icon={GitBranch} tooltip="Fork conversation" onclick={handleOpenForkDialog} />
 			{/if}
 
-			<ActionIcon icon={Trash2} tooltip="Delete" onclick={onDelete} />
+			<ActionIcon icon={Trash2} tooltip="Delete" onclick={messageActions.requestDelete} />
 		</div>
 	</div>
 
@@ -129,19 +109,19 @@
 </div>
 
 <DialogConfirmation
-	bind:open={showDeleteDialog}
+	open={messageActions.showDeleteDialog}
 	title="Delete Message"
-	description={deletionInfo && deletionInfo.totalCount > 1
-		? `This will delete ${deletionInfo.totalCount} messages including: ${deletionInfo.userMessages} user message${deletionInfo.userMessages > 1 ? 's' : ''} and ${deletionInfo.assistantMessages} assistant response${deletionInfo.assistantMessages > 1 ? 's' : ''}. All messages in this branch and their responses will be permanently removed. This action cannot be undone.`
+	description={messageActions.deletionInfo && messageActions.deletionInfo.totalCount > 1
+		? `This will delete ${messageActions.deletionInfo.totalCount} messages including: ${messageActions.deletionInfo.userMessages} user message${messageActions.deletionInfo.userMessages > 1 ? 's' : ''} and ${messageActions.deletionInfo.assistantMessages} assistant response${messageActions.deletionInfo.assistantMessages > 1 ? 's' : ''}. All messages in this branch and their responses will be permanently removed. This action cannot be undone.`
 		: 'Are you sure you want to delete this message? This action cannot be undone.'}
-	confirmText={deletionInfo && deletionInfo.totalCount > 1
-		? `Delete ${deletionInfo.totalCount} Messages`
+	confirmText={messageActions.deletionInfo && messageActions.deletionInfo.totalCount > 1
+		? `Delete ${messageActions.deletionInfo.totalCount} Messages`
 		: 'Delete'}
 	cancelText="Cancel"
 	variant="destructive"
 	icon={Trash2}
 	onConfirm={handleConfirmDelete}
-	onCancel={() => onShowDeleteDialogChange(false)}
+	onCancel={() => messageActions.setShowDeleteDialog(false)}
 />
 
 <DialogConfirmation

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionIcons/ChatMessageActionIconsBranchingControls.svelte

@@ -1,14 +1,17 @@
 <script lang="ts">
 	import { ChevronLeft, ChevronRight } from '@lucide/svelte';
 	import { ActionIcon } from '$lib/components/app';
+	import { getChatMessageActionsContext } from '$lib/contexts';
 
 	interface Props {
 		class?: string;
-		siblingInfo: ChatMessageSiblingInfo | null;
-		onNavigateToSibling?: (siblingId: string) => void;
 	}
 
-	let { class: className = '', onNavigateToSibling, siblingInfo }: Props = $props();
+	let { class: className = '' }: Props = $props();
+
+	const messageActions = getChatMessageActionsContext();
+
+	let siblingInfo = $derived(messageActions.siblingInfo);
 
 	let hasPrevious = $derived(siblingInfo && siblingInfo.currentIndex > 0);
 	let hasNext = $derived(siblingInfo && siblingInfo.currentIndex < siblingInfo.totalSiblings - 1);
@@ -31,7 +34,7 @@
 			tooltip="Previous version"
 			disabled={!hasPrevious}
 			class="h-5 w-5 p-0 {!hasPrevious ? '!cursor-not-allowed opacity-30' : ''}"
-			onclick={() => onNavigateToSibling?.(previousSiblingId!)}
+			onclick={() => messageActions.navigateToSibling(previousSiblingId!)}
 		/>
 
 		<span class="px-1 font-mono text-xs">
@@ -43,7 +46,7 @@
 			tooltip="Next version"
 			disabled={!hasNext}
 			class="h-5 w-5 p-0 {!hasNext ? 'opacity-30' : ''}"
-			onclick={() => onNavigateToSibling?.(nextSiblingId!)}
+			onclick={() => messageActions.navigateToSibling(nextSiblingId!)}
 		/>
 	</div>
 {/if}

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

@@ -8,21 +8,14 @@
 		MarkdownContent
 	} from '$lib/components/app';
 	import { AgenticSectionType, ChatMessageStatsView, ToolPermissionDecision } from '$lib/enums';
-	import {
-		agenticExecutingToolCallId,
-		agenticLastError,
-		agenticPendingContinueRequest,
-		agenticPendingPermissionRequest,
-		agenticResolveContinue,
-		agenticResolvePermission
-	} from '$lib/stores/agentic.svelte';
-	import { config } from '$lib/stores/settings.svelte';
+	import { agenticStore, settingsStore } from '$lib/stores';
 	import type {
+		AgenticSection,
 		ChatMessageAgenticTimings,
 		ChatMessageAgenticTurnStats,
 		DatabaseMessage
 	} from '$lib/types';
-	import { type AgenticSection, deriveAgenticSections } from '$lib/utils';
+	import { deriveAgenticSections } from '$lib/utils';
 
 	interface Props {
 		message: DatabaseMessage;
@@ -40,19 +33,25 @@
 
 	let expandedStates: Record<number, boolean> = $state({});
 
-	const showThoughtInProgress = $derived(Boolean(config().showThoughtInProgress));
-	const alwaysShowToolCallContent = $derived(Boolean(config().alwaysShowToolCallContent));
-	const showMessageStats = $derived(Boolean(config().showMessageStats));
-	const showAgenticTurnStats = $derived(showMessageStats && Boolean(config().showAgenticTurnStats));
+	const showThoughtInProgress = $derived(Boolean(settingsStore.config.showThoughtInProgress));
+	const alwaysShowToolCallContent = $derived(
+		Boolean(settingsStore.config.alwaysShowToolCallContent)
+	);
+	const showMessageStats = $derived(Boolean(settingsStore.config.showMessageStats));
+	const showAgenticTurnStats = $derived(
+		showMessageStats && Boolean(settingsStore.config.showAgenticTurnStats)
+	);
 
 	const hasReasoningError = $derived(
-		isLastAssistantMessage ? !!agenticLastError(message.convId) : false
+		isLastAssistantMessage ? !!agenticStore.lastError(message.convId) : false
 	);
 
 	let permissionDismissed = $state(false);
 
 	const pendingPermission = $derived(
-		isStreaming && isLastAssistantMessage ? agenticPendingPermissionRequest(message.convId) : null
+		isStreaming && isLastAssistantMessage
+			? agenticStore.pendingPermissionRequest(message.convId)
+			: null
 	);
 
 	let prevPendingRef: typeof pendingPermission = null;
@@ -68,13 +67,15 @@
 
 	function handlePermission(decision: ToolPermissionDecision) {
 		permissionDismissed = true;
-		agenticResolvePermission(message.convId, decision);
+		agenticStore.resolvePermission(message.convId, decision);
 	}
 
 	let continueDismissed = $state(false);
 
 	const pendingContinue = $derived(
-		isStreaming && isLastAssistantMessage ? agenticPendingContinueRequest(message.convId) : false
+		isStreaming && isLastAssistantMessage
+			? agenticStore.pendingContinueRequest(message.convId)
+			: false
 	);
 
 	let prevContinueRef = false;
@@ -90,13 +91,13 @@
 
 	function handleContinue(shouldContinue: boolean) {
 		continueDismissed = true;
-		agenticResolveContinue(message.convId, shouldContinue);
+		agenticStore.resolveContinue(message.convId, shouldContinue);
 	}
 
 	const sections = $derived(deriveAgenticSections(message, toolMessages, [], isStreaming));
 
 	const currentlyExecutingToolCallId = $derived(
-		isStreaming ? agenticExecutingToolCallId(message.convId) : null
+		isStreaming ? agenticStore.executingToolCallId(message.convId) : null
 	);
 
 	type TurnGroup = {

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

@@ -3,12 +3,12 @@
 	import { ChatForm, DialogConfirmation } from '$lib/components/app';
 	import { Button } from '$lib/components/ui/button';
 	import { Switch } from '$lib/components/ui/switch';
-	import { getMessageEditContext } from '$lib/contexts';
+	import { getChatMessageEditContext } from '$lib/contexts';
 	import { KeyboardKey, MessageRole } from '$lib/enums';
-	import { chatStore } from '$lib/stores/chat.svelte';
+	import { chatStore } from '$lib/stores';
 	import { processFilesToChatUploaded } from '$lib/utils/browser-only';
 
-	const editCtx = getMessageEditContext();
+	const editCtx = getChatMessageEditContext();
 
 	let saveWithoutRegenerate = $state(false);
 	let showDiscardDialog = $state(false);

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

@@ -3,9 +3,8 @@
 	import { CollapsibleContentBlock, MarkdownContent } from '$lib/components/app';
 	import { REASONING_SCROLL_AT_BOTTOM_THRESHOLD_PX } from '$lib/constants';
 	import { AgenticSectionType } from '$lib/enums';
-	import { config } from '$lib/stores/settings.svelte';
-	import type { DatabaseMessageExtra } from '$lib/types';
-	import type { AgenticSection } from '$lib/utils';
+	import { settingsStore } from '$lib/stores';
+	import type { AgenticSection, DatabaseMessageExtra } from '$lib/types';
 
 	interface Props {
 		section: AgenticSection;
@@ -25,7 +24,7 @@
 		section
 	}: Props = $props();
 
-	const currentConfig = config();
+	const currentConfig = settingsStore.config;
 
 	const REASONING_HEADER = 'Reasoning';
 	const REASONING_HEADER_PENDING = 'Reasoning...';

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

@@ -1,22 +1,8 @@
 <script lang="ts">
 	import { ChatMessage, ChatMessageUserPending } from '$lib/components/app';
-	import { setChatActionsContext } from '$lib/contexts';
 	import { MessageRole } from '$lib/enums';
-	import {
-		agenticClearSteeringMessage,
-		agenticInjectSteeringMessage,
-		agenticPendingSteeringMessageContent,
-		agenticPendingSteeringMessageExtras
-	} from '$lib/stores/agentic.svelte';
-	import { chatStore } from '$lib/stores/chat.svelte';
-	import {
-		chatClearPendingMessage,
-		chatInjectPendingMessage,
-		chatPendingMessageContent,
-		chatPendingMessageExtras
-	} from '$lib/stores/chat.svelte';
-	import { activeConversation, conversationsStore } from '$lib/stores/conversations.svelte';
-	import { config } from '$lib/stores/settings.svelte';
+	import { agenticStore, chatStore, conversationsStore, settingsStore } from '$lib/stores';
+	import type { ChatMessageActions } from '$lib/types';
 	import {
 		buildSiblingInfoMap,
 		copyToClipboard,
@@ -34,9 +20,9 @@
 
 	let allConversationMessages = $state<DatabaseMessage[]>([]);
 
-	const currentConfig = config();
+	const currentConfig = settingsStore.config;
 
-	setChatActionsContext({
+	const chatActions: ChatMessageActions = {
 		continueAssistantMessage: async (message: DatabaseMessage) => {
 			onUserAction?.();
 			await chatStore.continueAssistantMessage(message.id);
@@ -105,10 +91,10 @@
 			await chatStore.regenerateMessageWithBranching(message.id, modelOverride);
 			refreshAllMessages();
 		}
-	});
+	};
 
 	function refreshAllMessages() {
-		const conversation = activeConversation();
+		const conversation = conversationsStore.activeConversation;
 
 		if (conversation) {
 			conversationsStore.getConversationMessages(conversation.id).then((messages) => {
@@ -121,7 +107,7 @@
 
 	// Refresh messages whenever the active conversation changes
 	$effect(() => {
-		if (activeConversation()) {
+		if (conversationsStore.activeConversation) {
 			refreshAllMessages();
 		}
 	});
@@ -242,6 +228,7 @@
 	{#each displayMessages as { isLastAssistantMessage, isLastUserMessage, message, nextAssistantMessage, siblingInfo, toolMessages } (message.id)}
 		<ChatMessage
 			class="mx-auto mt-12 w-full max-w-3xl"
+			{chatActions}
 			{message}
 			{toolMessages}
 			{isLastAssistantMessage}
@@ -251,32 +238,33 @@
 		/>
 	{/each}
 
-	{#if activeConversation() && agenticPendingSteeringMessageContent(activeConversation()!.id)}
-		{@const convId = activeConversation()!.id}
-		{@const pendingContent = agenticPendingSteeringMessageContent(convId)}
+	{#if conversationsStore.activeConversation && agenticStore.pendingSteeringMessageContent(conversationsStore.activeConversation!.id)}
+		{@const convId = conversationsStore.activeConversation!.id}
+		{@const pendingContent = agenticStore.pendingSteeringMessageContent(convId)}
 
 		{#if pendingContent}
 			<ChatMessageUserPending
 				class="mx-auto mt-12 w-full max-w-[48rem]"
 				content={pendingContent}
-				extras={agenticPendingSteeringMessageExtras(convId)}
+				extras={agenticStore.pendingSteeringMessageExtras(convId)}
 				onSendImmediately={() => chatStore.abortCurrentFlow(convId)}
-				onEdit={(newContent, extras) => agenticInjectSteeringMessage(convId, newContent, extras)}
-				onDelete={() => agenticClearSteeringMessage(convId)}
+				onEdit={(newContent, extras) =>
+					agenticStore.injectSteeringMessage(convId, newContent, extras)}
+				onDelete={() => agenticStore.clearSteeringMessage(convId)}
 			/>
 		{/if}
-	{:else if activeConversation() && chatPendingMessageContent(activeConversation()!.id)}
-		{@const convId = activeConversation()!.id}
-		{@const pendingContent = chatPendingMessageContent(convId)}
+	{:else if conversationsStore.activeConversation && chatStore.pendingMessageContent(conversationsStore.activeConversation!.id)}
+		{@const convId = conversationsStore.activeConversation!.id}
+		{@const pendingContent = chatStore.pendingMessageContent(convId)}
 
 		{#if pendingContent}
 			<ChatMessageUserPending
 				class="mx-auto mt-12 w-full max-w-[48rem]"
 				content={pendingContent}
-				extras={chatPendingMessageExtras(convId)}
+				extras={chatStore.pendingMessageExtras(convId)}
 				onSendImmediately={() => chatStore.abortCurrentFlow(convId)}
-				onEdit={(newContent, extras) => chatInjectPendingMessage(convId, newContent, extras)}
-				onDelete={() => chatClearPendingMessage(convId)}
+				onEdit={(newContent, extras) => chatStore.injectPendingMessage(convId, newContent, extras)}
+				onDelete={() => chatStore.clearPendingMessage(convId)}
 			/>
 		{/if}
 	{/if}

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

@@ -20,26 +20,20 @@
 	import { useKeyboardShortcuts } from '$lib/hooks/use-keyboard-shortcuts.svelte';
 	import {
 		chatStore,
-		errorDialog,
-		isChatStreaming,
-		isEditing,
-		isLoading
-	} from '$lib/stores/chat.svelte';
-	import {
-		activeConversation,
-		activeMessages,
-		conversationsStore
-	} from '$lib/stores/conversations.svelte';
-	import { device } from '$lib/stores/device.svelte';
-	import { serverError, serverLoading } from '$lib/stores/server.svelte';
-	import { config } from '$lib/stores/settings.svelte';
-	import { isMobile } from '$lib/stores/viewport.svelte';
+		conversationsStore,
+		device,
+		isMobile,
+		serverStore,
+		settingsStore
+	} from '$lib/stores';
 	import { parseFilesToMessageExtras } from '$lib/utils/browser-only';
 	import { onDestroy, onMount, tick } from 'svelte';
 
 	let { showCenteredEmpty = false } = $props();
 
-	let disableAutoScroll = $derived(Boolean(config().disableAutoScroll) || isMobile.current);
+	let disableAutoScroll = $derived(
+		Boolean(settingsStore.config.disableAutoScroll) || isMobile.current
+	);
 	let isMobileUserScrolledUp = $state(false);
 	let mobileScrollDownHint = $state(false);
 	let mobileScrollDownHintLockedUntil = $state(0);
@@ -48,12 +42,15 @@
 	let showDeleteDialog = $state(false);
 	let showEmptyFileDialog = $state(false);
 	let isEmpty = $derived(
-		showCenteredEmpty && !activeConversation() && activeMessages().length === 0 && !isLoading()
+		showCenteredEmpty &&
+			!conversationsStore.activeConversation &&
+			conversationsStore.activeMessages.length === 0 &&
+			!chatStore.isLoading
 	);
-	let activeErrorDialog = $derived(errorDialog());
-	let isServerLoading = $derived(serverLoading());
-	let hasPropsError = $derived(!!serverError());
-	let isCurrentConversationLoading = $derived(isLoading() || isChatStreaming());
+	let activeErrorDialog = $derived(chatStore.errorDialogState);
+	let isServerLoading = $derived(serverStore.loading);
+	let hasPropsError = $derived(!!serverStore.error);
+	let isCurrentConversationLoading = $derived(chatStore.isLoading || chatStore.isStreaming());
 	let chatFormBottomPosition = $derived.by(() => {
 		if (!isMobile.current) return '1rem';
 
@@ -80,7 +77,7 @@
 	});
 	const { handleKeydown } = useKeyboardShortcuts({
 		deleteActiveConversation: () => {
-			if (activeConversation()) {
+			if (conversationsStore.activeConversation) {
 				showDeleteDialog = true;
 			}
 		}
@@ -100,7 +97,7 @@
 	}
 
 	async function handleDeleteConfirm() {
-		const conversation = activeConversation();
+		const conversation = conversationsStore.activeConversation;
 
 		if (conversation) {
 			await conversationsStore.deleteConversation(conversation.id);
@@ -148,7 +145,7 @@
 	async function handleMessagesReady(messageCount: number) {
 		if (messageCount === 0) return;
 
-		const id = activeConversation()?.id ?? null;
+		const id = conversationsStore.activeConversation?.id ?? null;
 
 		if (!id || id === lastScrolledConversationId) return;
 
@@ -168,7 +165,7 @@
 		const settle = () => {
 			if (autoScroll.userScrolledUp) return;
 
-			if (activeConversation()?.id !== id) return;
+			if (conversationsStore.activeConversation?.id !== id) return;
 
 			autoScroll.scrollToBottom();
 			const height = container.scrollHeight;
@@ -246,7 +243,7 @@
 
 	$effect(() => {
 		const shouldDisableAutoScroll =
-			config().disableAutoScroll || (isMobile.current && isCurrentConversationLoading);
+			settingsStore.config.disableAutoScroll || (isMobile.current && isCurrentConversationLoading);
 
 		autoScroll.setDisabled(shouldDisableAutoScroll);
 
@@ -310,7 +307,7 @@
 	>
 		{#if !isEmpty}
 			<ChatMessages
-				messages={activeMessages()}
+				messages={conversationsStore.activeMessages}
 				onMessagesReady={handleMessagesReady}
 				onUserAction={() => {
 					handleSendLikeScroll();
@@ -354,7 +351,7 @@
 
 			<ChatScreenForm
 				class="pointer-events-auto conversation-chat-form"
-				disabled={hasPropsError || isEditing()}
+				disabled={hasPropsError || chatStore.isEditing()}
 				{initialMessage}
 				isLoading={isCurrentConversationLoading}
 				onFileRemove={fileUpload.handleFileRemove}

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

@@ -3,7 +3,7 @@
 	import { page } from '$app/state';
 	import { ChatForm } from '$lib/components/app';
 	import { useDraftMessages } from '$lib/hooks/use-draft-messages.svelte';
-	import { isMobile } from '$lib/stores/viewport.svelte';
+	import { isMobile } from '$lib/stores';
 	import { onMount } from 'svelte';
 
 	interface Props {

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

@@ -1,5 +1,5 @@
 <script lang="ts">
-	import { serverStore } from '$lib/stores/server.svelte';
+	import { serverStore } from '$lib/stores';
 
 	interface Props {
 		isEmpty: boolean;

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

@@ -2,10 +2,10 @@
 	import { AlertTriangle, Loader2, RefreshCw } from '@lucide/svelte';
 	import * as Alert from '$lib/components/ui/alert';
 	import { ICON_CLASS_DEFAULT } from '$lib/constants';
-	import { serverError, serverLoading, serverStatus, serverStore } from '$lib/stores/server.svelte';
+	import { serverStore } from '$lib/stores';
 
-	let hasError = $derived(!!serverError());
-	let isLoadingModel = $derived(serverStatus() === 503);
+	let hasError = $derived(!!serverStore.error);
+	let isLoadingModel = $derived(serverStore.status === 503);
 </script>
 
 {#if hasError}
@@ -23,17 +23,17 @@
 				{#if !isLoadingModel}
 					<button
 						onclick={() => serverStore.fetch()}
-						disabled={serverLoading()}
+						disabled={serverStore.loading}
 						class="flex items-center gap-1.5 rounded-lg bg-destructive/20 px-2 py-1 text-xs font-medium hover:bg-destructive/30 disabled:opacity-50"
 					>
-						<RefreshCw class="h-3 w-3 {serverLoading() ? 'animate-spin' : ''}" />
-						{serverLoading() ? 'Retrying...' : 'Retry'}
+						<RefreshCw class="h-3 w-3 {serverStore.loading ? 'animate-spin' : ''}" />
+						{serverStore.loading ? 'Retrying...' : 'Retry'}
 					</button>
 				{/if}
 			</Alert.Title>
 
 			{#if !isLoadingModel}
-				<Alert.Description>{serverError()}</Alert.Description>
+				<Alert.Description>{serverStore.error}</Alert.Description>
 			{/if}
 		</Alert.Root>
 	</div>

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

@@ -1,7 +1,7 @@
 <script lang="ts">
 	import { Loader2 } from '@lucide/svelte';
 	import { StreamConnectionState } from '$lib/enums';
-	import { chatStore } from '$lib/stores/chat.svelte';
+	import { chatStore } from '$lib/stores';
 
 	let state = $derived(chatStore.streamConnectionState);
 </script>

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

@@ -120,7 +120,7 @@ export { default as ChatAttachmentsPreviewCurrentItem } from './ChatAttachments/
  * Used by ChatScreenForm and ChatMessageEditForm for both new conversations and message editing.
  *
  * **Architecture:**
- * - Composes ChatFormTextarea (or ChatFormContenteditable for messages with
+ * - Composes ChatFormTextarea (or ChatFormContentEditable for messages with
  *   file mention links), ChatFormActions, and ChatFormPickerMcpPrompts
  * - Manages file upload state via `uploadedFiles` bindable prop
  * - Integrates with ModelsSelectorDropdown for model selection in router mode
@@ -272,7 +272,7 @@ export { default as ChatFormMcpResourcesList } from './ChatForm/ChatFormMcpResou
  * source string. ChatForm swaps it in once a mention link lands in the
  * buffer. Shares the focus()/resetHeight()/caret handle with the textarea.
  */
-export { default as ChatFormContenteditable } from './ChatForm/ChatFormContenteditable.svelte';
+export { default as ChatFormContentEditable } from './ChatForm/ChatFormContentEditable.svelte';
 
 /**
  * Plain auto-resizing textarea with IME composition support. Default input

tools/ui/src/lib/components/app/content/MarkdownContent/MarkdownContent.svelte

@@ -1,5 +1,5 @@
 <script lang="ts">
-	import '$styles/katex-custom.scss';
+	import '$lib/styles/katex-custom.scss';
 	import {
 		getCodeInfoFromTarget,
 		getHastNodeId,
@@ -45,7 +45,7 @@
 	import { ColorMode, UrlProtocol } from '$lib/enums';
 	import { FileTypeText } from '$lib/enums/files.enums';
 	import { createAutoScrollController } from '$lib/hooks/use-auto-scroll.svelte';
-	import { config } from '$lib/stores/settings.svelte';
+	import { settingsStore } from '$lib/stores';
 	import type { DatabaseMessageExtra } from '$lib/types/database';
 	import {
 		copyCodeToClipboard,
@@ -864,7 +864,7 @@
 <div
 	bind:this={containerRef}
 	onclick={handleMermaidClick}
-	class="markdown-content {className}{config()[SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS]
+	class="markdown-content {className}{settingsStore.config[SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS]
 		? ' full-height-code-blocks'
 		: ''}"
 >

tools/ui/src/lib/components/app/content/MarkdownContent/plugins/rehype/file-badge.ts

@@ -16,8 +16,7 @@ import {
 	PATH_SEPARATOR,
 	SETTINGS_KEYS
 } from '$lib/constants';
-import { settingsStore } from '$lib/stores/settings.svelte';
-import { toolsStore } from '$lib/stores/tools.svelte';
+import { settingsStore, toolsStore } from '$lib/stores';
 import { decodeFileLinkPath, getMentionBadgeIconPaths, getMentionBadgeLabel } from '$lib/utils';
 import type { Element, Root } from 'hast';
 import type { Plugin } from 'unified';

tools/ui/src/lib/components/app/content/MentionBadge.svelte

@@ -1,7 +1,6 @@
 <script lang="ts">
 	import { SETTINGS_KEYS } from '$lib/constants';
-	import { settingsStore } from '$lib/stores/settings.svelte';
-	import { toolsStore } from '$lib/stores/tools.svelte';
+	import { settingsStore, toolsStore } from '$lib/stores';
 	import {
 		getMentionBadgeIconPaths,
 		getMentionBadgeLabel,

tools/ui/src/lib/components/app/dialogs/DialogMcpResourcePreview.svelte

@@ -5,7 +5,7 @@
 	import * as Dialog from '$lib/components/ui/dialog';
 	import { DEFAULT_RESOURCE_FILENAME, MIME_TYPE_SUBSTRINGS } from '$lib/constants';
 	import { MimeTypeText } from '$lib/enums';
-	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import { mcpStore } from '$lib/stores';
 	import type { DatabaseMessageExtraMcpResource } from '$lib/types';
 	import {
 		downloadResourceContent,