server: Ignore empty / non-existing Origin headers (#25756)

Otherwise this gives lots of unnecessary warnings:

  W srv    operator(): (CORS) skip non-localhost origin:

.github/workflows/release.yml

@@ -1651,6 +1651,9 @@ jobs:
 
             </details>
 
+            **Website:**
+            - <https://llama.app>
+
             **macOS/iOS:**
             - [macOS Apple Silicon (arm64)](https://github.com/ggml-org/llama.cpp/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-macos-arm64.tar.gz)
             - macOS Apple Silicon (arm64, KleidiAI enabled) [DISABLED](https://github.com/ggml-org/llama.cpp/pull/23780)

common/arg.cpp

@@ -2932,7 +2932,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
     ).set_examples({LLAMA_EXAMPLE_TOKENIZE}));
     add_opt(common_arg(
         {"--stdin"},
-        string_format("read the prompt from stdin (mutually exclusive with -f/--file and -p/--prompt) (default: %s)", params.tokenize_stdin ? "true" : "false"),
+        string_format("read the prompt from stdin (takes precedence over -f/--file and -p/--prompt) (default: %s)", params.tokenize_stdin ? "true" : "false"),
         [](common_params & params) {
             params.tokenize_stdin = true;
         }

common/speculative.cpp

@@ -260,7 +260,10 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
     bool process(const llama_batch & batch) override {
         auto * ctx_dft = params.ctx_dft;
 
-        const int ret = llama_decode(ctx_dft, batch);
+        llama_batch batch_dft = batch;
+        batch_dft.logits = nullptr;
+
+        const int ret = llama_decode(ctx_dft, batch_dft);
 
         if (ret != 0) {
             SPC_ERR("failed to decode draft batch, ret = %d\n", ret);

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

@@ -65,6 +65,7 @@
 #include "ggml-cuda/tri.cuh"
 #include "ggml-cuda/cumsum.cuh"
 #include "ggml-cuda/fill.cuh"
+#include "ggml-cuda/lightning-indexer.cuh"
 #include "ggml.h"
 
 #include <algorithm>
@@ -247,7 +248,11 @@ static ggml_cuda_device_info ggml_cuda_init() {
 
         info.default_tensor_split[id] = total_vram;
         total_vram += prop.totalGlobalMem;
+#if defined(GGML_USE_HIP)
+        info.devices[id].integrated = prop.integrated;
+#else
         info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034)
+#endif
         info.devices[id].nsm        = prop.multiProcessorCount;
         info.devices[id].smpb       = prop.sharedMemPerBlock;
         info.devices[id].warp_size  = prop.warpSize;
@@ -2257,6 +2262,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
         case GGML_OP_FILL:
             ggml_cuda_op_fill(ctx, dst);
             break;
+        case GGML_OP_LIGHTNING_INDEXER:
+            ggml_cuda_lightning_indexer(ctx, dst);
+            break;
         default:
             return false;
     }
@@ -4987,6 +4995,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
         case GGML_OP_DIAG:
         case GGML_OP_SOLVE_TRI:
             return true;
+        case GGML_OP_LIGHTNING_INDEXER:
+            return ggml_cuda_lightning_indexer_supported(dev_ctx->device, op);
 
         default:
             return false;

ggml/src/ggml-cuda/lightning-indexer.cu

@@ -0,0 +1,588 @@
+#include "common.cuh"
+#include "lightning-indexer.cuh"
+#include "fattn-common.cuh"
+#include "convert.cuh"
+
+#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+#if defined(TURING_MMA_AVAILABLE)
+
+typedef union {
+    int2 i2;
+    half2 h2[2];
+} half4;
+
+// TODO add support for AMD cards via rocWMMA
+#include <mma.h>
+namespace wmma = nvcuda::wmma;
+
+template <int WARPS_PER_BLOCK, int K_VECS_PER_BLOCK, int64_t N_EMBD, int64_t N_HEAD, ggml_type TYPE_K>
+static __global__ void lightning_indexer_kernel_wmma(
+        const float * Q, const char * K, const float * W, const half * M, float * dst,
+        int64_t n_stream, int64_t n_batch, int64_t n_kv,
+        size_t nb1, size_t nb2, size_t nb3,
+        size_t nbq1, size_t nbq2, size_t nbq3,
+        size_t nbk1, size_t nbk2, size_t nbk3,
+        size_t nbw1, size_t nbw2, size_t nbw3,
+        size_t nbm1, size_t nbm2, size_t nbm3,
+        int64_t nem3
+    ) {
+
+    constexpr int THREADS_PER_BLOCK = WARPS_PER_BLOCK * WARP_SIZE;
+    constexpr int HEADS_PER_INNER_LOOP = 8;
+    constexpr int K_EMBD_PER_INNER_LOOP = 16;
+    constexpr int N_EMBD_PADDED = N_EMBD + 8;
+
+    const int i_batch  = blockIdx.y;
+    const int i_stream = blockIdx.z;
+    const int i_warp   = threadIdx.y;
+    const int i_lane   = threadIdx.x;
+    const int tid      = i_warp * WARP_SIZE + i_lane;
+
+    // each block processes K_VECS_PER_BLOCK K vectors
+    const int start_kv = blockIdx.x * K_VECS_PER_BLOCK;
+
+    const char  * q_base = (const char  *)                 Q + i_batch*nbq2 + i_stream*nbq3;
+    const float * w_base = (const float *) ((const char *) W + i_batch*nbw1 + i_stream*nbw3);
+
+    // phase 1 - load weights and first Q tile to shared memory
+
+    __shared__ float w_shared[N_HEAD];
+    __shared__ int2  q_shared_h[HEADS_PER_INNER_LOOP][N_EMBD_PADDED / 4];
+
+    if (tid < N_HEAD) {
+        w_shared[tid] = w_base[tid];
+    }
+
+    // total number of half4 elements in HEADS_PER_INNER_LOOP x N_EMBD Q tile
+    constexpr int N_Q_TILE = HEADS_PER_INNER_LOOP * (N_EMBD / 4);
+    // number of registers needed in each thread to store Q tile in thread block
+    constexpr int N_Q_NEXT = (N_Q_TILE + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
+
+#pragma unroll
+    for (int i_q = tid; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) {
+        const int i_head = i_q / (N_EMBD / 4);
+        const int i_embd = i_q % (N_EMBD / 4);
+        const float4 q = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4));
+        half4 q_packed;
+        q_packed.h2[0] = __float22half2_rn(make_float2(q.x, q.y));
+        q_packed.h2[1] = __float22half2_rn(make_float2(q.z, q.w));
+        q_shared_h[i_head][i_embd] = q_packed.i2;
+    }
+
+    // phase 2 - load (and dequantize if needed) K to shared mem
+
+    __shared__ half2 k_shared_h[K_VECS_PER_BLOCK][N_EMBD_PADDED / 4][2];
+
+    constexpr int n_k = K_VECS_PER_BLOCK * (N_EMBD / 4);
+
+    if constexpr (TYPE_K == GGML_TYPE_F16) {
+#pragma unroll
+        for (int i_k = tid; i_k < n_k; i_k += THREADS_PER_BLOCK) {
+            const int i_k_vec = i_k / (N_EMBD / 4);
+            const int i_embd = i_k % (N_EMBD / 4);
+            const int i_kv = start_kv + i_k_vec;
+            if (i_kv < n_kv) {
+                const int2 * k_base = (const int2 *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3);
+                *(int2*) &k_shared_h[i_k_vec][i_embd] = k_base[i_embd];
+            } else {
+                *(int2*) &k_shared_h[i_k_vec][i_embd] = make_int2(0, 0);
+            }
+        }
+    } else {
+        constexpr dequantize_V_t dequantize_k = get_dequantize_V<TYPE_K, half, 4>();
+#pragma unroll
+        for (int i_k = tid; i_k < n_k; i_k += THREADS_PER_BLOCK) {
+            const int i_k_vec = i_k / (N_EMBD / 4);
+            const int i_embd = i_k % (N_EMBD / 4);
+            const int i_kv = start_kv + i_k_vec;
+            if (i_kv < n_kv) {
+                const void * k_base = (const void *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3);
+                dequantize_k(k_base, &k_shared_h[i_k_vec][i_embd][0], i_embd * 4);
+            } else {
+                *(int2*) &k_shared_h[i_k_vec][i_embd] = make_int2(0, 0);
+            }
+        }
+    }
+
+    __syncthreads();
+
+    // phase 3 - calculate lightning indexer scores
+
+    __shared__ float qk_shared[WARPS_PER_BLOCK][HEADS_PER_INNER_LOOP][K_VECS_PER_BLOCK];
+
+    // load K fragment
+    wmma::fragment<wmma::matrix_b, HEADS_PER_INNER_LOOP, K_VECS_PER_BLOCK, K_EMBD_PER_INNER_LOOP, half, wmma::col_major> frag_k;
+    wmma::load_matrix_sync(frag_k, (half*) &k_shared_h[0][i_warp * K_EMBD_PER_INNER_LOOP / 4], N_EMBD_PADDED);
+
+    float score_k = 0.0f;
+
+    for (int i_head_0 = 0; i_head_0 < N_HEAD; i_head_0 += HEADS_PER_INNER_LOOP) {
+        const int i_head_next = i_head_0 + HEADS_PER_INNER_LOOP;
+
+        // we don't use accumulator for anything, fill it with zeros
+        wmma::fragment<wmma::accumulator, HEADS_PER_INNER_LOOP, K_VECS_PER_BLOCK, K_EMBD_PER_INNER_LOOP, float> frag_acc;
+        wmma::fill_fragment(frag_acc, 0.0f);
+
+        // load Q fragment
+        wmma::fragment<wmma::matrix_a, HEADS_PER_INNER_LOOP, K_VECS_PER_BLOCK, K_EMBD_PER_INNER_LOOP, half, wmma::row_major> frag_q;
+        wmma::load_matrix_sync(frag_q, (half*) &q_shared_h[0][i_warp * K_EMBD_PER_INNER_LOOP / 4], N_EMBD_PADDED);
+
+        // preload next Q tile to registers during matrix multiplication
+        float4 q_next[N_Q_NEXT];
+
+        if (i_head_next < N_HEAD) {
+#pragma unroll
+            for (int i_q = tid, i_q_next = 0; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) {
+                const int i_head = i_head_next + i_q / (N_EMBD / 4);
+                const int i_embd =               i_q % (N_EMBD / 4);
+                q_next[i_q_next++] = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4));
+            }
+        }
+
+        // perform matrix multiplication
+        wmma::mma_sync(frag_acc, frag_q, frag_k, frag_acc);
+        wmma::store_matrix_sync((float*) &qk_shared[i_warp][0][0], frag_acc, K_VECS_PER_BLOCK, wmma::mem_row_major);
+
+        // make sure all threads finished using q_shared_h so we can store next tile
+        __syncthreads();
+
+        // write preloaded Q tile to shared memory
+        if (i_head_next < N_HEAD) {
+#pragma unroll
+            for (int i_q = tid, i_q_next = 0; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) {
+                const int i_head = i_q / (N_EMBD / 4);
+                const int i_embd = i_q % (N_EMBD / 4);
+                half4 q_packed;
+                q_packed.h2[0] = __float22half2_rn(make_float2(q_next[i_q_next].x, q_next[i_q_next].y));
+                q_packed.h2[1] = __float22half2_rn(make_float2(q_next[i_q_next].z, q_next[i_q_next].w));
+                q_shared_h[i_head][i_embd] = q_packed.i2;
+                ++i_q_next;
+            }
+        }
+
+        // accumulate QK multiplication results from all block warps
+        // (there are 256 threads in block and 256 matmul outputs)
+        // TODO it will break if WARP_SIZE is not 32
+        const int h = tid / K_VECS_PER_BLOCK;
+        const int k = tid % K_VECS_PER_BLOCK;
+        const float w_val = w_shared[i_head_0 + h];
+
+        float sum = 0.0f;
+#pragma unroll
+        for (int w = 0; w < WARPS_PER_BLOCK; ++w) {
+            sum += qk_shared[w][h][k];
+        }
+
+        // ReLU, weight
+        sum = sum > 0.0f ? sum : 0.0f;
+        sum *= w_val;
+
+        // wait until qk_shared[0] is no longer used
+        __syncthreads();
+
+        // reuse qk_shared[0] for storing partial results
+        qk_shared[0][h][k] = sum;
+
+        // wait until all threads write their results
+        __syncthreads();
+
+        // accumulate result over heads
+        if (tid < K_VECS_PER_BLOCK) {
+#pragma unroll
+            for (int i_head = 0; i_head < HEADS_PER_INNER_LOOP; ++i_head) {
+                score_k += qk_shared[0][i_head][tid];
+            }
+        }
+
+        // make sure all threads finished using qk_shared
+        __syncthreads();
+    }
+
+    // phase 4 - store output to VRAM
+
+    if (tid < K_VECS_PER_BLOCK) {
+        const int i_kv = start_kv + tid;
+        if (i_kv < n_kv) {
+            const half * m_base = (const half *) ((const char *) M + i_batch*nbm1 + (i_stream%nem3)*nbm3);
+            float * dst_base = (float *) ((char *) dst + i_batch*nb1 + i_stream*nb3);
+            dst_base[i_kv] = score_k + __half2float(m_base[i_kv]);
+        }
+    }
+}
+
+#else // defined(TURING_MMA_AVAILABLE)
+
+template <int WARPS_PER_BLOCK, int K_VECS_PER_BLOCK, int64_t N_EMBD, int64_t N_HEAD, ggml_type TYPE_K>
+static __global__ void lightning_indexer_kernel_wmma(
+        const float * Q, const char * K, const float * W, const half * M, float * dst,
+        int64_t n_stream, int64_t n_batch, int64_t n_kv,
+        size_t nb1, size_t nb2, size_t nb3,
+        size_t nbq1, size_t nbq2, size_t nbq3,
+        size_t nbk1, size_t nbk2, size_t nbk3,
+        size_t nbw1, size_t nbw2, size_t nbw3,
+        size_t nbm1, size_t nbm2, size_t nbm3,
+        int64_t nem3
+    ) {
+    GGML_UNUSED_VARS(Q, K, W, M, dst,
+        n_stream, n_batch, n_kv,
+        nb1, nb2, nb3,
+        nbq1, nbq2, nbq3,
+        nbk1, nbk2, nbk3,
+        nbw1, nbw2, nbw3,
+        nem3);
+    NO_DEVICE_CODE;
+}
+
+#endif // defined(TURING_MMA_AVAILABLE)
+#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+
+// TODO there is one ugly assumption used in this kernel - that WARP_SIZE is equal to 32
+// thanks to that one warp operating on float4 processes whole indexer K/Q vectors
+// 32 * 4 = 128 (N_EMBD)
+
+template <int WARPS_PER_BLOCK, int K_VECS_PER_BLOCK, int64_t N_EMBD, int64_t N_HEAD, ggml_type TYPE_K>
+static __global__ void lightning_indexer_kernel_vec(
+        const float * Q, const char * K, const float * W, const half * M, float * dst,
+        int64_t n_stream, int64_t n_batch, int64_t n_kv,
+        size_t nb1, size_t nb2, size_t nb3,
+        size_t nbq1, size_t nbq2, size_t nbq3,
+        size_t nbk1, size_t nbk2, size_t nbk3,
+        size_t nbw1, size_t nbw2, size_t nbw3,
+        size_t nbm1, size_t nbm2, size_t nbm3,
+        int64_t nem3
+    ) {
+
+    constexpr int K_VECS_PER_WARP = K_VECS_PER_BLOCK / WARPS_PER_BLOCK;
+    constexpr int THREADS_PER_BLOCK = WARPS_PER_BLOCK * WARP_SIZE;
+
+    const int i_batch  = blockIdx.y;
+    const int i_stream = blockIdx.z;
+    const int i_warp   = threadIdx.y;
+    const int i_lane   = threadIdx.x;
+    const int tid      = i_warp * WARP_SIZE + i_lane;
+
+    // each warp processes K_VECS_PER_WARP K vectors
+    const int start_kv_block = blockIdx.x * K_VECS_PER_BLOCK;
+    const int start_kv = start_kv_block + i_warp * K_VECS_PER_WARP;
+
+    const char  * q_base = (const char  *)                 Q + i_batch*nbq2 + i_stream*nbq3;
+    const float * w_base = (const float *) ((const char *) W + i_batch*nbw1 + i_stream*nbw3);
+
+    // phase 1 - load (and dequantize if needed) K to registers
+
+    float4 k_reg_f[K_VECS_PER_WARP];
+
+    if constexpr (TYPE_K == GGML_TYPE_F32) {
+        // direct copy of float4
+#pragma unroll
+        for (int k = 0; k < K_VECS_PER_WARP; ++k) {
+            int i_kv = start_kv + k;
+            if (i_kv < n_kv) {
+                const float4 * k_base = (const float4 *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3);
+                k_reg_f[k] = k_base[i_lane];
+            } else {
+                k_reg_f[k] = make_float4(0, 0, 0, 0);
+            }
+        }
+    } else {
+        // dequantize remaining types to float
+        constexpr dequantize_V_t dequantize_k = get_dequantize_V<TYPE_K, float, 4>();
+#pragma unroll
+        for (int k = 0; k < K_VECS_PER_WARP; ++k) {
+            int i_kv = start_kv + k;
+            if (i_kv < n_kv) {
+                const void * k_base = (const void *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3);
+                dequantize_k(k_base, &k_reg_f[k], i_lane * 4);
+            } else {
+                k_reg_f[k] = make_float4(0, 0, 0, 0);
+            }
+        }
+    }
+
+    float score_k[K_VECS_PER_WARP] = { 0.0f };
+
+    // load weights and Q only for N_HEAD_INNER heads at once to reduce shared memory usage
+    constexpr int N_HEAD_INNER = N_HEAD / 4;
+
+    for (int i_head_0 = 0; i_head_0 < N_HEAD; i_head_0 += N_HEAD_INNER) {
+        // phase 2 - load weights and Q to shared memory
+
+        __shared__ float  w_shared[N_HEAD_INNER];
+        __shared__ float4 q_shared_f[N_HEAD_INNER][N_EMBD / 4];
+
+        if (tid < N_HEAD_INNER) {
+            w_shared[tid] = w_base[i_head_0 + tid];
+        }
+
+        constexpr int n_q = N_HEAD_INNER * (N_EMBD / 4);
+#pragma unroll
+        for (int i_q = tid; i_q < n_q; i_q += THREADS_PER_BLOCK) {
+            const int i_head_inner = i_q / (N_EMBD / 4);
+            const int i_head = i_head_0 + i_head_inner;
+            const int i_embd = i_q % (N_EMBD / 4);
+            q_shared_f[i_head_inner][i_embd] = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4));
+        }
+
+        __syncthreads();
+
+        // phase 3 - calculate lightning indexer scores
+
+        for (int i_head_inner = 0; i_head_inner < N_HEAD_INNER; ++i_head_inner) {
+            const float w_val = w_shared[i_head_inner];
+            float qk[K_VECS_PER_WARP] = { 0.0f };
+
+            // dot product of floats
+            const float4 q_vec = q_shared_f[i_head_inner][i_lane];
+
+#pragma unroll
+            for (int k = 0; k < K_VECS_PER_WARP; ++k) {
+                ggml_cuda_mad(qk[k], q_vec.x, k_reg_f[k].x);
+                ggml_cuda_mad(qk[k], q_vec.y, k_reg_f[k].y);
+                ggml_cuda_mad(qk[k], q_vec.z, k_reg_f[k].z);
+                ggml_cuda_mad(qk[k], q_vec.w, k_reg_f[k].w);
+            }
+
+#pragma unroll
+            for (int k = 0; k < K_VECS_PER_WARP; ++k) {
+                float sum = warp_reduce_sum(qk[k]);
+
+                // ReLU, weight
+                if (i_lane == 0) {
+                    sum = (sum > 0.0f) ? sum : 0.0f;
+                    score_k[k] += sum * w_val;
+                }
+            }
+        }
+
+        __syncthreads();
+    }
+
+    // phase 4 - store outputs to shared memory
+
+    __shared__ float dst_shared[K_VECS_PER_BLOCK];
+
+    if (i_lane == 0) {
+#pragma unroll
+        for (int k = 0; k < K_VECS_PER_WARP; ++k) {
+            dst_shared[i_warp * K_VECS_PER_WARP + k] = score_k[k];
+        }
+    }
+
+    __syncthreads();
+
+    // phase 5 - write from shared memory to VRAM in coalesced manner
+
+    if (tid < K_VECS_PER_BLOCK) {
+        int i_kv = start_kv_block + tid;
+        if (i_kv < n_kv) {
+            const half * m_base = (const half *) ((const char *) M + i_batch*nbm1 + (i_stream%nem3)*nbm3);
+            float * dst_base = (float *) ((char *) dst + i_batch*nb1 + i_stream*nb3);
+            dst_base[i_kv] = dst_shared[tid] + __half2float(m_base[i_kv]);
+        }
+    }
+}
+
+#define LIGHTNING_INDEXER_CASE(lightning_indexer_kernel, n_embd, n_head, K, type_K)         \
+    if (K->type == (type_K)) {                                                              \
+        lightning_indexer_kernel<WARPS_PER_BLOCK, K_VECS_PER_BLOCK, n_embd, n_head, type_K> \
+            <<<grid, block, 0, ctx.stream()>>>(                                             \
+            q_d, k_d, w_d, m_d, dst_d,                                                      \
+            n_stream, n_batch, n_kv,                                                        \
+            nb1, nb2, nb3,                                                                  \
+            nbq1, nbq2, nbq3,                                                               \
+            nbk1, nbk2, nbk3,                                                               \
+            nbw1, nbw2, nbw3,                                                               \
+            nbm1, nbm2, nbm3,                                                               \
+            nem3                                                                            \
+        );                                                                                  \
+    } else
+
+void ggml_cuda_lightning_indexer(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+    const ggml_tensor * q = dst->src[0];
+    const ggml_tensor * k = dst->src[1];
+    const ggml_tensor * w = dst->src[2]; // weights
+    const ggml_tensor * m = dst->src[3]; // mask
+
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+    GGML_ASSERT(  q->type == GGML_TYPE_F32);
+    GGML_ASSERT(  w->type == GGML_TYPE_F32);
+    GGML_ASSERT(  m->type == GGML_TYPE_F16);
+
+    GGML_TENSOR_LOCALS(int64_t, neq,  q, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbq,  q, nb)
+    GGML_TENSOR_LOCALS(int64_t, nek,  k, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbk,  k, nb)
+    GGML_TENSOR_LOCALS(int64_t, new,  w, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbw,  w, nb)
+    GGML_TENSOR_LOCALS(int64_t, nem,  m, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbm,  m, nb)
+    GGML_TENSOR_LOCALS(int64_t, ne, dst, ne)
+    GGML_TENSOR_LOCALS(size_t,  nb, dst, nb)
+
+    // input tensor rows must be contiguous
+    GGML_ASSERT(nbq0 == ggml_type_size(q->type));
+    GGML_ASSERT(nbk0 == ggml_type_size(k->type));
+    GGML_ASSERT(nbw0 == ggml_type_size(w->type));
+    GGML_ASSERT(nbm0 == ggml_type_size(m->type));
+
+    // dst cannot be transposed or permuted
+    GGML_ASSERT(nb0 == sizeof(float));
+    GGML_ASSERT(nb0 <= nb1);
+    GGML_ASSERT(nb1 <= nb2);
+    GGML_ASSERT(nb2 <= nb3);
+
+    const int n_embd   = q->ne[0];
+    const int n_head   = q->ne[1];
+    const int n_batch  = q->ne[2];
+    const int n_stream = q->ne[3];
+    const int n_kv     = k->ne[2];
+
+    const float *   q_d = (const float *)   q->data;
+    const char  *   k_d = (const char  *)   k->data;
+    const float *   w_d = (const float *)   w->data;
+    const half  *   m_d = (const half  *)   m->data;
+    float       * dst_d = (      float *) dst->data;
+
+    const int device = ggml_cuda_get_device();
+    const int cc     = ggml_cuda_info().devices[device].cc;
+
+    if (n_embd == 128 && n_head == 64) {
+#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+        if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && k->type != GGML_TYPE_F32 && k->type != GGML_TYPE_BF16) {
+            // use wmma kernel
+            constexpr int K_VECS_PER_BLOCK = 32;
+            constexpr int WARPS_PER_BLOCK = 8;
+
+            dim3 block(32, WARPS_PER_BLOCK);
+            int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK);
+            dim3 grid(num_kv_blocks, n_batch, n_stream);
+
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_F16)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q4_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q4_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q5_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q5_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q8_0)
+            GGML_ABORT("fatal error");
+        } else {
+#else // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+        {
+#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+            // use vector kernel
+            constexpr int K_VECS_PER_WARP = 8;
+            constexpr int WARPS_PER_BLOCK = 8;
+            constexpr int K_VECS_PER_BLOCK = K_VECS_PER_WARP * WARPS_PER_BLOCK;
+
+            dim3 block(32, WARPS_PER_BLOCK);
+            int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK);
+            dim3 grid(num_kv_blocks, n_batch, n_stream);
+
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_F16)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q4_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q4_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q5_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q5_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q8_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_BF16)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_F32)
+            GGML_ABORT("fatal error");
+        }
+    } else if (n_embd == 128 && n_head == 32) {
+#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+        if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && k->type != GGML_TYPE_F32 && k->type != GGML_TYPE_BF16) {
+            // use wmma kernel
+            constexpr int K_VECS_PER_BLOCK = 32;
+            constexpr int WARPS_PER_BLOCK = 8;
+
+            dim3 block(32, WARPS_PER_BLOCK);
+            int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK);
+            dim3 grid(num_kv_blocks, n_batch, n_stream);
+
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_F16)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q4_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q4_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q5_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q5_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q8_0)
+            GGML_ABORT("fatal error");
+        } else {
+#else // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+        {
+#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+            // use vector kernel
+            constexpr int K_VECS_PER_WARP = 8;
+            constexpr int WARPS_PER_BLOCK = 8;
+            constexpr int K_VECS_PER_BLOCK = K_VECS_PER_WARP * WARPS_PER_BLOCK;
+
+            dim3 block(32, WARPS_PER_BLOCK);
+            int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK);
+            dim3 grid(num_kv_blocks, n_batch, n_stream);
+
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_F16)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q4_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q4_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q5_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q5_1)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q8_0)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_BF16)
+            LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_F32)
+            GGML_ABORT("fatal error");
+        }
+    } else {
+        GGML_ABORT("fatal error");
+    }
+}
+
+bool ggml_cuda_lightning_indexer_supported(int device, const ggml_tensor * dst) {
+    GGML_UNUSED(device);
+
+    const ggml_tensor * q = dst->src[0];
+    const ggml_tensor * k = dst->src[1];
+    const ggml_tensor * w = dst->src[2]; // weights
+    const ggml_tensor * m = dst->src[3]; // mask
+
+    GGML_TENSOR_LOCALS(int64_t, neq,  q, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbq,  q, nb)
+    GGML_TENSOR_LOCALS(int64_t, nek,  k, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbk,  k, nb)
+    GGML_TENSOR_LOCALS(int64_t, new,  w, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbw,  w, nb)
+    GGML_TENSOR_LOCALS(int64_t, nem,  m, ne)
+    GGML_TENSOR_LOCALS(size_t,  nbm,  m, nb)
+    GGML_TENSOR_LOCALS(int64_t, ne, dst, ne)
+    GGML_TENSOR_LOCALS(size_t,  nb, dst, nb)
+
+    if (neq0 != 128) {
+        return false;
+    }
+
+    if (neq1 != 64 && neq1 != 32) {
+        return false;
+    }
+
+    // alignment checks
+    for (const ggml_tensor * t : {q, k}) {
+        if (ggml_is_quantized(t->type)) {
+            continue;
+        }
+        for (size_t i = 1; i < GGML_MAX_DIMS; ++i) {
+            if (t->nb[i] % 16 != 0) {
+                return false;
+            }
+        }
+    }
+
+    switch(k->type) {
+        case GGML_TYPE_F32:
+        case GGML_TYPE_BF16:
+        case GGML_TYPE_F16:
+        case GGML_TYPE_Q8_0:
+        case GGML_TYPE_Q5_1:
+        case GGML_TYPE_Q5_0:
+        case GGML_TYPE_Q4_1:
+        case GGML_TYPE_Q4_0:
+            return true;
+        default:
+            return false;
+    }
+}

ggml/src/ggml-cuda/lightning-indexer.cuh

@@ -0,0 +1,4 @@
+#include "common.cuh"
+
+void ggml_cuda_lightning_indexer(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
+bool ggml_cuda_lightning_indexer_supported(int device, const ggml_tensor * dst);

ggml/src/ggml-cuda/mmf.cu

@@ -85,7 +85,7 @@ void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * sr
         GGML_ASSERT(sis1 > 0);
 
         ggml_cuda_launch_mm_ids_helper(ids_d, ids_src_compact_dev.get(), ids_dst_compact_dev.get(), expert_bounds_dev.get(),
-            static_cast<int>(n_experts), static_cast<int>(n_tokens), static_cast<int>(n_expert_used), static_cast<int>(ne11), si1, sis1, ctx.stream());
+            static_cast<int>(n_experts), static_cast<int>(n_tokens), static_cast<int>(n_expert_used), static_cast<int>(ne11), si1, sis1, /*write_inverse =*/ false, ctx.stream());
         CUDA_CHECK(cudaGetLastError());
 
         ids_info.ids_src_compact   = ids_src_compact_dev.get();

ggml/src/ggml-cuda/mmid.cu

@@ -27,7 +27,7 @@ template <int n_expert_used_template>
 __launch_bounds__(ggml_cuda_get_physical_warp_size(), 1)
 static __global__ void mm_ids_helper(
         const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds,
-        const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1) {
+        const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, const bool write_inverse) {
     constexpr int warp_size = ggml_cuda_get_physical_warp_size();
     const int n_expert_used = n_expert_used_template == 0 ? n_expert_used_var : n_expert_used_template;
     const int expert = blockIdx.x;
@@ -98,8 +98,13 @@ static __global__ void mm_ids_helper(
         const mm_ids_helper_store store_it = store[itc];
         const int it       = store_it.it();
         const int iex_used = store_it.iex_used();
-        ids_src1[nex_prev + itc] = it*sis1          + iex_used % nchannels_y;
-        ids_dst [nex_prev + itc] = it*n_expert_used + iex_used;
+        ids_dst[nex_prev + itc] = it*n_expert_used + iex_used;
+        // ids_src1 holds the forward map, or the inverse map (token slot -> compact row) for quant dedup
+        if (write_inverse) {
+            ids_src1[it*n_expert_used + iex_used] = nex_prev + itc;
+        } else {
+            ids_src1[nex_prev + itc] = it*sis1 + iex_used % nchannels_y;
+        }
     }
 
     if (threadIdx.x != 0) {
@@ -118,7 +123,7 @@ static __global__ void mm_ids_helper(
 template <int n_expert_used_template>
 static void launch_mm_ids_helper(
         const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds,
-        const int n_experts, const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, cudaStream_t stream) {
+        const int n_experts, const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream) {
     GGML_ASSERT(n_tokens          < (1 << 22) && "too few bits in mm_ids_helper_store");
     GGML_ASSERT(n_expert_used_var < (1 << 10) && "too few bits in mm_ids_helper_store");
 
@@ -132,33 +137,33 @@ static void launch_mm_ids_helper(
     const size_t nbytes_shared = n_tokens*sizeof(mm_ids_helper_store);
     GGML_ASSERT(nbytes_shared <= smpbo);
     mm_ids_helper<n_expert_used_template><<<num_blocks, block_size, nbytes_shared, stream>>>
-        (ids, ids_src1, ids_dst, expert_bounds, n_tokens, n_expert_used_var, nchannels_y, si1, sis1);
+        (ids, ids_src1, ids_dst, expert_bounds, n_tokens, n_expert_used_var, nchannels_y, si1, sis1, write_inverse);
 }
 
 void ggml_cuda_launch_mm_ids_helper(
         const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds,
-        const int n_experts, const int n_tokens, const int n_expert_used, const int nchannels_y, const int si1, const int sis1, cudaStream_t stream) {
+        const int n_experts, const int n_tokens, const int n_expert_used, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream) {
     switch (n_expert_used) {
         case  2:
-            launch_mm_ids_helper< 2>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream);
+            launch_mm_ids_helper< 2>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream);
             break;
         case  4:
-            launch_mm_ids_helper< 4>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream);
+            launch_mm_ids_helper< 4>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream);
             break;
         case  6:
-            launch_mm_ids_helper< 6>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream);
+            launch_mm_ids_helper< 6>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream);
             break;
         case  8:
-            launch_mm_ids_helper< 8>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream);
+            launch_mm_ids_helper< 8>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream);
             break;
         case 16:
-            launch_mm_ids_helper<16>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream);
+            launch_mm_ids_helper<16>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream);
             break;
         case 32:
-            launch_mm_ids_helper<32>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream);
+            launch_mm_ids_helper<32>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream);
             break;
         default:
-            launch_mm_ids_helper< 0>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, stream);
+            launch_mm_ids_helper< 0>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream);
             break;
     }
 }

ggml/src/ggml-cuda/mmid.cuh

@@ -2,4 +2,4 @@
 
 void ggml_cuda_launch_mm_ids_helper(
         const int32_t * ids, int32_t * ids_src1, int32_t * ids_dst, int32_t * expert_bounds,
-        int n_experts, int n_tokens, int n_expert_used, int nchannels_y, int si1, int sis1, cudaStream_t stream);
+        int n_experts, int n_tokens, int n_expert_used, int nchannels_y, int si1, int sis1, bool write_inverse, cudaStream_t stream);

ggml/src/ggml-cuda/mmq-load-tiles.cuh

@@ -39,29 +39,37 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
         }
 
         const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbx;
-        const int qs_offset = 4*kqsx;
-        const int qs0 = bxi->qs[qs_offset + 0] | (bxi->qs[qs_offset + 1] << 8) |
-                        (bxi->qs[qs_offset + 2] << 16) | (bxi->qs[qs_offset + 3] << 24);
-
-        int unpacked_bytes[8];
-#pragma unroll
-        for (int j = 0; j < 8; ++j) {
-            const int shift = j * 4;
-            const int bits4 = (qs0 >> shift) & 0x0F;
-            const int b0 = (bits4 & 0x01) ? 1 : -1;
-            const int b1 = (bits4 & 0x02) ? 1 : -1;
-            const int b2 = (bits4 & 0x04) ? 1 : -1;
-            const int b3 = (bits4 & 0x08) ? 1 : -1;
-            unpacked_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24);
-        }
+        const int16_t    * qxi = (const int16_t *) bxi->qs + kqsx * 2;
 
         const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0;
 #pragma unroll
-        for (int j = 0; j < 8; ++j) {
+        for (int j = 0; j < 2; ++j) {
+            const int q  = qxi[j];
+
+            // unpack crumbs into nibble indices
+            const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8,  9, 12, 13]
+            const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); // [2, 3, 6, 7] [10, 11, 14, 15]
+            // unpack nibbles into byte values
+            const int s0 = __byte_perm(0x01FF, 0x01FF, n0 >>  0);
+            const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >>  0);
+            const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16);
+            const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16);
+            // unshuffle values
+            const int v0 = __byte_perm(s0, s1, 0x5410);
+            const int v1 = __byte_perm(s0, s1, 0x7632);
+            const int v2 = __byte_perm(s2, s3, 0x5410);
+            const int v3 = __byte_perm(s2, s3, 0x7632);
+
 #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
-            x_qs[i*sram_stride           + dst_offset + j] = unpacked_bytes[j];
+            x_qs[i*sram_stride           + dst_offset + j*4+0] = v0;
+            x_qs[i*sram_stride           + dst_offset + j*4+1] = v1;
+            x_qs[i*sram_stride           + dst_offset + j*4+2] = v2;
+            x_qs[i*sram_stride           + dst_offset + j*4+3] = v3;
 #else
-            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j] = unpacked_bytes[j];
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+0] = v0;
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+1] = v1;
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+2] = v2;
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+3] = v3;
 #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
         }
     }

ggml/src/ggml-cuda/mmq.cu

@@ -122,11 +122,12 @@ void ggml_cuda_mul_mat_q(
 
     const bool fallback = ne01 % 128 != 0;
 
-    // TODO: tighter pool buffer size vs q8 path
     const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4);
+    const size_t y_block_size       = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq);
+    const size_t y_values_per_block = use_native_fp4 ? QK_FP4_MMQ            : QK8_1_MMQ;
 
     if (!ids) {
-        const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1 +
+        const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * y_block_size/y_values_per_block +
             ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq);
         ggml_cuda_pool_alloc<char> src1_q8_1(ctx.pool(), nbytes_src1_q8_1);
 
@@ -148,7 +149,7 @@ void ggml_cuda_mul_mat_q(
 
         // Stride depends on quantization format
         const int64_t s12 = use_native_fp4 ?
-                                ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) :  // block_fp4_mmq holds 256 values
+                                ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_FP4_MMQ * sizeof(int)) :
                                 ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int));
         const int64_t s13 = ne12*s12;
 
@@ -174,17 +175,21 @@ void ggml_cuda_mul_mat_q(
     ggml_cuda_pool_alloc<int32_t> ids_dst(ctx.pool(), ne_get_rows);
     ggml_cuda_pool_alloc<int32_t> expert_bounds(ctx.pool(), ne02 + 1);
 
+    // gate/up activations are broadcast across experts (ne11 == 1): quantize each token once and
+    // scatter to its slots. ids_src1 then holds the inverse map (token slot -> compact row).
+    const bool dedup_bcast = ne11 == 1 && n_expert_used > 1;
+
     {
         GGML_ASSERT(ids->nb[0] == ggml_element_size(ids));
         const int si1  = ids->nb[1] / ggml_element_size(ids);
         const int sis1 = nb12 / nb11;
 
         ggml_cuda_launch_mm_ids_helper((const int32_t *) ids->data, ids_src1.get(), ids_dst.get(), expert_bounds.get(),
-            ne02, ne12, n_expert_used, ne11, si1, sis1, stream);
+            ne02, ne12, n_expert_used, ne11, si1, sis1, /*write_inverse =*/ dedup_bcast, stream);
         CUDA_CHECK(cudaGetLastError());
     }
 
-    const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * sizeof(block_q8_1)/QK8_1 +
+    const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * y_block_size/y_values_per_block +
         ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq);
     ggml_cuda_pool_alloc<char> src1_q8_1(ctx.pool(), nbytes_src1_q8_1);
 
@@ -197,7 +202,16 @@ void ggml_cuda_mul_mat_q(
         const int64_t s12 = src1->nb[2] / ts_src1;
         const int64_t s13 = src1->nb[3] / ts_src1;
 
-        if (use_native_fp4) {
+        if (dedup_bcast) {
+            // quantize each token once, scatter its block to all n_expert_used slots
+            if (use_native_fp4) {
+                quantize_scatter_mmq_fp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10,
+                                        /*stride_token=*/s12, ne10_padded, ne12, ne11_flat, n_expert_used, stream);
+            } else {
+                quantize_scatter_mmq_q8_1_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10,
+                                        /*stride_token=*/s12, ne10_padded, ne12, ne11_flat, n_expert_used, stream);
+            }
+        } else if (use_native_fp4) {
             quantize_mmq_fp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13,
                                     ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream);
         } else {
@@ -207,8 +221,8 @@ void ggml_cuda_mul_mat_q(
         CUDA_CHECK(cudaGetLastError());
     }
 
-    static_assert(QK_K == 8 * QK_MXFP4, "QK_K needs to be 8 * QK_MXFP4");
-    const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) :
+    static_assert(QK_FP4_MMQ == 8 * QK_MXFP4, "QK_FP4_MMQ needs to be 8 * QK_MXFP4");
+    const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_FP4_MMQ * sizeof(int)) :
                                          ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int));
     const int64_t s13 = ne12*s12;
 

ggml/src/ggml-cuda/mmq.cuh

@@ -21,6 +21,9 @@ enum mmq_q8_1_ds_layout {
     MMQ_Q8_1_DS_LAYOUT_D2S6,
 };
 
+static constexpr int QK8_1_MMQ  = 4*QK8_1;
+static constexpr int QK_FP4_MMQ = 2*QK8_1_MMQ;
+
 struct block_q8_1_mmq {
     // The y float data is converted to a data layout that can simply be copied to shared memory as a contiguous block.
     // The y float data is first grouped as blocks of 128 values.
@@ -39,7 +42,7 @@ struct block_q8_1_mmq {
         half  d2s6[8];  // 1 16 bit scale per 64 values + 1 16 bit partial sum per 16 values for the first 96 values,
                         //     stored as d0,d1,s1,s2,s3,s4,s5
     };
-    int8_t qs[4*QK8_1]; // 128 values quantized to 8 bit each
+    int8_t qs[QK8_1_MMQ];
 };
 
 // this struct is used for fp4 data types (currently only used for Blackwell)
@@ -47,10 +50,10 @@ struct block_q8_1_mmq {
 // nvfp4 has block size 16, each int32 of d4 contains 4 ue4m3 scales
 struct block_fp4_mmq {
     uint32_t d4[4];
-    int8_t   qs[4 * 32];  // 256 FP4 values packed as 4-bit pairs (2 per byte)
+    int8_t   qs[QK_FP4_MMQ / 2];
 };
 
-static_assert(sizeof(block_q8_1_mmq) == 4*QK8_1 + 4*sizeof(half2), "Unexpected block_q8_1_mmq size");
+static_assert(sizeof(block_q8_1_mmq) == QK8_1_MMQ + 4*sizeof(half2), "Unexpected block_q8_1_mmq size");
 static_assert(sizeof(block_q8_1_mmq) == 4*sizeof(block_q8_1),      "Unexpected block_q8_1_mmq size");
 static_assert(sizeof(block_fp4_mmq)  == sizeof(block_q8_1_mmq),    "Unexpected block_fp4_mmq size");
 
@@ -833,9 +836,9 @@ static __device__ __forceinline__ void mul_mat_q_process_tile(
 
 #if defined(BLACKWELL_MMA_AVAILABLE)
     // FP4 tile stores 8 blocks
-    constexpr int ne_block = (type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4) ? QK_K : 4 * QK8_1;
+    constexpr int ne_block = (type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4) ? QK_FP4_MMQ : QK8_1_MMQ;
 #else
-    constexpr int ne_block = 4 * QK8_1;
+    constexpr int ne_block = QK8_1_MMQ;
 #endif  // defined(BLACKWELL_MMA_AVAILABLE)
 
     constexpr int ITER_K          = ggml_cuda_mmq_get_K_vram(type, J, fallback);

ggml/src/ggml-cuda/quantize.cu

@@ -75,10 +75,12 @@ __device__ __forceinline__ uint8_t compute_e8m0_scale(float amax) {
 }
 
 
+// scatter: grid over tokens, quantize once, write to all the token's compact rows
+template <bool scatter>
 static __global__ void quantize_mmq_nvfp4(
         const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy,
         const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
-        const int64_t ne0, const int64_t ne1, const int64_t ne2) {
+        const int64_t ne0, const int64_t ne1, const int64_t ne2, const int n_expert_used) {
 #if defined(BLACKWELL_MMA_AVAILABLE)
 
     const int64_t i0_base = ((int64_t) blockDim.x * blockIdx.y + threadIdx.x) * QK_NVFP4_SUB;
@@ -86,25 +88,25 @@ static __global__ void quantize_mmq_nvfp4(
         return;
     }
 
-    const int64_t i1 = blockIdx.x;
-    const int64_t i2 = blockIdx.z % ne2;
-    const int64_t i3 = blockIdx.z / ne2;
-    const int64_t i01 = ids ? ids[i1] : i1;
-    const int64_t k_block = i0_base / QK_K;
-    const int64_t blocks_per_col = (ne0 + QK_K - 1) / QK_K;
+    const int64_t k_block = i0_base / QK_FP4_MMQ;
+    const int64_t blocks_per_col = (ne0 + QK_FP4_MMQ - 1) / QK_FP4_MMQ;
     if (k_block >= blocks_per_col) {
         return;
     }
+    const int sub = (i0_base % QK_FP4_MMQ) / QK_NVFP4_SUB;
 
-    const int64_t ib = blockIdx.z * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x;
-    block_fp4_mmq * y = (block_fp4_mmq *) vy;
-    block_fp4_mmq * yb = y + ib;
-
-    const int sub = (i0_base % QK_K) / QK_NVFP4_SUB;
+    int64_t base_idx;
+    if constexpr (scatter) {
+        base_idx = (int64_t) blockIdx.x * s02; // one physical row per token
+    } else {
+        const int64_t i2  = blockIdx.z % ne2;
+        const int64_t i3  = blockIdx.z / ne2;
+        const int64_t i01 = ids ? ids[blockIdx.x] : blockIdx.x;
+        base_idx = i3 * s03 + i2 * s02 + i01 * s01;
+    }
 
     float vals_raw[QK_NVFP4_SUB];
     float amax_raw = 0.0f;
-    const int64_t base_idx = i3 * s03 + i2 * s02 + i01 * s01;
 #pragma unroll
     for (int k = 0; k < QK_NVFP4_SUB; k++) {
         const int64_t i00 = i0_base + k;
@@ -160,11 +162,27 @@ static __global__ void quantize_mmq_nvfp4(
         q1 |= (uint32_t) ggml_cuda_float_to_fp4_e2m1(vals_raw[k + 12], inv_scale) << (8 * k + 4);
     }
 
-    uint32_t * yqs = reinterpret_cast<uint32_t *>(yb->qs);
-    yqs[2 * sub + 0] = q0;
-    yqs[2 * sub + 1] = q1;
-    reinterpret_cast<uint8_t *>(yb->d4)[sub] = fp8_code;
+    block_fp4_mmq * y = (block_fp4_mmq *) vy;
+    if constexpr (scatter) {
+#pragma unroll
+        for (int slot = 0; slot < n_expert_used; ++slot) {
+            const int64_t i = ids[(int64_t) blockIdx.x * n_expert_used + slot];
+            block_fp4_mmq * yb = y + (k_block * ne1 + i);
+            uint32_t * yqs = reinterpret_cast<uint32_t *>(yb->qs);
+            yqs[2 * sub + 0] = q0;
+            yqs[2 * sub + 1] = q1;
+            reinterpret_cast<uint8_t *>(yb->d4)[sub] = fp8_code;
+        }
+    } else {
+        block_fp4_mmq * yb = y + (blockIdx.z * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x);
+        uint32_t * yqs = reinterpret_cast<uint32_t *>(yb->qs);
+        yqs[2 * sub + 0] = q0;
+        yqs[2 * sub + 1] = q1;
+        reinterpret_cast<uint8_t *>(yb->d4)[sub] = fp8_code;
+    }
+    GGML_UNUSED(n_expert_used);
 #else
+    GGML_UNUSED(n_expert_used);
     NO_DEVICE_CODE; // This is for Blackwell NVFP4 activations only.
 #endif // defined(BLACKWELL_MMA_AVAILABLE)
 
@@ -172,6 +190,8 @@ static __global__ void quantize_mmq_nvfp4(
 
 // quantize values in the format mxfp4 is stored which is interleaved nibbles
 // i.e. a block a0-a31 is represented as a0a16,a1a17 ...a15a31
+// scatter: grid over tokens, quantize once, write to all the token's compact rows
+template <bool scatter>
 static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
                                           const int32_t * __restrict__ ids,
                                           void * __restrict__ vy,
@@ -181,7 +201,8 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
                                           const int64_t s03,
                                           const int64_t ne0,
                                           const int     ne1,
-                                          const int     ne2) {
+                                          const int     ne2,
+                                          const int     n_expert_used) {
     constexpr int vals_per_scale = 32;
     constexpr int vals_per_warp  = 2 * vals_per_scale;  // Each warp processes 2 blocks of 32 = 64 values
 
@@ -196,30 +217,27 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
         return;
     }
 
-    const int64_t i1 = blockIdx.x;
-    const int64_t i2 = blockIdx.z % ne2;
-    const int64_t i3 = blockIdx.z / ne2;
-
-    ggml_cuda_pdl_sync();
-    const int64_t i01 = ids ? ids[i1] : i1;
-    const int64_t i02 = i2;
-    const int64_t i03 = i3;
-
-    block_fp4_mmq * y = (block_fp4_mmq *) vy;
-
-    const int64_t block_fp4_mmq_size = 8 * QK_MXFP4;  // 256 values
-    const int64_t ib0                = blockIdx.z * ((int64_t) ne1 * (ne0 / block_fp4_mmq_size));
-    const int64_t ib = ib0 + (warp_start_offset / block_fp4_mmq_size) * ne1 + blockIdx.x;
+    const int64_t block_fp4_mmq_size = QK_FP4_MMQ;
+    const int64_t k_block            = warp_start_offset / block_fp4_mmq_size;
     const int64_t quad_idx_in_block  = (warp_start_offset % block_fp4_mmq_size) / vals_per_warp;
 
     const int group_id = lane_id_32 / 4;
     const int lane_in_group = lane_id_32 % 4;
     const int base = group_id * 2;
-    char2 * yqs2 = (char2 *) y[ib].qs;
 
-    const int64_t base_pos = i03 * s03 + i02 * s02 + i01 * s01;
+    ggml_cuda_pdl_sync();
+    int64_t base_pos;
+    if constexpr (scatter) {
+        base_pos = (int64_t) blockIdx.x * s02; // one physical row per token
+    } else {
+        const int64_t i2  = blockIdx.z % ne2;
+        const int64_t i3  = blockIdx.z / ne2;
+        const int64_t i01 = ids ? ids[blockIdx.x] : blockIdx.x;
+        base_pos = i3 * s03 + i2 * s02 + i01 * s01;
+    }
 
     uint8_t scales[2];
+    char2   packed[2];
 
 #pragma unroll
     for (int b = 0; b < 2; ++b) {
@@ -244,11 +262,8 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
         const float val2 = __shfl_sync(0xFFFFFFFF, scaled_val, base + 1, WARP_SIZE);
         const float val3 = __shfl_sync(0xFFFFFFFF, scaled_val, base + 17, WARP_SIZE);
 
-        if (lane_in_group == 0) {
-            __nv_fp4x4_e2m1 fp4_packed(make_float4(val0, val1, val2, val3));
-
-            yqs2[quad_idx_in_block * 16 + b * 8 + group_id] = *(char2 *) &fp4_packed;
-        }
+        __nv_fp4x4_e2m1 fp4_packed(make_float4(val0, val1, val2, val3));
+        packed[b] = *(char2 *) &fp4_packed;
 #else
         // Fallback: manual FP4 conversion using LUT
         const uint8_t q_val = ggml_cuda_float_to_fp4_e2m1(xi, inv_s);
@@ -258,26 +273,49 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x,
         const uint8_t q_hi_0 = __shfl_sync(0xFFFFFFFF, q_val, base + 16, WARP_SIZE);
         const uint8_t q_hi_1 = __shfl_sync(0xFFFFFFFF, q_val, base + 17, WARP_SIZE);
 
-        if (lane_in_group == 0) {
-            char2 q;
-            q.x = (q_hi_0 << 4) | q_lo_0;
-            q.y = (q_hi_1 << 4) | q_lo_1;
-            yqs2[quad_idx_in_block * 16 + b * 8 + group_id] = q;
-        }
+        char2 q;
+        q.x = (q_hi_0 << 4) | q_lo_0;
+        q.y = (q_hi_1 << 4) | q_lo_1;
+        packed[b] = q;
 #endif // CUDART_VERSION >= 12080
     }
 
-    if (lane_id_32 == 0) {
-        // Store 2 scales packed into 1 uint32
-        y[ib].d4[quad_idx_in_block] = (scales[1] << 8) | scales[0];
+    block_fp4_mmq * y = (block_fp4_mmq *) vy;
+    if constexpr (scatter) {
+#pragma unroll
+        for (int slot = 0; slot < n_expert_used; ++slot) {
+            const int64_t i = ids[(int64_t) blockIdx.x * n_expert_used + slot];
+            block_fp4_mmq * yb = y + (k_block * ne1 + i);
+            char2 * yqs2 = (char2 *) yb->qs;
+            if (lane_in_group == 0) {
+                yqs2[quad_idx_in_block * 16 + 0 * 8 + group_id] = packed[0];
+                yqs2[quad_idx_in_block * 16 + 1 * 8 + group_id] = packed[1];
+            }
+            if (lane_id_32 == 0) {
+                yb->d4[quad_idx_in_block] = (scales[1] << 8) | scales[0];
+            }
+        }
+    } else {
+        const int64_t ib0 = blockIdx.z * ((int64_t) ne1 * (ne0 / block_fp4_mmq_size));
+        block_fp4_mmq * yb = y + (ib0 + k_block * ne1 + blockIdx.x);
+        char2 * yqs2 = (char2 *) yb->qs;
+        if (lane_in_group == 0) {
+            yqs2[quad_idx_in_block * 16 + 0 * 8 + group_id] = packed[0];
+            yqs2[quad_idx_in_block * 16 + 1 * 8 + group_id] = packed[1];
+        }
+        if (lane_id_32 == 0) {
+            yb->d4[quad_idx_in_block] = (scales[1] << 8) | scales[0];
+        }
     }
+    GGML_UNUSED(n_expert_used);
 }
 
-template <mmq_q8_1_ds_layout ds_layout>
+// scatter: grid over tokens, quantize once, write to all the token's compact rows
+template <mmq_q8_1_ds_layout ds_layout, bool scatter>
 static __global__ void quantize_mmq_q8_1(
         const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy,
         const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
-        const int64_t ne0, const int ne1, const int ne2) {
+        const int64_t ne0, const int ne1, const int ne2, const int n_expert_used) {
 
     constexpr int vals_per_scale = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 64 : 32;
     constexpr int vals_per_sum   = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 16 : 32;
@@ -288,26 +326,27 @@ static __global__ void quantize_mmq_q8_1(
         return;
     }
 
-    const int64_t i1 = blockIdx.x;
-    const int64_t i2 = blockIdx.z % ne2;
-    const int64_t i3 = blockIdx.z / ne2;
-
     const int64_t i00 = i0;
     ggml_cuda_pdl_sync();
-    const int64_t i01 = ids ? ids[i1] : i1;
-    const int64_t i02 = i2;
-    const int64_t i03 = i3;
+
+    int64_t base_idx;
+    if constexpr (scatter) {
+        base_idx = (int64_t) blockIdx.x * s02; // one physical row per token
+    } else {
+        const int64_t i2  = blockIdx.z % ne2;
+        const int64_t i3  = blockIdx.z / ne2;
+        const int64_t i01 = ids ? ids[blockIdx.x] : blockIdx.x;
+        base_idx = i3*s03 + i2*s02 + i01*s01;
+    }
 
     const float4 * x4 = (const float4 *) x;
-
     block_q8_1_mmq * y = (block_q8_1_mmq *) vy;
 
-    const int64_t ib0 = blockIdx.z*((int64_t)gridDim.x*gridDim.y*blockDim.x/QK8_1); // first block of channel
-    const int64_t ib  = ib0 + (i0 / (4*QK8_1))*ne1 + blockIdx.x;                    // block index in channel
-    const int64_t iqs = i0 % (4*QK8_1);                                             // quant index in block
+    const int64_t k_block = i0 / QK8_1_MMQ; // column block in the channel
+    const int64_t iqs     = i0 % QK8_1_MMQ; // quant index in block
 
     // Load 4 floats per thread and calculate max. abs. value between them:
-    const float4 xi = i0 < ne00 ? x4[(i03*s03 + i02*s02 + i01*s01 + i00)/4] : make_float4(0.0f, 0.0f, 0.0f, 0.0f);
+    const float4 xi = i0 < ne00 ? x4[(base_idx + i00)/4] : make_float4(0.0f, 0.0f, 0.0f, 0.0f);
     float amax = fabsf(xi.x);
     amax = fmaxf(amax, fabsf(xi.y));
     amax = fmaxf(amax, fabsf(xi.z));
@@ -336,40 +375,41 @@ static __global__ void quantize_mmq_q8_1(
     q.y = roundf(xi.y*d_inv);
     q.z = roundf(xi.z*d_inv);
     q.w = roundf(xi.w*d_inv);
-
-    // Write back 4 int8 values as a single 32 bit value for better memory bandwidth:
-    char4 * yqs4 = (char4 *) y[ib].qs;
-    yqs4[iqs/4] = q;
-
-    if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6) {
-        if (iqs % 16 != 0 || iqs >= 96) {
-            return;
-        }
-
-        y[ib].d2s6[2 + iqs/16] = sum;
-
-        if (iqs % 64 != 0) {
-            return;
-        }
-
-        const float d = 1.0f / d_inv;
-
-        y[ib].d2s6[iqs/64] = d;
-
-        return;
-    }
-
-    if (iqs % 32 != 0) {
-        return;
-    }
-
     const float d = 1.0f / d_inv;
 
-    if (ds_layout == MMQ_Q8_1_DS_LAYOUT_DS4) {
-        y[ib].ds4[iqs/32] = make_half2(d, sum);
-    } else {
-        y[ib].d4[iqs/32]  = d;
+    // write the block once (normal) or to each of the token's compact rows (scatter)
+    const int nwrite = scatter ? n_expert_used : 1;
+#pragma unroll
+    for (int slot = 0; slot < nwrite; ++slot) {
+        int64_t ib;
+        if constexpr (scatter) {
+            const int64_t i = ids[(int64_t) blockIdx.x * n_expert_used + slot];
+            ib = k_block*ne1 + i;
+        } else {
+            const int64_t ib0 = blockIdx.z*((int64_t)gridDim.x*gridDim.y*blockDim.x/QK8_1); // first block of channel
+            ib = ib0 + k_block*ne1 + blockIdx.x;
+        }
+
+        // Write back 4 int8 values as a single 32 bit value for better memory bandwidth:
+        char4 * yqs4 = (char4 *) y[ib].qs;
+        yqs4[iqs/4] = q;
+
+        if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6) {
+            if (iqs % 16 == 0 && iqs < 96) {
+                y[ib].d2s6[2 + iqs/16] = sum;
+                if (iqs % 64 == 0) {
+                    y[ib].d2s6[iqs/64] = d;
+                }
+            }
+        } else if (iqs % 32 == 0) {
+            if (ds_layout == MMQ_Q8_1_DS_LAYOUT_DS4) {
+                y[ib].ds4[iqs/32] = make_half2(d, sum);
+            } else {
+                y[ib].d4[iqs/32]  = d;
+            }
+        }
     }
+    GGML_UNUSED(n_expert_used);
 }
 
 void quantize_row_q8_1_cuda(
@@ -394,7 +434,7 @@ void quantize_mmq_q8_1_cuda(
         const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
         const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) {
     GGML_ASSERT(ne00 % 4 == 0);
-    GGML_ASSERT(ne0 % (4*QK8_1) == 0);
+    GGML_ASSERT(ne0 % QK8_1_MMQ == 0);
 
     // ne1 tends to assume the highest values, therefore use it as the "x" dimension of the CUDA grid:
     const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ);
@@ -402,16 +442,16 @@ void quantize_mmq_q8_1_cuda(
     const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1);
     switch (mmq_get_q8_1_ds_layout(type_src0)) {
         case MMQ_Q8_1_DS_LAYOUT_D4:
-            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4>
-                <<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
+            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4, false>
+                <<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
             break;
         case MMQ_Q8_1_DS_LAYOUT_DS4:
-            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4>
-                <<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
+            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4, false>
+                <<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
             break;
         case MMQ_Q8_1_DS_LAYOUT_D2S6:
-            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D2S6>
-                <<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
+            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D2S6, false>
+                <<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
             break;
         default:
             GGML_ABORT("fatal error");
@@ -419,6 +459,62 @@ void quantize_mmq_q8_1_cuda(
     }
 }
 
+// scatter=true reuses the quant kernel: grid over tokens, ids = inverse map (token slot -> compact row)
+void quantize_scatter_mmq_q8_1_cuda(
+        const float * x, const int32_t * ids_src1_inv, void * vy, const ggml_type type_src0,
+        const int64_t ne00, const int64_t stride_token, const int64_t ne0,
+        const int64_t n_tokens, const int64_t nrows_dst, const int n_expert_used, cudaStream_t stream) {
+    GGML_ASSERT(ne00 % 4 == 0);
+    GGML_ASSERT(ne0 % QK8_1_MMQ == 0);
+
+    const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ);
+    const dim3 num_blocks(n_tokens, block_num_y, 1);
+    const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1);
+    switch (mmq_get_q8_1_ds_layout(type_src0)) {
+        case MMQ_Q8_1_DS_LAYOUT_D4:
+            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4, true><<<num_blocks, block_size, 0, stream>>>(
+                x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);
+            break;
+        case MMQ_Q8_1_DS_LAYOUT_DS4:
+            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4, true><<<num_blocks, block_size, 0, stream>>>(
+                x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);
+            break;
+        case MMQ_Q8_1_DS_LAYOUT_D2S6:
+            quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D2S6, true><<<num_blocks, block_size, 0, stream>>>(
+                x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);
+            break;
+        default:
+            GGML_ABORT("fatal error");
+            break;
+    }
+}
+
+// scatter=true reuses the quant kernels: grid over tokens, ids = inverse map (token slot -> compact row)
+void quantize_scatter_mmq_fp4_cuda(
+        const float * x, const int32_t * ids_src1_inv, void * vy, const ggml_type type_src0,
+        const int64_t ne00, const int64_t stride_token, const int64_t ne0,
+        const int64_t n_tokens, const int64_t nrows_dst, const int n_expert_used, cudaStream_t stream) {
+    GGML_ASSERT(ne0 > 0);
+    if (type_src0 == GGML_TYPE_NVFP4) {
+        GGML_ASSERT(ne00 % QK_NVFP4 == 0);
+        constexpr int nvfp4_block_size = 128;
+        const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size);
+        const dim3 block_size(nvfp4_block_size, 1, 1);
+        const dim3 num_blocks(n_tokens, block_num_y, 1);
+        quantize_mmq_nvfp4<true><<<num_blocks, block_size, 0, stream>>>(
+            x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/nrows_dst, /*ne2=*/1, n_expert_used);
+    } else {
+        GGML_ASSERT(type_src0 == GGML_TYPE_MXFP4);
+        constexpr int nwarps = 8;
+        constexpr int vals_per_block = nwarps * 2 * QK_MXFP4;
+        const int64_t block_num_y = (ne0 + vals_per_block - 1) / vals_per_block;
+        const dim3 block_size(WARP_SIZE, nwarps, 1);
+        const dim3 num_blocks(n_tokens, block_num_y, 1);
+        quantize_mmq_mxfp4<true><<<num_blocks, block_size, 0, stream>>>(
+            x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);
+    }
+}
+
 void quantize_mmq_fp4_cuda(
         const float * x, const int32_t * ids, void * vy, const ggml_type type_src0,
         const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
@@ -432,8 +528,8 @@ void quantize_mmq_fp4_cuda(
         const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size);
         const dim3 block_size(nvfp4_block_size, 1, 1);
         const dim3 num_blocks(ne1, block_num_y, ne2 * ne3);
-        quantize_mmq_nvfp4<<<num_blocks, block_size, 0, stream>>>(
-            x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
+        quantize_mmq_nvfp4<false><<<num_blocks, block_size, 0, stream>>>(
+            x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
     } else {
         GGML_ASSERT(ne0 % (2 * QK_MXFP4) == 0);
 
@@ -445,6 +541,6 @@ void quantize_mmq_fp4_cuda(
         const dim3    num_blocks(ne1, block_num_y, ne2 * ne3);
         const dim3    block_size(WARP_SIZE, nwarps, 1);
 
-        quantize_mmq_mxfp4<<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
+        quantize_mmq_mxfp4<false><<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
     }
 }

ggml/src/ggml-cuda/quantize.cuh

@@ -39,3 +39,28 @@ void quantize_mmq_fp4_cuda(const float *   x,
                              int64_t         ne2,
                              int64_t         ne3,
                              cudaStream_t    stream);
+
+// quantize each token once and scatter the block to its compact rows (via the inverse map)
+void quantize_scatter_mmq_fp4_cuda(const float *   x,
+                                   const int32_t * ids_src1_inv,
+                                   void *          vy,
+                                   ggml_type       type_src0,
+                                   int64_t         ne00,
+                                   int64_t         stride_token,
+                                   int64_t         ne0,
+                                   int64_t         n_tokens,
+                                   int64_t         nrows_dst,
+                                   int             n_expert_used,
+                                   cudaStream_t    stream);
+
+void quantize_scatter_mmq_q8_1_cuda(const float *   x,
+                                    const int32_t * ids_src1_inv,
+                                    void *          vy,
+                                    ggml_type       type_src0,
+                                    int64_t         ne00,
+                                    int64_t         stride_token,
+                                    int64_t         ne0,
+                                    int64_t         n_tokens,
+                                    int64_t         nrows_dst,
+                                    int             n_expert_used,
+                                    cudaStream_t    stream);

ggml/src/ggml-cuda/vecdotq.cuh

@@ -681,35 +681,40 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1(
     // Q8_1: 32 elements per block with individual scales
     // iqs selects which of the 4 chunks of 32 elements to process (0-3)
 
-    const float d1 = bq1_0->d;
+    const float     d1 = bq1_0->d;
+    const int16_t * qs = (const int16_t *) bq1_0->qs + iqs * 2;
 
     // Process only the chunk specified by iqs
     const block_q8_1 * bq8_1_chunk = bq8_1 + iqs;
 
-    // Load 32 bits (4 bytes) for this chunk from Q1_0
-    const int offset = iqs * 4;
-    const int v = bq1_0->qs[offset + 0] | (bq1_0->qs[offset + 1] << 8) |
-                  (bq1_0->qs[offset + 2] << 16) | (bq1_0->qs[offset + 3] << 24);
-
-    // Unpack 32 bits into 32 signed values (-1 or +1)
-    int vi_bytes[8];
-#pragma unroll
-    for (int j = 0; j < 8; ++j) {
-        const int shift = j * 4;
-        const int bits4 = (v >> shift) & 0x0F;
-        const int b0 = (bits4 & 0x01) ? 1 : -1;
-        const int b1 = (bits4 & 0x02) ? 1 : -1;
-        const int b2 = (bits4 & 0x04) ? 1 : -1;
-        const int b3 = (bits4 & 0x08) ? 1 : -1;
-        vi_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24);
-    }
-
-    // Compute dot product for this 32-element chunk
     int sumi = 0;
 #pragma unroll
-    for (int j = 0; j < 8; ++j) {
-        const int u = get_int_b4(bq8_1_chunk->qs, j);
-        sumi = ggml_cuda_dp4a(vi_bytes[j], u, sumi);
+    for (int j = 0; j < 2; ++j) {
+        const int q  = qs[j];
+
+        const int u0 = get_int_b4(bq8_1_chunk->qs, j*4+0);
+        const int u1 = get_int_b4(bq8_1_chunk->qs, j*4+1);
+        const int u2 = get_int_b4(bq8_1_chunk->qs, j*4+2);
+        const int u3 = get_int_b4(bq8_1_chunk->qs, j*4+3);
+
+        // unpack crumbs into nibble indices
+        const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); // [0, 1, 4, 5] [ 8,  9, 12, 13]
+        const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); // [2, 3, 6, 7] [10, 11, 14, 15]
+        // unpack nibbles into byte values
+        const int s0 = __byte_perm(0x01FF, 0x01FF, n0 >>  0);
+        const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >>  0);
+        const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16);
+        const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16);
+        // unshuffle values
+        const int v0 = __byte_perm(s0, s1, 0x5410);
+        const int v1 = __byte_perm(s0, s1, 0x7632);
+        const int v2 = __byte_perm(s2, s3, 0x5410);
+        const int v3 = __byte_perm(s2, s3, 0x7632);
+
+        sumi = ggml_cuda_dp4a(v0, u0, sumi);
+        sumi = ggml_cuda_dp4a(v1, u1, sumi);
+        sumi = ggml_cuda_dp4a(v2, u2, sumi);
+        sumi = ggml_cuda_dp4a(v3, u3, sumi);
     }
 
     // Apply Q1_0's single scale and this chunk's Q8_1 scale

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -114,6 +114,7 @@ enum GPU_FAMILY {
 
 enum ADRENO_GPU_GEN {
     ADRENO_UNKNOWN,
+    A6X,
     A7X,
     A8X,
     X1E,
@@ -122,6 +123,7 @@ enum ADRENO_GPU_GEN {
 
 enum ADRENO_CL_COMPILER_TYPE {
     E031,
+    E17,
     DX,
 };
 
@@ -243,6 +245,19 @@ static ggml_cl_version get_opencl_c_version(ggml_cl_version platform_version, cl
 }
 
 static ADRENO_GPU_GEN get_adreno_gpu_gen(const char *device_name) {
+    if (strstr(device_name, "610") || strstr(device_name, "612") ||
+        strstr(device_name, "613") || strstr(device_name, "615") ||
+        strstr(device_name, "616") || strstr(device_name, "618") ||
+        strstr(device_name, "619") || strstr(device_name, "620") ||
+        strstr(device_name, "630") || strstr(device_name, "640") ||
+        strstr(device_name, "642") || strstr(device_name, "643") ||
+        strstr(device_name, "644") || strstr(device_name, "650") ||
+        strstr(device_name, "660") || strstr(device_name, "663") ||
+        strstr(device_name, "680") || strstr(device_name, "685") ||
+        strstr(device_name, "690")) {
+        return ADRENO_GPU_GEN::A6X;
+    }
+
     if (strstr(device_name, "730") ||
         strstr(device_name, "740") ||
         strstr(device_name, "750")) {
@@ -250,7 +265,8 @@ static ADRENO_GPU_GEN get_adreno_gpu_gen(const char *device_name) {
     }
 
     if (strstr(device_name, "830") ||
-        strstr(device_name, "840")) {
+        strstr(device_name, "840") ||
+        strstr(device_name, "850")) {
         return ADRENO_GPU_GEN::A8X;
     }
 
@@ -274,6 +290,17 @@ static ggml_cl_compiler_version get_adreno_cl_compiler_version(const char *drive
     size_t compiler_minor_offset = 8;
     size_t compiler_patch_offset = 11;
 
+    if (compiler_ver_pos == std::string::npos) {
+        compiler_ver_pos = driver_ver_str.find("E17");
+        if (compiler_ver_pos != std::string::npos) {
+            type = ADRENO_CL_COMPILER_TYPE::E17;
+            compiler_ver_len = 12;
+            compiler_major_offset = 4;
+            compiler_minor_offset = 7;
+            compiler_patch_offset = 10;
+        }
+    }
+
     if (compiler_ver_pos == std::string::npos) {
         compiler_ver_pos = driver_ver_str.find("DX");
         if (compiler_ver_pos == std::string::npos) {
@@ -282,6 +309,8 @@ static ggml_cl_compiler_version get_adreno_cl_compiler_version(const char *drive
         type = ADRENO_CL_COMPILER_TYPE::DX;
         compiler_ver_len = 11;
         compiler_major_offset = 3;
+        compiler_minor_offset = 6;
+        compiler_patch_offset = 9;
     }
 
     std::string compiler_ver_str = driver_ver_str.substr(compiler_ver_pos, compiler_ver_len);
@@ -1641,6 +1670,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
     // those compiler versions since it is anyway not used for Adreno.
     if (backend_ctx->gpu_family != ADRENO ||
         backend_ctx->adreno_cl_compiler_version.newer_than_or_same(E031, 38, 11, 0) ||
+        backend_ctx->adreno_cl_compiler_version.type == E17 ||
         backend_ctx->adreno_cl_compiler_version.type == DX) {
 #ifdef GGML_OPENCL_EMBED_KERNELS
         const std::string kernel_src {
@@ -6932,8 +6962,31 @@ inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, c
     return threashold_ok;
 }
 
+static bool adreno_e17_compiler_quirks(const ggml_backend_opencl_context *backend_ctx) {
+    if (!backend_ctx || backend_ctx->gpu_family != GPU_FAMILY::ADRENO ||
+        backend_ctx->adreno_cl_compiler_version.type != ADRENO_CL_COMPILER_TYPE::E17) {
+        return false;
+    }
+    const char * env = getenv("GGML_OPENCL_ADRENO_E17_QUIRKS");
+    return !(env && env[0] == '0');
+}
+
 inline bool use_adreno_moe_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
-    GGML_UNUSED(backend_ctx);
+    // The moe weight repack kernels *_trans4_ns alias a private ushort8 through a uchar*.
+    // Certain compilers (found with some A7x and A6x) miscompiles this, corrupting the weights.
+    // So, exclude A6x and A7x from using Adreno MoE kernels for now.
+    // The quants that have a general mul_mat_id kernel fallback to the general version; the
+    // rest fallback to CPU.
+    if (backend_ctx && (backend_ctx->adreno_gen == ADRENO_GPU_GEN::A6X ||
+                        backend_ctx->adreno_gen == ADRENO_GPU_GEN::A7X ||
+                        backend_ctx->adreno_gen == ADRENO_GPU_GEN::ADRENO_UNKNOWN)) {
+        return false;
+    }
+
+    if (adreno_e17_compiler_quirks(backend_ctx)) {
+        return false;
+    }
+
     int ne01 = tensor->ne[1];
     return (((strstr(tensor->name, "ffn") != NULL) && (strstr(tensor->name, "exps") != NULL)) || (strstr(tensor->name, "as") != NULL)) && (ne01 % 32 == 0);
 }
@@ -7266,6 +7319,10 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
         case GGML_OP_MEAN:
             return op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_FLASH_ATTN_EXT: {
+            // The E17 compilers segfault while building FA kernels, skip E17 for now
+            if (adreno_e17_compiler_quirks(backend_ctx)) {
+                return false;
+            }
             const ggml_tensor * q = op->src[0];
             const ggml_tensor * k = op->src[1];
             const ggml_tensor * v = op->src[2];
@@ -7319,6 +7376,14 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
                 return false;
             }
 
+            // Some compilers for A7x (Adreno 740, compiler E031.41) crashes when
+            // building FA kernels with mixed or quant types (f32_f16, f32_q8_0, f32_q4_0)
+            // Here we skip all A7x for these kernels to avoid crash
+            if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::A7X &&
+                (is_f32_f16 || is_f32_q8_0 || is_f32_q4_0)) {
+                return false;
+            }
+
             if (dk == 512) {
                 if (backend_ctx->gpu_family == INTEL) {
                     return false;

ggml/src/ggml-opencl/kernels/flash_attn_f32_f16.cl

@@ -30,6 +30,10 @@
 #elif defined(cl_qcom_subgroup_shuffle)
 #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
 #define HAS_SUBGROUP_SHUFFLE 1
+// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR
+// name, so calling it is an implicit declaration and the program fails to build.
+// Route it to the qcom builtin.
+#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f)
 #endif
 
 #define ACC_TYPE float

ggml/src/ggml-opencl/kernels/flash_attn_f32_q4_0.cl

@@ -10,6 +10,10 @@
 #elif defined(cl_qcom_subgroup_shuffle)
 #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
 #define HAS_SUBGROUP_SHUFFLE 1
+// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR
+// name, so calling it is an implicit declaration and the program fails to build.
+// Route it to the qcom builtin.
+#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f)
 #endif
 
 // Flash attention: Q=f32, K=q4_0, V=q4_0.

ggml/src/ggml-opencl/kernels/flash_attn_f32_q8_0.cl

@@ -10,6 +10,10 @@
 #elif defined(cl_qcom_subgroup_shuffle)
 #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
 #define HAS_SUBGROUP_SHUFFLE 1
+// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR
+// name, so calling it is an implicit declaration and the program fails to build.
+// Route it to the qcom builtin.
+#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f)
 #endif
 
 // Flash attention: Q=f32, K=q8_0, V=q8_0.

ggml/src/ggml-opencl/kernels/gemm_moe_mxfp4_f32_ns.cl

@@ -274,8 +274,9 @@ kernel void kernel_gemm_moe_mxfp4_f32_ns(
         shared_b[b_local_offset.y] = bx8_f16.hi;
 
         // Dequantization
-        reg_a.lo = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.lo)) * s;
-        reg_a.hi = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.hi)) * s;
+        // Cast the e8m0 scale to half to satisfy E17 compilers
+        reg_a.lo = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.lo)) * (half)s;
+        reg_a.hi = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.hi)) * (half)s;
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 
@@ -304,8 +305,9 @@ kernel void kernel_gemm_moe_mxfp4_f32_ns(
         shared_b[b_local_offset.y] = bx8_f16.hi;
 
         // Dequantization
-        reg_a.lo = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.lo)) * s;
-        reg_a.hi = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.hi)) * s;
+        // Cast the e8m0 scale to half to satisfy E17 compilers
+        reg_a.lo = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.lo)) * (half)s;
+        reg_a.hi = mxfp4_to_fp16_packed8(as_ushort2(mxfp4x16.hi)) * (half)s;
 
         sub_group_barrier(CLK_LOCAL_MEM_FENCE);
 

ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl

@@ -24,6 +24,10 @@
 #elif defined(cl_qcom_subgroup_shuffle)
 #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
 #define HAS_SUBGROUP_SHUFFLE 1
+// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR
+// name, so calling it is an implicit declaration and the program fails to build.
+// Route it to the qcom builtin.
+#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f)
 #endif
 
 // Assumes row size (ne00) is a multiple of 4

ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl

@@ -1,3 +1,5 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
 #ifdef cl_intel_required_subgroup_size
 #pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
 #define INTEL_GPU 1

src/llama-quant.cpp

@@ -673,7 +673,7 @@ static ggml_type llama_tensor_get_type(quantize_state_impl & qs, const llama_mod
     ggml_type new_type = default_type;
 
     // get more optimal quantization type based on the tensor shape, layer, etc.
-    if (!params->pure && ggml_is_quantized(default_type)) {
+    if (ggml_is_quantized(default_type)) {
         // if the user provided tensor types - use those
         bool manual = false;
         if (!qs.tensor_type_patterns.empty()) {
@@ -692,7 +692,7 @@ static ggml_type llama_tensor_get_type(quantize_state_impl & qs, const llama_mod
         }
 
         // if not manual - use the standard logic for choosing the quantization type based on the selected mixture
-        if (!manual) {
+        if (!manual && !params->pure) {
             new_type = llama_tensor_get_type_impl(qs, new_type, tensor, params->ftype, tm.category);
         }
 

tools/server/server-http.cpp

@@ -283,9 +283,9 @@ bool server_http_context::init(const common_params & params) {
         } else if (params.cors_origins == "localhost") {
             // special case: only reflect the Origin header if it is a localhost origin
             std::string origin = req.get_header_value("Origin");
-            if (origin_is_localhost(origin)) {
+            if (!origin.empty() && origin_is_localhost(origin)) {
                 res.set_header("Access-Control-Allow-Origin", origin);
-            } else {
+            } else if (!origin.empty()) {
                 SRV_WRN("(CORS) skip non-localhost origin: %s\n", origin.c_str());
             }
         } else {

tools/server/server-tools.cpp

@@ -289,7 +289,7 @@ struct server_tool_read_file : server_tool {
             {"function", {
                 {"name", name},
                 {"description", "Read the contents of a file. Optionally specify a 1-based line range. "
-                                "If append_loc is true, each line is prefixed with its line number (e.g. \"1\u2192 ...\")."},
+                                "If append_loc is true, each line is prefixed with its line number (e.g. \"1\u2192...\")."},
                 {"parameters", {
                     {"type", "object"},
                     {"properties", {
@@ -339,7 +339,7 @@ struct server_tool_read_file : server_tool {
 
             std::string out_line;
             if (append_loc) {
-                out_line = std::to_string(lineno) + "\u2192 " + line + "\n";
+                out_line = std::to_string(lineno) + "\u2192" + line + "\n";
             } else {
                 out_line = line + "\n";
             }

tools/server/tests/unit/test_speculative.py

@@ -12,8 +12,9 @@ def create_server():
     server = ServerPreset.stories15m_moe()
     # set default values
     server.model_draft = download_file(MODEL_DRAFT_FILE_URL)
-    server.draft_min = 4
-    server.draft_max = 8
+    server.spec_type = "draft-simple"
+    server.spec_draft_n_min = 4
+    server.spec_draft_n_max = 8
     server.fa = "off"
 
 
@@ -25,6 +26,7 @@ def fixture_create_server():
 def test_with_and_without_draft():
     global server
     server.model_draft = None  # disable draft model
+    server.spec_type = None
     server.start()
     res = server.make_request("POST", "/completion", data={
         "prompt": "I believe the meaning of life is",
@@ -46,6 +48,7 @@ def test_with_and_without_draft():
         "n_predict": 16,
     })
     assert res.status_code == 200
+    assert res.body["timings"]["draft_n"] > 0
     content_draft = res.body["content"]
 
     assert content_no_draft == content_draft
@@ -63,8 +66,8 @@ def test_different_draft_min_draft_max():
     last_content = None
     for draft_min, draft_max in test_values:
         server.stop()
-        server.draft_min = draft_min
-        server.draft_max = draft_max
+        server.spec_draft_n_min = draft_min
+        server.spec_draft_n_max = draft_max
         server.start()
         res = server.make_request("POST", "/completion", data={
             "prompt": "I believe the meaning of life is",

tools/server/tests/utils.py

@@ -95,6 +95,7 @@ class ServerProcess:
     no_models_autoload: bool | None = None
     lora_files: List[str] | None = None
     enable_ctx_shift: int | None = False
+    spec_type: str | None = None
     spec_draft_n_min: int | None = None
     spec_draft_n_max: int | None = None
     no_ui: bool | None = None
@@ -226,6 +227,8 @@ class ServerProcess:
                 server_args.extend(["--lora", lora_file])
         if self.enable_ctx_shift:
             server_args.append("--context-shift")
+        if self.spec_type:
+            server_args.extend(["--spec-type", self.spec_type])
         if self.api_key:
             server_args.extend(["--api-key", self.api_key])
         if self.spec_draft_n_max:

tools/tokenize/tokenize.cpp

@@ -103,19 +103,11 @@ int main(int argc, char ** argv) {
         return 1;
     }
 
-    // which prompt source was requested?
-    // -p/--prompt and -f/--file both end up in params.prompt (common's -f also
-    // strips a single trailing newline), but -f additionally records the path
-    // in params.prompt_file, so we use that to tell them apart.
+    // -f and -p both land in params.prompt; -f also sets prompt_file. -f and -p
+    // resolve like the other tools (no mutual exclusion), --stdin takes precedence.
     const bool use_stdin = params.tokenize_stdin;
     const bool use_file  = !params.prompt_file.empty();
 
-    // sanity check: --stdin is mutually exclusive with -f/--file and -p/--prompt
-    if (use_stdin && (use_file || !params.prompt.empty())) {
-        LOG_ERR("error: --stdin is mutually exclusive with --file and --prompt\n");
-        return 1;
-    }
-
     // must have some prompt
     if (!use_stdin && !use_file && params.prompt.empty()) {
         LOG_ERR("error: must specify one of: --stdin, --file or --prompt\n");

tools/ui/eslint.config.js

@@ -29,6 +29,13 @@ export default ts.config(
 			// This app uses hash-based routing (#/) where resolve() from $app/paths does not apply
 			'svelte/no-navigation-without-resolve': 'off',
 
+			// Snippet bodies often ignore one or more of the parent's params
+			// (e.g. `{#snippet children(_meta, ctx)}` when only ctx is read).
+			'@typescript-eslint/no-unused-vars': [
+				'error',
+				{ argsIgnorePattern: '^_', varsIgnorePattern: '^_' }
+			],
+
 			// Enforce empty line at end of file
 			'eol-last': 'error'
 		}

tools/ui/src/app.css

@@ -193,6 +193,33 @@
 		-ms-overflow-style: none;
 		scrollbar-width: none;
 	}
+
+	.shimmer-text {
+		background: linear-gradient(
+			90deg,
+			var(--muted-foreground),
+			var(--foreground),
+			var(--muted-foreground)
+		);
+		background-size: 200% 100%;
+		background-clip: text;
+		-webkit-background-clip: text;
+		-webkit-text-fill-color: transparent;
+		font-weight: 500;
+		animation: shimmer 1s linear infinite;
+	}
+
+	@keyframes shimmer {
+		to {
+			background-position: -200% 0;
+		}
+	}
+
+	@media (prefers-reduced-motion: reduce) {
+		.shimmer-text {
+			animation: none;
+		}
+	}
 }
 
 .mermaidTooltip {

tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Copy } from '@lucide/svelte';
 	import { copyToClipboard } from '$lib/utils';
 	import ActionIcon from './ActionIcon.svelte';
@@ -11,7 +12,7 @@
 <ActionIcon
 	icon={Copy}
 	tooltip={ariaLabel}
-	iconSize="h-4 w-4"
+	iconSize={ICON_CLASS_DEFAULT}
 	disabled={!canCopy}
 	onclick={() => canCopy && copyToClipboard(text)}
 />

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { X, Music, Video } from '@lucide/svelte';
 	import {
 		formatFileSize,
@@ -109,9 +110,9 @@
 		class="flex h-8 w-8 items-center justify-center rounded bg-primary/10 text-xs font-medium text-primary"
 	>
 		{#if isAudio}
-			<Music class="h-4 w-4 text-white/70" />
+			<Music class="{ICON_CLASS_DEFAULT} text-white/70" />
 		{:else if isVideo}
-			<Video class="h-4 w-4 text-white/70" />
+			<Video class="{ICON_CLASS_DEFAULT} text-white/70" />
 		{:else}
 			{fileTypeLabel}
 		{/if}

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import type { ChatAttachmentDisplayItem } from '$lib/types';
 	import { FileText, Eye, Info } from '@lucide/svelte';
 	import { Button } from '$lib/components/ui/button';
@@ -88,7 +89,7 @@
 		onclick={() => (pdfViewMode = PdfViewMode.TEXT)}
 		disabled={pdfImagesLoading}
 	>
-		<FileText class="mr-1 h-4 w-4" />
+		<FileText class="mr-1 {ICON_CLASS_DEFAULT}" />
 		Text
 	</Button>
 
@@ -100,10 +101,10 @@
 	>
 		{#if pdfImagesLoading}
 			<div
-				class="mr-1 h-4 w-4 animate-spin rounded-full border-2 border-current border-t-transparent"
+				class="mr-1 {ICON_CLASS_DEFAULT} animate-spin rounded-full border-2 border-current border-t-transparent"
 			></div>
 		{:else}
-			<Eye class="mr-1 h-4 w-4" />
+			<Eye class="mr-1 {ICON_CLASS_DEFAULT}" />
 		{/if}
 		Pages
 	</Button>
@@ -111,7 +112,7 @@
 
 {#if !hasVisionModality && activeModelId && currentItem}
 	<Alert.Root class="mb-4 max-w-4xl">
-		<Info class="h-4 w-4" />
+		<Info class={ICON_CLASS_DEFAULT} />
 		<Alert.Title>Preview only</Alert.Title>
 		<Alert.Description>
 			<span class="inline-flex">

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Music, Video, FileText } from '@lucide/svelte';
 	import { HorizontalScrollCarousel } from '$lib/components/app/misc';
 
@@ -49,11 +50,11 @@
 							class="bg-foreground-muted/50 flex h-12 w-12 flex-col items-center justify-center gap-0.5 py-1"
 						>
 							{#if item.isAudio}
-								<Music class="h-4 w-4 text-white/70" />
+								<Music class="{ICON_CLASS_DEFAULT} text-white/70" />
 							{:else if item.isVideo}
-								<Video class="h-4 w-4 text-white/70" />
+								<Video class="{ICON_CLASS_DEFAULT} text-white/70" />
 							{:else}
-								<FileText class="h-4 w-4 text-white/70" />
+								<FileText class="{ICON_CLASS_DEFAULT} text-white/70" />
 							{/if}
 
 							<span class="font-mono text-[9px] text-white/60">{getFileExtension(item.name)}</span>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Plus } from '@lucide/svelte';
 	import { Button } from '$lib/components/ui/button';
 	import * as Tooltip from '$lib/components/ui/tooltip';
@@ -23,7 +24,7 @@
 		>
 			<span class="sr-only">{ATTACHMENT_TOOLTIP_TEXT}</span>
 
-			<Plus class="h-4 w-4" />
+			<Plus class={ICON_CLASS_DEFAULT} />
 		</Button>
 	</Tooltip.Trigger>
 

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Plus, File, MessageSquare, Zap, FolderOpen } from '@lucide/svelte';
 	import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
 	import * as Tooltip from '$lib/components/ui/tooltip';
@@ -83,7 +84,7 @@
 					>
 						<span class="sr-only">{ATTACHMENT_TOOLTIP_TEXT}</span>
 
-						<Plus class="h-4 w-4" />
+						<Plus class={ICON_CLASS_DEFAULT} />
 					</DropdownMenu.Trigger>
 				{/snippet}
 			</Tooltip.Trigger>
@@ -100,7 +101,7 @@
 
 			<DropdownMenu.Sub>
 				<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
-					<File class="h-4 w-4" />
+					<File class={ICON_CLASS_DEFAULT} />
 
 					<span>Add files</span>
 				</DropdownMenu.SubTrigger>
@@ -113,7 +114,7 @@
 								class="{item.class ?? ''} flex cursor-pointer items-center gap-2"
 								onclick={() => attachmentMenu.callbacks[item.action]()}
 							>
-								<item.icon class="h-4 w-4" />
+								<item.icon class={ICON_CLASS_DEFAULT} />
 
 								<span>{item.label}</span>
 							</DropdownMenu.Item>
@@ -126,7 +127,7 @@
 												class="{item.class ?? ''} flex items-center gap-2"
 												disabled
 											>
-												<item.icon class="h-4 w-4" />
+												<item.icon class={ICON_CLASS_DEFAULT} />
 
 												<span>{item.label}</span>
 											</DropdownMenu.Item>
@@ -147,7 +148,7 @@
 				class="flex cursor-pointer items-center gap-2"
 				onclick={onSystemPromptClick}
 			>
-				<MessageSquare class="h-4 w-4" />
+				<MessageSquare class={ICON_CLASS_DEFAULT} />
 
 				<span>System Message</span>
 			</DropdownMenu.Item>
@@ -163,7 +164,7 @@
 					class="flex cursor-pointer items-center gap-2"
 					onclick={onMcpPromptClick}
 				>
-					<Zap class="h-4 w-4" />
+					<Zap class={ICON_CLASS_DEFAULT} />
 
 					<span>MCP Prompt</span>
 				</DropdownMenu.Item>
@@ -174,7 +175,7 @@
 					class="flex cursor-pointer items-center gap-2"
 					onclick={onMcpResourcesClick}
 				>
-					<FolderOpen class="h-4 w-4" />
+					<FolderOpen class={ICON_CLASS_DEFAULT} />
 
 					<span>MCP Resources</span>
 				</DropdownMenu.Item>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Settings, Plus } from '@lucide/svelte';
 	import { Switch } from '$lib/components/ui/switch';
 	import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
@@ -60,7 +61,7 @@
 <DropdownMenu.Root>
 	<DropdownMenu.Sub onOpenChange={handleMcpSubMenuOpen}>
 		<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
-			<McpLogo class="h-4 w-4" />
+			<McpLogo class={ICON_CLASS_DEFAULT} />
 
 			<span>MCP Servers</span>
 		</DropdownMenu.SubTrigger>
@@ -92,7 +93,7 @@
 										<McpServerIdentity
 											{displayName}
 											{faviconUrl}
-											iconClass="h-4 w-4"
+											iconClass={ICON_CLASS_DEFAULT}
 											iconRounded="rounded-sm"
 											showVersion={false}
 											nameClass="text-sm"
@@ -123,7 +124,7 @@
 							class="flex cursor-pointer items-center gap-2"
 							onclick={handleMcpSettingsClick}
 						>
-							<Settings class="h-4 w-4" />
+							<Settings class={ICON_CLASS_DEFAULT} />
 
 							<span>Manage MCP Servers</span>
 						</DropdownMenu.Item>
@@ -140,7 +141,7 @@
 					class="flex cursor-pointer items-center gap-2"
 					onclick={handleMcpSettingsClick}
 				>
-					<Plus class="h-4 w-4" />
+					<Plus class={ICON_CLASS_DEFAULT} />
 
 					<span>Add MCP Servers</span>
 				</DropdownMenu.Item>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Lightbulb, LightbulbOff, Check, Info } from '@lucide/svelte';
 	import * as DropdownMenu from '$lib/components/ui/dropdown-menu';
 	import * as Tooltip from '$lib/components/ui/tooltip';
@@ -13,9 +14,9 @@
 	<DropdownMenu.Sub bind:open={subOpen}>
 		<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
 			{#if reasoning.thinkingEnabled}
-				<Lightbulb class="h-4 w-4 shrink-0 text-amber-400" />
+				<Lightbulb class="{ICON_CLASS_DEFAULT} shrink-0 text-amber-400" />
 			{:else}
-				<LightbulbOff class="h-4 w-4 shrink-0 text-muted-foreground" />
+				<LightbulbOff class="{ICON_CLASS_DEFAULT} shrink-0 text-muted-foreground" />
 			{/if}
 
 			<span
@@ -46,9 +47,9 @@
 					}}
 				>
 					{#if reasoning.isSelected(level)}
-						<Check class="h-4 w-4 shrink-0 text-foreground" />
+						<Check class="{ICON_CLASS_DEFAULT} shrink-0 text-foreground" />
 					{:else}
-						<div class="h-4 w-4 shrink-0"></div>
+						<div class="{ICON_CLASS_DEFAULT} shrink-0"></div>
 					{/if}
 
 					<span class="flex-1">{level.label}</span>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import type { Snippet } from 'svelte';
 	import * as Tooltip from '$lib/components/ui/tooltip';
 	import * as Sheet from '$lib/components/ui/sheet';
@@ -108,15 +109,15 @@
 					>
 						<Collapsible.Trigger class={sheetItemClass}>
 							{#if reasoningExpanded}
-								<ChevronDown class="h-4 w-4 shrink-0" />
+								<ChevronDown class="{ICON_CLASS_DEFAULT} shrink-0" />
 							{:else}
-								<ChevronRight class="h-4 w-4 shrink-0" />
+								<ChevronRight class="{ICON_CLASS_DEFAULT} shrink-0" />
 							{/if}
 
 							{#if reasoning.thinkingEnabled}
-								<Lightbulb class="h-4 w-4 shrink-0 text-amber-400" />
+								<Lightbulb class="{ICON_CLASS_DEFAULT} shrink-0 text-amber-400" />
 							{:else}
-								<LightbulbOff class="h-4 w-4 shrink-0 text-muted-foreground" />
+								<LightbulbOff class="{ICON_CLASS_DEFAULT} shrink-0 text-muted-foreground" />
 							{/if}
 
 							<span class="flex-1">Reasoning</span>
@@ -138,9 +139,9 @@
 									>
 										<div class="flex min-w-0 items-center gap-3">
 											{#if reasoning.isSelected(level)}
-												<Check class="h-4 w-4 shrink-0 text-foreground" />
+												<Check class="{ICON_CLASS_DEFAULT} shrink-0 text-foreground" />
 											{:else}
-												<div class="h-4 w-4 shrink-0"></div>
+												<div class="{ICON_CLASS_DEFAULT} shrink-0"></div>
 											{/if}
 
 											<span class="text-sm">{level.label}</span>
@@ -161,12 +162,12 @@
 				<Collapsible.Root open={filesExpanded} onOpenChange={(open) => (filesExpanded = open)}>
 					<Collapsible.Trigger class={sheetItemClass}>
 						{#if filesExpanded}
-							<ChevronDown class="h-4 w-4 shrink-0" />
+							<ChevronDown class="{ICON_CLASS_DEFAULT} shrink-0" />
 						{:else}
-							<ChevronRight class="h-4 w-4 shrink-0" />
+							<ChevronRight class="{ICON_CLASS_DEFAULT} shrink-0" />
 						{/if}
 
-						<File class="h-4 w-4 shrink-0" />
+						<File class="{ICON_CLASS_DEFAULT} shrink-0" />
 
 						<span class="flex-1">Add files</span>
 					</Collapsible.Trigger>
@@ -181,7 +182,7 @@
 										class={sheetItemClass}
 										onclick={() => attachmentMenu.callbacks[item.action]()}
 									>
-										<item.icon class="h-4 w-4 shrink-0" />
+										<item.icon class="{ICON_CLASS_DEFAULT} shrink-0" />
 
 										<span>{item.label}</span>
 									</button>
@@ -189,7 +190,7 @@
 									<Tooltip.Root delayDuration={TOOLTIP_DELAY_DURATION}>
 										<Tooltip.Trigger>
 											<button type="button" class={sheetItemClass} disabled>
-												<item.icon class="h-4 w-4 shrink-0" />
+												<item.icon class="{ICON_CLASS_DEFAULT} shrink-0" />
 
 												<span>{item.label}</span>
 											</button>
@@ -208,12 +209,12 @@
 				<Collapsible.Root open={mcpExpanded} onOpenChange={(open) => (mcpExpanded = open)}>
 					<Collapsible.Trigger class={sheetItemClass}>
 						{#if mcpExpanded}
-							<ChevronDown class="h-4 w-4 shrink-0" />
+							<ChevronDown class="{ICON_CLASS_DEFAULT} shrink-0" />
 						{:else}
-							<ChevronRight class="h-4 w-4 shrink-0" />
+							<ChevronRight class="{ICON_CLASS_DEFAULT} shrink-0" />
 						{/if}
 
-						<McpLogo class="inline h-4 w-4 shrink-0" />
+						<McpLogo class="inline {ICON_CLASS_DEFAULT} shrink-0" />
 
 						<span class="flex-1">MCP Servers</span>
 
@@ -242,7 +243,7 @@
 											<img
 												src={faviconUrl}
 												alt=""
-												class="h-4 w-4 shrink-0 rounded-sm"
+												class="{ICON_CLASS_DEFAULT} shrink-0 rounded-sm"
 												onerror={(e) => {
 													(e.currentTarget as HTMLImageElement).style.display = 'none';
 												}}
@@ -280,12 +281,12 @@
 					<Collapsible.Root open={toolsExpanded} onOpenChange={(open) => (toolsExpanded = open)}>
 						<Collapsible.Trigger class={sheetItemClass}>
 							{#if toolsExpanded}
-								<ChevronDown class="h-4 w-4 shrink-0" />
+								<ChevronDown class="{ICON_CLASS_DEFAULT} shrink-0" />
 							{:else}
-								<ChevronRight class="h-4 w-4 shrink-0" />
+								<ChevronRight class="{ICON_CLASS_DEFAULT} shrink-0" />
 							{/if}
 
-							<PencilRuler class="inline h-4 w-4 shrink-0" />
+							<PencilRuler class="inline {ICON_CLASS_DEFAULT} shrink-0" />
 
 							<span class="flex-1">Tools</span>
 
@@ -310,7 +311,7 @@
 											<img
 												src={favicon}
 												alt=""
-												class="h-4 w-4 shrink-0 rounded-sm"
+												class="{ICON_CLASS_DEFAULT} shrink-0 rounded-sm"
 												onerror={(e) => {
 													(e.currentTarget as HTMLImageElement).style.display = 'none';
 												}}
@@ -325,7 +326,7 @@
 
 										<Checkbox
 											{checked}
-											class="h-4 w-4 shrink-0"
+											class="{ICON_CLASS_DEFAULT} shrink-0"
 											onclick={(e) => e.stopPropagation()}
 											onCheckedChange={() => toolsPanel.toggleGroupByLabel(group.label)}
 										/>
@@ -341,7 +342,7 @@
 					class={sheetItemClass}
 					onclick={() => attachmentMenu.callbacks[AttachmentAction.SYSTEM_PROMPT_CLICK]()}
 				>
-					<MessageSquare class="h-4 w-4 shrink-0" />
+					<MessageSquare class="{ICON_CLASS_DEFAULT} shrink-0" />
 
 					<span>System Message</span>
 				</button>
@@ -352,7 +353,7 @@
 						class={sheetItemClass}
 						onclick={() => attachmentMenu.callbacks[AttachmentAction.MCP_PROMPT_CLICK]()}
 					>
-						<Zap class="h-4 w-4 shrink-0" />
+						<Zap class="{ICON_CLASS_DEFAULT} shrink-0" />
 
 						<span>MCP Prompt</span>
 					</button>
@@ -364,7 +365,7 @@
 						class={sheetItemClass}
 						onclick={() => attachmentMenu.callbacks[AttachmentAction.MCP_RESOURCES_CLICK]()}
 					>
-						<FolderOpen class="h-4 w-4 shrink-0" />
+						<FolderOpen class="{ICON_CLASS_DEFAULT} shrink-0" />
 
 						<span>MCP Resources</span>
 					</button>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { PencilRuler, ChevronDown, ChevronRight, Loader2, Info, Check } from '@lucide/svelte';
 	import { Checkbox } from '$lib/components/ui/checkbox';
 	import * as Collapsible from '$lib/components/ui/collapsible';
@@ -15,7 +16,7 @@
 
 <DropdownMenu.Sub onOpenChange={(open) => open && toolsPanel.handleOpen()}>
 	<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
-		<PencilRuler class="h-4 w-4" />
+		<PencilRuler class={ICON_CLASS_DEFAULT} />
 
 		<span>Tools</span>
 	</DropdownMenu.SubTrigger>
@@ -24,14 +25,14 @@
 		{#if toolsPanel.totalToolCount === 0}
 			{#if toolsStore.loading}
 				<div class="px-3 py-4 text-center text-sm text-muted-foreground">
-					<Loader2 class="mx-auto mb-1 h-4 w-4 animate-spin" />
+					<Loader2 class="mx-auto mb-1 {ICON_CLASS_DEFAULT} animate-spin" />
 
 					Loading tools...
 				</div>
 			{:else if toolsStore.isToolsEndpointUnreachable}
 				<div class="grid gap-2.5 px-3 py-4 text-sm text-muted-foreground">
 					<span class="flex gap-2">
-						<Info class="mt-0.5 h-4 w-4 shrink-0" />
+						<Info class="mt-0.5 {ICON_CLASS_DEFAULT} shrink-0" />
 
 						<span>
 							Run llama-server with <code>{CLI_FLAGS.TOOLS}</code> flag to enable
@@ -41,7 +42,7 @@
 					</span>
 
 					<span class="flex gap-2">
-						<Info class="mt-0.5 h-4 w-4 shrink-0" />
+						<Info class="mt-0.5 {ICON_CLASS_DEFAULT} shrink-0" />
 
 						<span>
 							{hasMcpServersAvailable ? 'Enable' : 'Add'} MCP Server(s) to access
@@ -54,7 +55,7 @@
 				<div class="px-3 py-4 text-center text-sm text-muted-foreground">Failed to load tools</div>
 			{:else if toolsPanel.noToolsInfoMessage}
 				<div class="flex gap-2 px-3 py-4 text-sm text-muted-foreground">
-					<Info class="mt-0.5 h-4 w-4 shrink-0" />
+					<Info class="mt-0.5 {ICON_CLASS_DEFAULT} shrink-0" />
 
 					<span>{toolsPanel.noToolsInfoMessage}</span>
 				</div>
@@ -87,7 +88,7 @@
 										<img
 											src={favicon}
 											alt=""
-											class="h-4 w-4 shrink-0 rounded-sm"
+											class="{ICON_CLASS_DEFAULT} shrink-0 rounded-sm"
 											onerror={(e) => {
 												(e.currentTarget as HTMLImageElement).style.display = 'none';
 											}}
@@ -109,7 +110,7 @@
 											{...props}
 											{checked}
 											onCheckedChange={() => toolsPanel.toggleGroupByLabel(group.label)}
-											class="mr-2 h-4 w-4 shrink-0"
+											class="mr-2 {ICON_CLASS_DEFAULT} shrink-0"
 										/>
 									{/snippet}
 								</Tooltip.Trigger>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Mic, Square } from '@lucide/svelte';
 	import { Button } from '$lib/components/ui/button';
 	import * as Tooltip from '$lib/components/ui/tooltip';
@@ -36,9 +37,9 @@
 				<span class="sr-only">{isRecording ? 'Stop recording' : 'Start recording'}</span>
 
 				{#if isRecording}
-					<Square class="h-4 w-4 animate-pulse fill-white" />
+					<Square class="{ICON_CLASS_DEFAULT} animate-pulse fill-white" />
 				{:else}
-					<Mic class="h-4 w-4" />
+					<Mic class={ICON_CLASS_DEFAULT} />
 				{/if}
 			</Button>
 		</Tooltip.Trigger>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Square, SkipForward } from '@lucide/svelte';
 	import { Button } from '$lib/components/ui/button';
 	import { ChatService } from '$lib/services';
@@ -185,7 +186,9 @@
 		>
 			<span class="sr-only">Skip reasoning</span>
 
-			<SkipForward class="h-4 w-4 stroke-muted-foreground group-hover:stroke-foreground" />
+			<SkipForward
+				class="{ICON_CLASS_DEFAULT} stroke-muted-foreground group-hover:stroke-foreground"
+			/>
 		</Button>
 	{/if}
 

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

@@ -6,7 +6,7 @@
 	import { REASONING_EFFORT_TOKENS } from '$lib/constants/reasoning-effort-tokens';
 	import { REASONING_EFFORT_LEVELS } from '$lib/constants/reasoning-effort';
 	import type { ReasoningEffortLevel } from '$lib/types';
-	import { DIALOG_SUBMENU_CONTENT } from '$lib/constants/css-classes';
+	import { DIALOG_SUBMENU_CONTENT, ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import {
 		modelsStore,
 		checkModelSupportsThinking,
@@ -74,9 +74,9 @@
 	<DropdownMenu.Sub bind:open={subOpen}>
 		<DropdownMenu.SubTrigger class="flex cursor-pointer items-center gap-2">
 			{#if thinkingEnabled}
-				<Lightbulb class="h-4 w-4 shrink-0 text-amber-400" />
+				<Lightbulb class="{ICON_CLASS_DEFAULT} shrink-0 text-amber-400" />
 			{:else}
-				<LightbulbOff class="h-4 w-4 shrink-0 text-muted-foreground" />
+				<LightbulbOff class="{ICON_CLASS_DEFAULT} shrink-0 text-muted-foreground" />
 			{/if}
 
 			<span class="flex-1">Thinking</span>
@@ -118,7 +118,7 @@
 					{/if}
 
 					{#if isSelected(level)}
-						<Check class="h-4 w-4 shrink-0 text-foreground" />
+						<Check class="{ICON_CLASS_DEFAULT} shrink-0 text-foreground" />
 					{/if}
 				</button>
 			{/each}

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

@@ -384,7 +384,6 @@
 			{isLastAssistantMessage}
 			{message}
 			{toolMessages}
-			messageContent={message.content}
 			onConfirmDelete={handleConfirmDelete}
 			onContinue={handleContinue}
 			onCopy={handleCopy}

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

@@ -2,23 +2,20 @@
 	import {
 		ChatMessageAgenticContent,
 		ChatMessageActionIcons,
-		ChatMessageEditForm,
-		ChatMessageStatistics,
-		ModelBadge,
-		ModelsSelectorDropdown
+		ChatMessageAssistantModel,
+		ChatMessageAssistantProcessingInfo,
+		ChatMessageAssistantRawOutput,
+		ChatMessageAssistantStatistics,
+		ChatMessageEditForm
 	} from '$lib/components/app';
 	import { getMessageEditContext } from '$lib/contexts';
 	import { useProcessingState } from '$lib/hooks/use-processing-state.svelte';
 	import { isLoading, isChatStreaming } from '$lib/stores/chat.svelte';
-	import { copyToClipboard, deriveAgenticSections, modelLoadProgressText } from '$lib/utils';
-	import { AgenticSectionType, ChatMessageStatisticsMode } from '$lib/enums';
-	import { REASONING_TAGS } from '$lib/constants/agentic';
-	import { fade } from 'svelte/transition';
+	import { modelLoadProgressText } from '$lib/utils';
 	import { MessageRole } from '$lib/enums';
 	import { config } from '$lib/stores/settings.svelte';
 	import { isRouterMode } from '$lib/stores/server.svelte';
 	import { modelsStore } from '$lib/stores/models.svelte';
-	import { ServerModelStatus } from '$lib/enums';
 
 	import { hasAgenticContent } from '$lib/utils';
 
@@ -33,7 +30,6 @@
 		isLastAssistantMessage?: boolean;
 		message: DatabaseMessage;
 		toolMessages?: DatabaseMessage[];
-		messageContent: string | undefined;
 		onCopy: () => void;
 		onConfirmDelete: () => void;
 		onContinue?: () => void;
@@ -54,7 +50,6 @@
 		isLastAssistantMessage = false,
 		message,
 		toolMessages = [],
-		messageContent,
 		onConfirmDelete,
 		onContinue,
 		onCopy,
@@ -77,55 +72,11 @@
 
 	let currentConfig = $derived(config());
 	let isRouter = $derived(isRouterMode());
+
 	let showRawOutput = $state(false);
 
-	let rawOutputContent = $derived.by(() => {
-		const sections = deriveAgenticSections(message, toolMessages, [], false);
-		const parts: string[] = [];
-
-		for (const section of sections) {
-			switch (section.type) {
-				case AgenticSectionType.REASONING:
-				case AgenticSectionType.REASONING_PENDING:
-					parts.push(`${REASONING_TAGS.START}\n${section.content}\n${REASONING_TAGS.END}`);
-					break;
-
-				case AgenticSectionType.TEXT:
-					parts.push(section.content);
-					break;
-
-				case AgenticSectionType.TOOL_CALL:
-				case AgenticSectionType.TOOL_CALL_PENDING:
-				case AgenticSectionType.TOOL_CALL_STREAMING: {
-					const callObj: Record<string, unknown> = { name: section.toolName };
-
-					if (section.toolArgs) {
-						try {
-							callObj.arguments = JSON.parse(section.toolArgs);
-						} catch {
-							callObj.arguments = section.toolArgs;
-						}
-					}
-
-					parts.push(JSON.stringify(callObj, null, 2));
-
-					if (section.toolResult) {
-						parts.push(`[Tool Result]\n${section.toolResult}`);
-					}
-
-					break;
-				}
-			}
-		}
-
-		return parts.join('\n\n\n');
-	});
-
 	let displayedModel = $derived(message.model ?? null);
 
-	// model being switched to while it loads, so the selector bar tracks it
-	let pendingModel = $state<string | null>(null);
-
 	let isCurrentlyLoading = $derived(isLoading());
 	let isStreaming = $derived(isChatStreaming());
 	let hasNoContent = $derived(!message?.content?.trim());
@@ -189,10 +140,6 @@
 		};
 	});
 
-	function handleCopyModel() {
-		void copyToClipboard(displayedModel ?? '');
-	}
-
 	$effect(() => {
 		if (showProcessingInfoTop || showProcessingInfoBottom) {
 			processingState.startMonitoring();
@@ -211,23 +158,14 @@
 	aria-label="Assistant message with actions"
 >
 	{#if showProcessingInfoTop}
-		<div class="mt-6 w-full max-w-3xl" in:fade>
-			<div class="processing-container">
-				<span class="processing-text">
-					{modelLoadingText ??
-						processingState.getPromptProgressText() ??
-						processingState.getProcessingMessage() ??
-						'Processing...'}
-				</span>
-			</div>
-		</div>
+		<ChatMessageAssistantProcessingInfo {modelLoadingText} {processingState} position="top" />
 	{/if}
 
 	{#if editCtx.isEditing}
 		<ChatMessageEditForm />
-	{:else if message.role === MessageRole.ASSISTANT}
+	{:else}
 		{#if showRawOutput}
-			<pre class="raw-output">{rawOutputContent || ''}</pre>
+			<ChatMessageAssistantRawOutput {message} {toolMessages} />
 		{:else}
 			<ChatMessageAgenticContent
 				{message}
@@ -236,78 +174,28 @@
 				{isLastAssistantMessage}
 			/>
 		{/if}
-	{:else}
-		<div class="text-sm whitespace-pre-wrap">
-			{messageContent}
-		</div>
 	{/if}
 
 	{#if showProcessingInfoBottom}
-		<div class="mt-4 w-full max-w-3xl" in:fade>
-			<div class="processing-container">
-				<span class="processing-text">
-					{modelLoadingText ??
-						processingState.getPromptProgressText() ??
-						processingState.getProcessingMessage() ??
-						'Processing...'}
-				</span>
-			</div>
-		</div>
+		<ChatMessageAssistantProcessingInfo {modelLoadingText} {processingState} position="bottom" />
 	{/if}
 
 	<div class="info my-6 grid gap-4 tabular-nums">
 		{#if displayedModel}
 			<div class="inline-flex flex-wrap items-start gap-2 text-xs text-muted-foreground">
-				{#if isRouter}
-					<ModelsSelectorDropdown
-						currentModel={pendingModel ?? displayedModel}
-						disabled={isLoading()}
-						onModelChange={async (modelId: string, modelName: string) => {
-							const status = modelsStore.getModelStatus(modelId);
+				<ChatMessageAssistantModel
+					{displayedModel}
+					isLoading={isLoading()}
+					{isRouter}
+					{onRegenerate}
+				/>
 
-							if (status !== ServerModelStatus.LOADED) {
-								pendingModel = modelId;
-
-								try {
-									await modelsStore.loadModel(modelId);
-								} finally {
-									pendingModel = null;
-								}
-							}
-
-							onRegenerate(modelName);
-							return true;
-						}}
-					/>
-				{:else}
-					<ModelBadge model={displayedModel || undefined} onclick={handleCopyModel} />
-				{/if}
-
-				{#if currentConfig.showMessageStats && message.timings && message.timings.predicted_n && message.timings.predicted_ms}
-					{@const agentic = message.timings.agentic}
-					<ChatMessageStatistics
-						mode={ChatMessageStatisticsMode.GENERATION}
-						promptTokens={agentic ? agentic.llm.prompt_n : message.timings.prompt_n}
-						promptMs={agentic ? agentic.llm.prompt_ms : message.timings.prompt_ms}
-						predictedTokens={agentic ? agentic.llm.predicted_n : message.timings.predicted_n}
-						predictedMs={agentic ? agentic.llm.predicted_ms : message.timings.predicted_ms}
-						agenticTimings={agentic}
-					/>
-				{:else if isLoading() && currentConfig.showMessageStats}
-					{@const liveStats = processingState.getLiveProcessingStats()}
-					{@const genStats = processingState.getLiveGenerationStats()}
-
-					{#if genStats}
-						<ChatMessageStatistics
-							mode={ChatMessageStatisticsMode.GENERATION}
-							isLive
-							promptTokens={liveStats?.tokensProcessed}
-							promptMs={liveStats?.timeMs}
-							predictedTokens={genStats.tokensGenerated}
-							predictedMs={genStats.timeMs}
-						/>
-					{/if}
-				{/if}
+				<ChatMessageAssistantStatistics
+					{message}
+					isLoading={isLoading()}
+					{processingState}
+					showMessageStats={currentConfig.showMessageStats}
+				/>
 			</div>
 		{/if}
 	</div>
@@ -353,47 +241,4 @@
 			);
 		}
 	}
-
-	.processing-container {
-		display: flex;
-		flex-direction: column;
-		align-items: flex-start;
-		gap: 0.5rem;
-	}
-
-	.processing-text {
-		background: linear-gradient(
-			90deg,
-			var(--muted-foreground),
-			var(--foreground),
-			var(--muted-foreground)
-		);
-		background-size: 200% 100%;
-		background-clip: text;
-		-webkit-background-clip: text;
-		-webkit-text-fill-color: transparent;
-		animation: shine 1s linear infinite;
-		font-weight: 500;
-		font-size: 0.875rem;
-	}
-
-	@keyframes shine {
-		to {
-			background-position: -200% 0;
-		}
-	}
-
-	.raw-output {
-		width: 100%;
-		max-width: 48rem;
-		margin-top: 1.5rem;
-		padding: 1rem 1.25rem;
-		border-radius: 1rem;
-		background: hsl(var(--muted) / 0.3);
-		color: var(--foreground);
-		font-size: 0.875rem;
-		line-height: 1.6;
-		white-space: pre-wrap;
-		word-break: break-word;
-	}
 </style>

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

@@ -0,0 +1,46 @@
+<script lang="ts">
+	import { ModelBadge, ModelsSelectorDropdown } from '$lib/components/app';
+	import { copyToClipboard } from '$lib/utils';
+	import { modelsStore } from '$lib/stores/models.svelte';
+	import { ServerModelStatus } from '$lib/enums';
+
+	interface Props {
+		displayedModel: string | null;
+		isRouter: boolean;
+		isLoading: boolean;
+		onRegenerate: (modelOverride?: string) => void;
+	}
+
+	let { displayedModel, isRouter, isLoading, onRegenerate }: Props = $props();
+
+	let pendingModel = $state<string | null>(null);
+
+	function handleCopyModel() {
+		void copyToClipboard(displayedModel ?? '');
+	}
+</script>
+
+{#if isRouter}
+	<ModelsSelectorDropdown
+		currentModel={pendingModel ?? displayedModel}
+		disabled={isLoading}
+		onModelChange={async (modelId: string, modelName: string) => {
+			const status = modelsStore.getModelStatus(modelId);
+
+			if (status !== ServerModelStatus.LOADED) {
+				pendingModel = modelId;
+
+				try {
+					await modelsStore.loadModel(modelId);
+				} finally {
+					pendingModel = null;
+				}
+			}
+
+			onRegenerate(modelName);
+			return true;
+		}}
+	/>
+{:else}
+	<ModelBadge model={displayedModel || undefined} onclick={handleCopyModel} />
+{/if}

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

@@ -0,0 +1,25 @@
+<script lang="ts">
+	import { fade } from 'svelte/transition';
+	import type { UseProcessingStateReturn } from '$lib/hooks/use-processing-state.svelte';
+
+	interface Props {
+		modelLoadingText: string | null;
+		processingState: UseProcessingStateReturn;
+		position: 'top' | 'bottom';
+	}
+
+	let { modelLoadingText, processingState, position }: Props = $props();
+
+	const marginClass = position === 'top' ? 'mt-6' : 'mt-4';
+</script>
+
+<div class="{marginClass} w-full max-w-3xl" in:fade>
+	<div class="flex flex-col items-start gap-2">
+		<span class="shimmer-text text-sm">
+			{modelLoadingText ??
+				processingState.getPromptProgressText() ??
+				processingState.getProcessingMessage() ??
+				'Processing...'}
+		</span>
+	</div>
+</div>

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

@@ -0,0 +1,33 @@
+<script lang="ts">
+	import { deriveAgenticSections, buildAssistantRawOutput } from '$lib/utils';
+
+	interface Props {
+		message: DatabaseMessage;
+		toolMessages?: DatabaseMessage[];
+	}
+
+	let { message, toolMessages = [] }: Props = $props();
+
+	let rawOutputContent = $derived.by(() => {
+		const sections = deriveAgenticSections(message, toolMessages, [], false);
+		return buildAssistantRawOutput(sections);
+	});
+</script>
+
+<pre class="raw-output">{rawOutputContent || ''}</pre>
+
+<style>
+	.raw-output {
+		width: 100%;
+		max-width: 48rem;
+		margin-top: 1.5rem;
+		padding: 1rem 1.25rem;
+		border-radius: 1rem;
+		background: hsl(var(--muted) / 0.3);
+		color: var(--foreground);
+		font-size: 0.875rem;
+		line-height: 1.6;
+		white-space: pre-wrap;
+		word-break: break-word;
+	}
+</style>

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

@@ -0,0 +1,40 @@
+<script lang="ts">
+	import { ChatMessageStatistics } from '$lib/components/app';
+	import { ChatMessageStatisticsMode } from '$lib/enums';
+	import type { UseProcessingStateReturn } from '$lib/hooks/use-processing-state.svelte';
+
+	interface Props {
+		message: DatabaseMessage;
+		isLoading: boolean;
+		processingState: UseProcessingStateReturn;
+		showMessageStats: boolean;
+	}
+
+	let { message, isLoading, processingState, showMessageStats }: Props = $props();
+</script>
+
+{#if showMessageStats && message.timings && message.timings.predicted_n && message.timings.predicted_ms}
+	{@const agentic = message.timings.agentic}
+	<ChatMessageStatistics
+		mode={ChatMessageStatisticsMode.GENERATION}
+		promptTokens={agentic ? agentic.llm.prompt_n : message.timings.prompt_n}
+		promptMs={agentic ? agentic.llm.prompt_ms : message.timings.prompt_ms}
+		predictedTokens={agentic ? agentic.llm.predicted_n : message.timings.predicted_n}
+		predictedMs={agentic ? agentic.llm.predicted_ms : message.timings.predicted_ms}
+		agenticTimings={agentic}
+	/>
+{:else if isLoading && showMessageStats}
+	{@const liveStats = processingState.getLiveProcessingStats()}
+	{@const genStats = processingState.getLiveGenerationStats()}
+
+	{#if genStats}
+		<ChatMessageStatistics
+			mode={ChatMessageStatisticsMode.GENERATION}
+			isLive
+			promptTokens={liveStats?.tokensProcessed}
+			promptMs={liveStats?.timeMs}
+			predictedTokens={genStats.tokensGenerated}
+			predictedMs={genStats.timeMs}
+		/>
+	{/if}
+{/if}

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

@@ -0,0 +1,66 @@
+<script lang="ts">
+	import { BuiltInTool } from '$lib/enums';
+	import {
+		extractSearchQuery,
+		extractSearchResults,
+		isWebSearchToolName,
+		type AgenticSection
+	} from '$lib/utils';
+	import type { DatabaseMessageExtra } from '$lib/types';
+	import ChatMessageToolCallBlockDefault from './ChatMessageToolCallBlockDefault.svelte';
+	import ChatMessageToolCallBlockEditFile from './ChatMessageToolCallBlockEditFile.svelte';
+	import ChatMessageToolCallBlockExecShellCommand from './ChatMessageToolCallBlockExecShellCommand.svelte';
+	import ChatMessageToolCallBlockFileGlobSearch from './ChatMessageToolCallBlockFileGlobSearch.svelte';
+	import ChatMessageToolCallBlockGetDatetime from './ChatMessageToolCallBlockGetDatetime.svelte';
+	import ChatMessageToolCallBlockGrepSearch from './ChatMessageToolCallBlockGrepSearch.svelte';
+	import ChatMessageToolCallBlockReadFile from './ChatMessageToolCallBlockReadFile.svelte';
+	import ChatMessageToolCallBlockRunJavascript from './ChatMessageToolCallBlockRunJavascript.svelte';
+	import ChatMessageToolCallBlockSearchResults from './ChatMessageToolCallBlockSearchResults.svelte';
+	import ChatMessageToolCallBlockWriteFile from './ChatMessageToolCallBlockWriteFile.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		attachments?: DatabaseMessageExtra[];
+		open: boolean;
+		isStreaming: boolean;
+		isExecuting?: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, attachments, open, isStreaming, isExecuting, onToggle }: Props = $props();
+
+	const searchResults = $derived(extractSearchResults(section.toolResult));
+	const searchQuery = $derived(extractSearchQuery(section.toolArgs));
+	const isSearchCall = $derived(
+		searchResults.length > 0 || (searchQuery.length > 0 && isWebSearchToolName(section.toolName))
+	);
+</script>
+
+{#if isSearchCall}
+	<ChatMessageToolCallBlockSearchResults {section} {open} {isStreaming} {onToggle} />
+{:else if section.toolName === BuiltInTool.GET_DATETIME}
+	<ChatMessageToolCallBlockGetDatetime {section} {isStreaming} />
+{:else if section.toolName === BuiltInTool.READ_FILE}
+	<ChatMessageToolCallBlockReadFile {section} {open} {isStreaming} {onToggle} />
+{:else if section.toolName === BuiltInTool.EDIT_FILE}
+	<ChatMessageToolCallBlockEditFile {section} {open} {isStreaming} {onToggle} />
+{:else if section.toolName === BuiltInTool.WRITE_FILE}
+	<ChatMessageToolCallBlockWriteFile {section} {open} {isStreaming} {onToggle} />
+{:else if section.toolName === BuiltInTool.EXEC_SHELL_COMMAND}
+	<ChatMessageToolCallBlockExecShellCommand
+		{section}
+		{open}
+		{isStreaming}
+		{isExecuting}
+		{attachments}
+		{onToggle}
+	/>
+{:else if section.toolName === BuiltInTool.FILE_GLOB_SEARCH}
+	<ChatMessageToolCallBlockFileGlobSearch {section} {open} {isStreaming} {onToggle} />
+{:else if section.toolName === BuiltInTool.GREP_SEARCH}
+	<ChatMessageToolCallBlockGrepSearch {section} {open} {isStreaming} {onToggle} />
+{:else if section.toolName === BuiltInTool.RUN_JAVASCRIPT}
+	<ChatMessageToolCallBlockRunJavascript {section} {open} {isStreaming} {onToggle} />
+{:else}
+	<ChatMessageToolCallBlockDefault {section} {open} {isStreaming} {attachments} {onToggle} />
+{/if}

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

@@ -0,0 +1,124 @@
+<script lang="ts">
+	// Fall-through renderer for tool calls without a dedicated block.
+	// Renders section.toolArgs / section.toolResult directly using the
+	// shared chrome shell.
+
+	import { Loader2 } from '@lucide/svelte';
+	import { MarkdownContent, SyntaxHighlightedCode } from '$lib/components/app';
+	import { FileTypeText, ToolResultKind } from '$lib/enums';
+	import { MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
+	import {
+		classifyToolResult,
+		formatJsonPretty,
+		parseToolResultWithImages,
+		type AgenticSection,
+		type ToolResultLine
+	} from '$lib/utils';
+	import { getBuiltinToolUi } from '$lib/constants/built-in-tools';
+	import type { DatabaseMessageExtra } from '$lib/types';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		attachments?: DatabaseMessageExtra[];
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, attachments, onToggle }: Props = $props();
+
+	const title = $derived(getBuiltinToolUi(section.toolName)?.label ?? section.toolName ?? '');
+
+	const parsedLines: ToolResultLine[] = $derived(
+		section.toolResult ? parseToolResultWithImages(section.toolResult, attachments) : []
+	);
+	const outputKind = $derived(classifyToolResult(section.toolResult));
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={null} {title} {onToggle}>
+	{#snippet children(_meta, ctx)}
+		{#if ctx.isStreamingCall}
+			<div class="mb-2 flex items-center gap-2 text-xs text-muted-foreground/70">
+				<span>Input</span>
+				{#if ctx.isStreaming}
+					<Loader2 class="h-3 w-3 animate-spin" />
+				{/if}
+			</div>
+			{#if section.toolArgs}
+				<SyntaxHighlightedCode
+					code={formatJsonPretty(section.toolArgs)}
+					language={FileTypeText.JSON}
+					maxHeight={MAX_HEIGHT_CODE_BLOCK}
+					streaming={ctx.isCodeStreaming}
+				/>
+			{:else if ctx.isStreaming}
+				<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">
+					Receiving arguments...
+				</div>
+			{:else}
+				<div
+					class="rounded bg-yellow-500/10 p-2 text-xs text-yellow-600 italic dark:text-yellow-400"
+				>
+					Response was truncated
+				</div>
+			{/if}
+		{:else}
+			{@const showInput = Boolean(section.toolArgs)}
+			{#if showInput}
+				<div class="mb-1.5 flex items-center gap-2 text-xs text-muted-foreground/70">
+					<span>Input</span>
+				</div>
+				<SyntaxHighlightedCode
+					code={formatJsonPretty(section.toolArgs ?? '')}
+					language={FileTypeText.JSON}
+					maxHeight={MAX_HEIGHT_CODE_BLOCK}
+					streaming={ctx.isCodeStreaming}
+				/>
+			{/if}
+			<div
+				class={showInput
+					? 'mt-4 mb-1.5 flex items-center gap-2 text-xs text-muted-foreground/70'
+					: 'mb-1.5 flex items-center gap-2 text-xs text-muted-foreground/70'}
+			>
+				<span>Output</span>
+				{#if ctx.isPending}
+					<Loader2 class="h-3 w-3 animate-spin" />
+				{/if}
+			</div>
+			{#if ctx.isPending}
+				<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">
+					Waiting for result...
+				</div>
+			{:else if section.toolResult}
+				{#if outputKind === ToolResultKind.JSON}
+					<SyntaxHighlightedCode
+						code={formatJsonPretty(section.toolResult)}
+						language={FileTypeText.JSON}
+						maxHeight={MAX_HEIGHT_CODE_BLOCK}
+					/>
+				{:else if outputKind === ToolResultKind.MARKDOWN}
+					<MarkdownContent content={section.toolResult} {attachments} />
+				{:else}
+					<div class="overflow-auto">
+						{#each parsedLines as line, i (i)}
+							<div class="font-mono text-[11px] leading-relaxed whitespace-pre-wrap">
+								{line.text}
+							</div>
+							{#if line.image}
+								<img
+									src={line.image.base64Url}
+									alt={line.image.name}
+									class="mt-2 mb-2 h-auto max-w-full rounded-lg"
+									loading="lazy"
+								/>
+							{/if}
+						{/each}
+					</div>
+				{/if}
+			{:else}
+				<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">No output</div>
+			{/if}
+		{/if}
+	{/snippet}
+</ToolCallBlock>

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

@@ -0,0 +1,166 @@
+<script lang="ts">
+	import { XCircle } from '@lucide/svelte';
+	import { MAX_HEIGHT_CODE_BLOCK, RESULT_STAT_SEPARATOR } from '$lib/constants';
+	import { computeLineDiff, prefixFor, type AgenticSection } from '$lib/utils';
+	import { parseEditFileMeta } from './parsers/edit-file';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, onToggle }: Props = $props();
+
+	const editFileMeta = $derived(parseEditFileMeta(section));
+
+	const editDiffs = $derived(
+		(editFileMeta?.edits ?? []).map((edit) => computeLineDiff(edit.oldText, edit.newText))
+	);
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={editFileMeta} {onToggle}>
+	{#snippet titleSnippet()}
+		<span class="text-muted-foreground">Edit file </span>
+		<span class="font-mono">{editFileMeta?.filePath}</span>
+		{#if editFileMeta?.errorMessage}
+			<span class="ml-1 text-xs italic text-muted-foreground/70">(failed)</span>
+		{/if}
+	{/snippet}
+
+	{#snippet children(meta, _ctx)}
+		{#if meta?.errorMessage}
+			<div
+				class="flex items-start gap-2 rounded bg-red-500/10 p-2 text-xs text-red-600 italic dark:text-red-400"
+			>
+				<XCircle class="mt-0.5 h-3 w-3 shrink-0" />
+				<span>{meta.errorMessage}</span>
+			</div>
+		{:else if meta && meta.edits.length > 0}
+			{#each editDiffs as diffLines, ei (ei)}
+				<div class={ei === 0 ? '' : 'mt-3'}>
+					<div class="mb-1.5 text-xs text-muted-foreground/70 italic">
+						Edit {ei + 1}&nbsp;of&nbsp;{meta.edits.length}
+					</div>
+					<div class="diff-block" style:max-height={MAX_HEIGHT_CODE_BLOCK}>
+						<div class="diff-pre">
+							{#each diffLines as line, li (li)}
+								<div class="diff-line diff-{line.kind}">
+									<span class="diff-old-num">{line.oldLine ?? ''}</span>
+									<span class="diff-marker">{prefixFor(line.kind)}</span>
+									<span class="diff-new-num">{line.newLine ?? ''}</span>
+									<span class="diff-text">{line.text || ' '}</span>
+								</div>
+							{/each}
+						</div>
+					</div>
+				</div>
+			{/each}
+			<div class="mt-1.5 text-xs text-muted-foreground/70 italic">
+				{#if meta.resultMessage}
+					{meta.resultMessage}{meta.editsApplied != null ? RESULT_STAT_SEPARATOR : ''}{/if}
+				{#if meta.editsApplied != null}
+					<span class="font-mono">{meta.editsApplied}</span>
+					{meta.editsApplied === 1 ? 'edit' : 'edits'}&nbsp;applied
+				{/if}
+			</div>
+		{:else}
+			<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">No edits</div>
+		{/if}
+	{/snippet}
+</ToolCallBlock>
+
+<style>
+	.diff-block {
+		overflow: auto;
+		border-radius: 0.75rem;
+		border-width: 1px;
+		border-color: color-mix(in oklch, var(--border) 30%, transparent);
+		background: var(--code-background);
+		box-shadow: 0 1px 2px 0 rgb(0 0 0 / 0.05);
+	}
+
+	:global(.dark) .diff-block {
+		border-color: color-mix(in oklch, var(--border) 20%, transparent);
+	}
+
+	/* Each row is a 4-column grid: old-line#, marker, new-line#, text.
+	 * The gutters stay fixed-width so the text column lines up unversally. */
+	.diff-line {
+		display: grid;
+		grid-template-columns: 3.25rem 1.5rem 3.25rem 1fr;
+		font-family: var(--font-mono);
+		font-size: 11px;
+		line-height: 1.65;
+		align-items: stretch;
+	}
+
+	.diff-old-num,
+	.diff-new-num {
+		text-align: right;
+		padding-right: 0.5rem;
+		user-select: none;
+		color: color-mix(in oklch, var(--muted-foreground) 70%, transparent);
+		font-variant-numeric: tabular-nums;
+	}
+
+	.diff-marker {
+		text-align: center;
+		color: color-mix(in oklch, var(--muted-foreground) 70%, transparent);
+		user-select: none;
+	}
+
+	.diff-line.diff-add {
+		background-color: #f0fff4;
+		color: #22863a;
+	}
+	.diff-line.diff-add .diff-new-num,
+	.diff-line.diff-add .diff-marker {
+		color: #22863a;
+	}
+
+	.diff-line.diff-remove {
+		background-color: #ffeef0;
+		color: #b31d28;
+	}
+	.diff-line.diff-remove .diff-old-num,
+	.diff-line.diff-remove .diff-marker {
+		color: #b31d28;
+	}
+
+	.diff-line.diff-add .diff-old-num,
+	.diff-line.diff-remove .diff-new-num {
+		/* Empty gutter columns for add/remove rows mirror git unification
+		 * (added lines don't have an old number, removed lines don't have a
+		 * new number). Keep them visible so columns stay aligned across
+		 * mixed rows. */
+		opacity: 0;
+	}
+
+	.diff-text {
+		padding-left: 0.4rem;
+		padding-right: 0.5rem;
+		white-space: pre;
+		overflow-x: auto;
+		min-width: 0;
+	}
+
+	:global(.dark) .diff-line.diff-add {
+		background-color: #033a16;
+		color: #aff5b4;
+	}
+	:global(.dark) .diff-line.diff-add .diff-new-num,
+	:global(.dark) .diff-line.diff-add .diff-marker {
+		color: #aff5b4;
+	}
+	:global(.dark) .diff-line.diff-remove {
+		background-color: #67060c;
+		color: #ffdcd7;
+	}
+	:global(.dark) .diff-line.diff-remove .diff-old-num,
+	:global(.dark) .diff-line.diff-remove .diff-marker {
+		color: #ffdcd7;
+	}
+</style>

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

@@ -0,0 +1,293 @@
+<script lang="ts">
+	// Block for `exec_shell_command`. Unlike the other tools, this
+	// renderer uses CollapsibleTerminalBlock (terminal-style frame)
+	// and treats "live" output chunks as active even after the call
+	// resolved, so the spinner stays on while stdout is still flowing.
+	// The scroll-to-bottom auto-scroll logic mirrors what was here
+	// before extraction.
+
+	import { Check, Loader2, XCircle, AlertTriangle } from '@lucide/svelte';
+	import { CollapsibleTerminalBlock } from '$lib/components/app';
+	import { SETTINGS_KEYS } from '$lib/constants';
+	import { config } from '$lib/stores/settings.svelte';
+	import { TOOL_RUNTIME_SCROLL_AT_BOTTOM_THRESHOLD_PX } from '$lib/constants/auto-scroll';
+	import {
+		highlightCode,
+		isExitCodeSummaryLine,
+		parseExecShellCommandError,
+		parseExecShellCommandExitStatus,
+		parseToolResultWithImages,
+		type AgenticSection,
+		type ExecShellExitStatus,
+		type ToolResultLine
+	} from '$lib/utils';
+	import { parseExecShellCommandMeta } from './parsers/exec-shell-command';
+	import type { DatabaseMessageExtra } from '$lib/types';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		/** True while the agentic loop is streaming output chunks for THIS
+		 *  tool call. Drives max-height + auto-scroll while true; releases
+		 *  them when the loop reports this call as done. */
+		isExecuting?: boolean;
+		attachments?: DatabaseMessageExtra[];
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, isExecuting = false, attachments, onToggle }: Props = $props();
+
+	// `isLive` covers all in-flight phases: pre-chunk spinner and
+	// streaming itself. Frozen output (tool done while agent continues)
+	// is not live.
+	const isLive = $derived(isExecuting);
+
+	const execShellMeta = $derived(parseExecShellCommandMeta(section));
+	const execShellError = $derived(parseExecShellCommandError(section.toolResult));
+	const execShellExitStatus: ExecShellExitStatus | undefined = $derived(
+		parseExecShellCommandExitStatus(section.toolResult)
+	);
+
+	const parsedLines: ToolResultLine[] = $derived(
+		section.toolResult ? parseToolResultWithImages(section.toolResult, attachments) : []
+	);
+
+	// Drop the trailing "[exit code: N]" line - rendered as a colored
+	// badge below. During streaming we keep it so a partial stream still
+	// shows the status once the final chunk lands.
+	const outputLines: ToolResultLine[] = $derived(
+		execShellExitStatus && parsedLines.length > 0
+			? parsedLines.slice(0, parsedLines.length - 1)
+			: parsedLines
+	);
+
+	const isExitCodeFinalLine = $derived(
+		execShellExitStatus !== undefined &&
+			parsedLines.length > 0 &&
+			isExitCodeSummaryLine(parsedLines[parsedLines.length - 1].text, execShellExitStatus)
+	);
+
+	// Highlight just the command for the title; the (typically large)
+	// output blob uses bare monospace to skip hljs per-line highlighting.
+	const highlightedCommandHtml = $derived(
+		execShellMeta ? highlightCode(execShellMeta.command, 'bash') : ''
+	);
+
+	const exitBadgeClass = $derived(
+		execShellExitStatus?.timedOut
+			? 'exit-badge warning'
+			: execShellExitStatus?.code === 0
+				? 'exit-badge success'
+				: 'exit-badge failure'
+	);
+
+	const useFullHeightCodeBlocks = $derived(
+		Boolean(config()[SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS])
+	);
+
+	const autoScroll = $derived(isLive && !useFullHeightCodeBlocks);
+
+	const SCROLL_BOTTOM_THRESHOLD_PX = TOOL_RUNTIME_SCROLL_AT_BOTTOM_THRESHOLD_PX;
+
+	let scrollEl: HTMLDivElement | undefined = $state();
+	let userScrolledUp = $state(false);
+	let lastScrollTop = 0;
+	let pendingFrame: number | null = null;
+
+	function isAtBottom(): boolean {
+		if (!scrollEl) return false;
+		return (
+			scrollEl.scrollHeight - scrollEl.clientHeight - scrollEl.scrollTop <=
+			SCROLL_BOTTOM_THRESHOLD_PX
+		);
+	}
+
+	function scrollToBottomOnFrame() {
+		if (pendingFrame !== null || !scrollEl || userScrolledUp) return;
+		pendingFrame = requestAnimationFrame(() => {
+			pendingFrame = null;
+
+			// Re-check on rAF - user may scroll between scheduling and paint.
+			if (scrollEl && !userScrolledUp) {
+				scrollEl.scrollTop = scrollEl.scrollHeight;
+			}
+		});
+	}
+
+	function handleScrollEvent() {
+		if (!scrollEl) return;
+		const isScrollingUp = scrollEl.scrollTop < lastScrollTop;
+		if (isScrollingUp && !isAtBottom()) {
+			userScrolledUp = true;
+		} else if (isAtBottom()) {
+			userScrolledUp = false;
+		}
+		lastScrollTop = scrollEl.scrollTop;
+	}
+
+	$effect(() => {
+		void section.toolResult;
+		if (!scrollEl || !autoScroll) return;
+		scrollToBottomOnFrame();
+	});
+
+	$effect(() => {
+		// Catch layout changes that don't touch toolResult (line-wrap
+		// reflow, image attaches, hljs settle).
+		if (!scrollEl || !autoScroll) return;
+
+		const observer = new MutationObserver(() => scrollToBottomOnFrame());
+		observer.observe(scrollEl, {
+			childList: true,
+			subtree: true,
+			characterData: true
+		});
+
+		return () => observer.disconnect();
+	});
+
+	$effect(() => {
+		// Reset on stream end so the next render (full-height) starts
+		// pinned.
+		if (!isLive) {
+			userScrolledUp = false;
+			lastScrollTop = 0;
+		}
+	});
+</script>
+
+{#snippet execShellTitle()}
+	{#if highlightedCommandHtml}
+		<span class="font-mono">{@html highlightedCommandHtml}</span>
+	{:else}
+		<span class="font-mono">{execShellMeta?.command}</span>
+	{/if}
+{/snippet}
+
+<ToolCallBlock
+	{section}
+	{open}
+	{isStreaming}
+	meta={execShellMeta ? { errorMessage: execShellError } : null}
+	wrapper={CollapsibleTerminalBlock}
+	extraLiveStreaming={isLive}
+	spinIconWhenActive={true}
+	{onToggle}
+>
+	{#snippet titleSnippet()}
+		{@render execShellTitle()}
+	{/snippet}
+
+	{#snippet children(_meta, ctx)}
+		{#if ctx.isPending}
+			<div class="flex items-start gap-2 text-xs text-muted-foreground/70">
+				<Loader2 class="h-3 w-3 animate-spin" />
+				Running...
+			</div>
+		{:else if execShellError}
+			<div class="flex items-start gap-2 text-xs text-red-600 italic dark:text-red-400">
+				<XCircle class="mt-0.5 h-3 w-3 shrink-0" />
+				<span>{execShellError}</span>
+			</div>
+		{:else if section.toolResult}
+			<div
+				bind:this={scrollEl}
+				class="terminal-output"
+				class:is-clamped={!useFullHeightCodeBlocks}
+				onscroll={handleScrollEvent}
+			>
+				{#each outputLines as line, i (i)}
+					<div class="font-mono text-[11px] leading-relaxed whitespace-pre-wrap">{line.text}</div>
+					{#if line.image}
+						<img
+							src={line.image.base64Url}
+							alt={line.image.name}
+							class="mt-2 mb-2 h-auto max-w-full rounded-lg"
+							loading="lazy"
+						/>
+					{/if}
+				{/each}
+
+				{#if isExitCodeFinalLine && execShellExitStatus}
+					<div class={exitBadgeClass}>
+						{#if execShellExitStatus.timedOut}
+							<AlertTriangle class="h-3 w-3" />
+							<span>timed out</span>
+							<span class="exit-sep">&middot;</span>
+							<span>exit {execShellExitStatus.code}</span>
+						{:else if execShellExitStatus.code === 0}
+							<Check class="h-3 w-3" />
+							<span>exit 0</span>
+						{:else}
+							<XCircle class="h-3 w-3" />
+							<span>exit {execShellExitStatus.code}</span>
+						{/if}
+					</div>
+				{/if}
+			</div>
+		{/if}
+	{/snippet}
+</ToolCallBlock>
+
+<style>
+	.terminal-output {
+		overscroll-behavior: contain;
+	}
+
+	.terminal-output.is-clamped {
+		max-height: 28rem;
+		overflow-y: auto;
+		scrollbar-gutter: stable;
+		padding-right: 0.25rem;
+	}
+
+	.exit-badge {
+		display: inline-flex;
+		align-items: center;
+		gap: 0.35rem;
+		margin-top: 0.5rem;
+		padding: 0.2rem 0.55rem;
+		border-radius: 0.375rem;
+		font-family: var(--font-mono);
+		font-size: 11px;
+		font-weight: 500;
+		letter-spacing: 0.01em;
+		line-height: 1;
+	}
+
+	.exit-badge.success {
+		background: color-mix(in oklch, var(--color-green-500, #22c55e) 14%, transparent);
+		color: var(--color-green-700, #15803d);
+	}
+
+	:global(.dark) .exit-badge.success {
+		background: color-mix(in oklch, var(--color-green-400, #4ade80) 18%, transparent);
+		color: var(--color-green-300, #86efac);
+	}
+
+	.exit-badge.failure {
+		background: color-mix(in oklch, var(--color-red-500, #ef4444) 14%, transparent);
+		color: var(--color-red-700, #b91c1c);
+	}
+
+	:global(.dark) .exit-badge.failure {
+		background: color-mix(in oklch, var(--color-red-400, #f87171) 18%, transparent);
+		color: var(--color-red-300, #fca5a5);
+	}
+
+	.exit-badge.warning {
+		background: color-mix(in oklch, var(--color-amber-500, #f59e0b) 14%, transparent);
+		color: var(--color-amber-700, #b45309);
+	}
+
+	:global(.dark) .exit-badge.warning {
+		background: color-mix(in oklch, var(--color-amber-400, #fbbf24) 18%, transparent);
+		color: var(--color-amber-300, #fcd34d);
+	}
+
+	.exit-sep {
+		opacity: 0.45;
+	}
+</style>

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

@@ -0,0 +1,61 @@
+<script lang="ts">
+	import { XCircle } from '@lucide/svelte';
+	import { type AgenticSection } from '$lib/utils';
+	import { parseFileGlobSearchMeta } from './parsers/file-glob-search';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, onToggle }: Props = $props();
+
+	const fileGlobMeta = $derived(parseFileGlobSearchMeta(section));
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={fileGlobMeta} {onToggle}>
+	{#snippet titleSnippet()}
+		{#if fileGlobMeta}
+			<span class="text-muted-foreground"
+				>{fileGlobMeta.include === '**' ? 'List files' : 'Search files'}&nbsp;</span
+			>
+			{#if fileGlobMeta.include !== '**'}
+				<span class="font-mono">{fileGlobMeta.include}</span>
+			{/if}
+			<span class="text-muted-foreground">&nbsp;in&nbsp;</span>
+			<span class="font-mono">{fileGlobMeta.path}</span>
+		{/if}
+	{/snippet}
+
+	{#snippet children(meta, ctx)}
+		{#if ctx.isPending}
+			<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">
+				Searching...
+			</div>
+		{:else if meta?.errorMessage}
+			<div
+				class="flex items-start gap-2 rounded bg-red-500/10 p-2 text-xs text-red-600 italic dark:text-red-400"
+			>
+				<XCircle class="mt-0.5 h-3 w-3 shrink-0" />
+				<span>{meta.errorMessage}</span>
+			</div>
+		{:else if meta && meta.matches.length > 0}
+			<div class="max-h-96 overflow-auto">
+				{#each meta.matches as match, i (i)}
+					<div class="font-mono text-[11px] leading-relaxed whitespace-pre-wrap">{match}</div>
+				{/each}
+			</div>
+			<div class="mt-1.5 text-xs text-muted-foreground/70 italic">
+				Total matches: <span class="font-mono">{meta.totalMatches ?? meta.matches.length}</span>
+			</div>
+		{:else}
+			<div class="text-xs text-muted-foreground/70 italic">No matches</div>
+			<div class="mt-1.5 text-xs text-muted-foreground/70 italic">
+				Total matches: <span class="font-mono">{meta?.totalMatches ?? 0}</span>
+			</div>
+		{/if}
+	{/snippet}
+</ToolCallBlock>

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

@@ -0,0 +1,57 @@
+<script lang="ts">
+	import { Clock, Loader2 } from '@lucide/svelte';
+	import { AgenticSectionType } from '$lib/enums';
+	import type { AgenticSection } from '$lib/utils';
+
+	interface Props {
+		section: AgenticSection;
+		isStreaming?: boolean;
+	}
+
+	let { section, isStreaming = false }: Props = $props();
+
+	const isPending = $derived(section.type === AgenticSectionType.TOOL_CALL_PENDING);
+	const isStreamingCall = $derived(section.type === AgenticSectionType.TOOL_CALL_STREAMING);
+	const showSpinner = $derived(isPending || (isStreamingCall && isStreaming));
+
+	type GetDatetimeMeta = {
+		dateString?: string;
+		errorMessage?: string;
+	};
+
+	function parseGetDatetimeMeta(toolResultString: string | undefined): GetDatetimeMeta {
+		if (!toolResultString) return {};
+
+		try {
+			const parsed: unknown = JSON.parse(toolResultString);
+			if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) {
+				const obj = parsed as Record<string, unknown>;
+				if (typeof obj.error === 'string') return { errorMessage: obj.error };
+				if (typeof obj.result === 'string') return { dateString: obj.result.trim() };
+			}
+		} catch {
+			return { dateString: toolResultString.trim() };
+		}
+
+		return {};
+	}
+
+	const dateMeta = $derived(parseGetDatetimeMeta(section.toolResult));
+</script>
+
+<div class="text-muted-foreground flex items-center gap-2 py-1.5">
+	<Clock class="text-muted-foreground/60 h-3.5 w-3.5 shrink-0" />
+	{#if showSpinner}
+		<span class="text-foreground/80 text-sm font-medium">Current time</span>
+		<Loader2 class="text-muted-foreground/70 h-3 w-3 animate-spin" />
+	{:else if dateMeta.errorMessage}
+		<span class="text-foreground/80 text-sm font-medium">Current time&nbsp;</span>
+		<span class="text-red-600 text-xs italic dark:text-red-400">-&nbsp;{dateMeta.errorMessage}</span
+		>
+	{:else if dateMeta.dateString}
+		<span class="text-foreground/80 text-sm font-medium">Current time is&nbsp;</span>
+		<span class="font-mono text-foreground/90 text-sm">{dateMeta.dateString}</span>
+	{:else}
+		<span class="text-foreground/80 text-sm font-medium">Current time</span>
+	{/if}
+</div>

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

@@ -0,0 +1,67 @@
+<script lang="ts">
+	import { XCircle } from '@lucide/svelte';
+	import { type AgenticSection } from '$lib/utils';
+	import { parseGrepSearchMeta } from './parsers/grep-search';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, onToggle }: Props = $props();
+
+	const grepMeta = $derived(parseGrepSearchMeta(section));
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={grepMeta} {onToggle}>
+	{#snippet titleSnippet()}
+		{#if grepMeta}
+			<span class="text-muted-foreground">Search for&nbsp;</span>
+			<span class="font-mono">{grepMeta.pattern}</span>
+			<span class="text-muted-foreground">&nbsp;in&nbsp;</span>
+			<span class="font-mono">{grepMeta.path}</span>
+		{/if}
+	{/snippet}
+
+	{#snippet children(meta, ctx)}
+		{#if ctx.isPending}
+			<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">
+				Searching...
+			</div>
+		{:else if meta?.errorMessage}
+			<div
+				class="flex items-start gap-2 rounded bg-red-500/10 p-2 text-xs text-red-600 italic dark:text-red-400"
+			>
+				<XCircle class="mt-0.5 h-3 w-3 shrink-0" />
+				<span>{meta.errorMessage}</span>
+			</div>
+		{:else if meta && meta.matches.length > 0}
+			<div class="max-h-96 overflow-auto">
+				{#each meta.matches as match, mi (mi)}
+					<div class="font-mono text-[11px] leading-relaxed">
+						<span class="text-muted-foreground/70">{match.file}</span>
+						{#if meta.showLineNumbers && match.line != null}
+							<span class="text-muted-foreground/70">:{match.line}</span>
+						{/if}
+						<span class="text-muted-foreground/70">:</span>
+						<span>{match.content}</span>
+					</div>
+				{/each}
+			</div>
+			<div class="mt-1.5 text-xs text-muted-foreground/70 italic">
+				Total matches: <span class="font-mono">{meta.totalMatches ?? meta.matches.length}</span>
+				{#if meta.showLineNumbers}
+					&nbsp;<span class="italic">(with line numbers)</span>
+				{/if}
+			</div>
+		{:else}
+			<div class="text-xs text-muted-foreground/70 italic">No matches</div>
+			<div class="mt-1.5 text-xs text-muted-foreground/70 italic">
+				Total matches: <span class="font-mono">{meta?.totalMatches ?? 0}</span>
+			</div>
+		{/if}
+	{/snippet}
+</ToolCallBlock>

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

@@ -0,0 +1,44 @@
+<script lang="ts">
+	import { SyntaxHighlightedCode } from '$lib/components/app';
+	import { DEFAULT_LANGUAGE, MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
+	import { type AgenticSection } from '$lib/utils';
+	import { parseReadFileMeta } from './parsers/read-file';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, onToggle }: Props = $props();
+
+	const readFileMeta = $derived(parseReadFileMeta(section));
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={readFileMeta} {onToggle}>
+	{#snippet titleSnippet()}
+		<span class="text-muted-foreground">Read file </span>
+		<span class="font-mono">{readFileMeta?.fileName}</span>
+		{#if readFileMeta?.lineRange}
+			<span class="text-muted-foreground"
+				>&nbsp;(lines {readFileMeta.lineRange.start}-{readFileMeta.lineRange.end})</span
+			>
+		{/if}
+	{/snippet}
+
+	{#snippet children(_meta, _ctx)}
+		{#if section.toolResult}
+			<SyntaxHighlightedCode
+				code={section.toolResult}
+				language={readFileMeta?.language ?? DEFAULT_LANGUAGE}
+				maxHeight={MAX_HEIGHT_CODE_BLOCK}
+			/>
+		{:else}
+			<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">
+				Waiting for file content...
+			</div>
+		{/if}
+	{/snippet}
+</ToolCallBlock>

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

@@ -0,0 +1,69 @@
+<script lang="ts">
+	import { XCircle, Terminal } from '@lucide/svelte';
+	import { SyntaxHighlightedCode } from '$lib/components/app';
+	import { FileTypeText } from '$lib/enums';
+	import { MAX_HEIGHT_CODE_BLOCK } from '$lib/constants';
+	import { getBuiltinToolUi, type AgenticSection } from '$lib/utils';
+	import { parseRunJavascriptMeta } from './parsers/run-javascript';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, onToggle }: Props = $props();
+
+	const runJsMeta = $derived(parseRunJavascriptMeta(section));
+	const title = $derived(getBuiltinToolUi(section.toolName)?.label ?? section.toolName ?? '');
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={runJsMeta} {title} {onToggle}>
+	{#snippet children(meta, ctx)}
+		{#if ctx.isPending}
+			<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">Running...</div>
+		{:else if meta?.errorMessage}
+			<div
+				class="flex items-start gap-2 rounded bg-red-500/10 p-2 text-xs text-red-600 italic dark:text-red-400"
+			>
+				<XCircle class="mt-0.5 h-3 w-3 shrink-0" />
+				<span>{meta.errorMessage}</span>
+			</div>
+			<div class="mt-3">
+				<SyntaxHighlightedCode
+					code={meta.code}
+					language={FileTypeText.JAVASCRIPT}
+					maxHeight={MAX_HEIGHT_CODE_BLOCK}
+					streaming={ctx.isCodeStreaming}
+				/>
+			</div>
+		{:else if meta}
+			<SyntaxHighlightedCode
+				code={meta.code}
+				language={FileTypeText.JAVASCRIPT}
+				maxHeight={MAX_HEIGHT_CODE_BLOCK}
+				streaming={ctx.isCodeStreaming}
+			/>
+			<div class="mb-2 mt-3 flex items-center gap-2 text-xs text-muted-foreground/70">
+				<Terminal class="h-3 w-3" />
+				<span>Console</span>
+				{#if meta.timeoutMs != null}
+					<span class="font-mono">&middot;&nbsp;timeout&nbsp;{meta.timeoutMs}&nbsp;ms</span>
+				{/if}
+			</div>
+			{#if section.toolResult}
+				<div class="mt-1">
+					<SyntaxHighlightedCode
+						code={section.toolResult}
+						language={FileTypeText.JAVASCRIPT}
+						maxHeight={MAX_HEIGHT_CODE_BLOCK}
+					/>
+				</div>
+			{:else}
+				<div class="rounded bg-muted/20 p-2 text-xs text-muted-foreground/70 italic">No output</div>
+			{/if}
+		{/if}
+	{/snippet}
+</ToolCallBlock>

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

@@ -0,0 +1,167 @@
+<script lang="ts">
+	import { ICON_CLASS_DEFAULT, ICON_CLASS_SPIN } from '$lib/constants/css-classes';
+	import { Globe, Loader2 } from '@lucide/svelte';
+	import { CollapsibleContentBlock } from '$lib/components/app';
+	import * as HoverCard from '$lib/components/ui/hover-card';
+	import { AgenticSectionType } from '$lib/enums';
+	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import {
+		extractSearchResults,
+		extractSearchQuery,
+		faviconForUrl,
+		sanitizeExternalUrl,
+		type SearchResult,
+		type AgenticSection
+	} from '$lib/utils';
+
+	interface Props {
+		section: AgenticSection;
+		open?: boolean;
+		isStreaming?: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, open = $bindable(false), isStreaming = false, onToggle }: Props = $props();
+
+	const isPending = $derived(section.type === AgenticSectionType.TOOL_CALL_PENDING);
+	const isStreamingCall = $derived(section.type === AgenticSectionType.TOOL_CALL_STREAMING);
+	const showSpinner = $derived(isPending || (isStreamingCall && isStreaming));
+
+	const results: SearchResult[] = $derived(extractSearchResults(section.toolResult));
+	const query = $derived(extractSearchQuery(section.toolArgs));
+
+	// Same icon-resolution chain as ChatMessageToolCallBlockDefault so
+	// MCP-server branding is consistent across both views. Spinner wins
+	// while the call is in flight so the user sees execution status.
+	const iconUrl = $derived(showSpinner ? null : mcpStore.getServerFaviconForTool(section.toolName));
+	const icon = $derived(showSpinner ? Loader2 : undefined);
+	const iconClass = $derived(showSpinner ? ICON_CLASS_SPIN : ICON_CLASS_DEFAULT);
+
+	// Verb reflects state: "Searching" while the call is in flight, "Searched"
+	// once results (or a definitive empty response) have arrived. Lets the
+	// heading read as a live progress indicator rather than a completed
+	// retrospective.
+	const title = $derived.by(() => {
+		const verb = showSpinner ? 'Searching' : 'Searched';
+		return query ? `${verb} web for "${query}"` : `${verb} web`;
+	});
+
+	function hideBrokenIcon(event: Event) {
+		(event.currentTarget as HTMLImageElement).style.display = 'none';
+	}
+
+	function formatPublishDate(iso: string | undefined): string | null {
+		if (!iso) return null;
+		try {
+			const date = new Date(iso);
+			if (Number.isNaN(date.getTime())) return iso;
+			return date.toLocaleDateString(undefined, {
+				year: 'numeric',
+				month: 'short',
+				day: 'numeric'
+			});
+		} catch {
+			return iso;
+		}
+	}
+
+	function hostFor(url: string): string | null {
+		try {
+			return new URL(url).host;
+		} catch {
+			return null;
+		}
+	}
+
+	function hasDetails(result: SearchResult): boolean {
+		return Boolean(result.highlights || result.published || result.author);
+	}
+</script>
+
+{#snippet pill(result: SearchResult)}
+	{@const faviconUrl = faviconForUrl(result.url)}
+	{@const safeUrl = sanitizeExternalUrl(result.url)}
+	{@const showHoverCard = safeUrl !== null && hasDetails(result)}
+	{#if safeUrl}
+		<HoverCard.Root openDelay={150} closeDelay={100}>
+			<HoverCard.Trigger
+				href={safeUrl}
+				target="_blank"
+				rel="noopener noreferrer"
+				class="hover:bg-muted/80 focus-visible:ring-ring inline-flex max-w-full items-center gap-1.5 rounded-full border bg-muted px-2.5 py-1 text-xs transition-colors outline-none focus-visible:ring-2"
+			>
+				{#if faviconUrl}
+					<img
+						src={faviconUrl}
+						alt=""
+						class="h-3 w-3 shrink-0 rounded-sm"
+						onerror={hideBrokenIcon}
+					/>
+				{:else}
+					<Globe class="text-muted-foreground/70 h-3 w-3 shrink-0" />
+				{/if}
+				<span class="truncate font-medium text-foreground/80">{result.title}</span>
+			</HoverCard.Trigger>
+			{#if showHoverCard}
+				{@const publishDate = formatPublishDate(result.published)}
+				{@const host = hostFor(safeUrl)}
+				<HoverCard.Content
+					side="top"
+					align="start"
+					sideOffset={6}
+					class="bg-popover text-popover-foreground z-50 w-80 max-w-[90vw] rounded-lg border p-0 shadow-lg"
+				>
+					<div class="flex flex-col gap-2 p-3">
+						<a
+							href={safeUrl}
+							target="_blank"
+							rel="noopener noreferrer"
+							class="line-clamp-3 text-sm font-medium leading-snug hover:underline"
+							>{result.title}</a
+						>
+						{#if publishDate || result.author}
+							<div class="text-muted-foreground flex items-center gap-1.5 text-[11px]">
+								{#if publishDate}
+									<span>{publishDate}</span>
+								{/if}
+								{#if publishDate && result.author}
+									<span class="opacity-50">&middot;</span>
+								{/if}
+								{#if result.author}
+									<span class="truncate">{result.author}</span>
+								{/if}
+							</div>
+						{/if}
+						{#if result.highlights}
+							<p
+								class="text-popover-foreground/85 line-clamp-5 text-xs leading-relaxed whitespace-pre-line"
+							>
+								{result.highlights}
+							</p>
+						{/if}
+						{#if host}
+							<div class="text-muted-foreground/80 truncate text-[11px]">{host}</div>
+						{/if}
+					</div>
+				</HoverCard.Content>
+			{/if}
+		</HoverCard.Root>
+	{/if}
+{/snippet}
+
+<CollapsibleContentBlock {open} class="my-2" {icon} {iconClass} {iconUrl} {title} {onToggle}>
+	{#if results.length > 0}
+		<div class="flex flex-wrap items-center gap-2 pb-1">
+			{#each results as result (result.url)}
+				{@render pill(result)}
+			{/each}
+		</div>
+	{:else if showSpinner}
+		<div class="text-muted-foreground/70 flex items-center gap-2 py-1 text-xs italic">
+			<Loader2 class="h-3 w-3 animate-spin" />
+			<span>Searching...</span>
+		</div>
+	{:else}
+		<div class="text-muted-foreground/70 py-1 text-xs italic">No results</div>
+	{/if}
+</CollapsibleContentBlock>

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

@@ -0,0 +1,55 @@
+<script lang="ts">
+	import { XCircle } from '@lucide/svelte';
+	import { SyntaxHighlightedCode } from '$lib/components/app';
+	import { MAX_HEIGHT_CODE_BLOCK, RESULT_STAT_SEPARATOR } from '$lib/constants';
+	import { type AgenticSection } from '$lib/utils';
+	import { parseWriteFileMeta } from './parsers/write-file';
+	import ToolCallBlock from './ToolCallBlock.svelte';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		onToggle?: () => void;
+	}
+
+	let { section, open, isStreaming, onToggle }: Props = $props();
+
+	const writeFileMeta = $derived(parseWriteFileMeta(section));
+</script>
+
+<ToolCallBlock {section} {open} {isStreaming} meta={writeFileMeta} {onToggle}>
+	{#snippet titleSnippet()}
+		<span class="text-muted-foreground">Write file </span>
+		<span class="font-mono">{writeFileMeta?.filePath}</span>
+		{#if writeFileMeta?.errorMessage}
+			<span class="ml-1 text-xs italic text-muted-foreground/70">(failed)</span>
+		{/if}
+	{/snippet}
+
+	{#snippet children(meta, ctx)}
+		{#if meta?.errorMessage}
+			<div
+				class="flex items-start gap-2 rounded bg-red-500/10 p-2 text-xs text-red-600 italic dark:text-red-400"
+			>
+				<XCircle class="mt-0.5 h-3 w-3 shrink-0" />
+				<span>{meta.errorMessage}</span>
+			</div>
+		{:else if meta}
+			<SyntaxHighlightedCode
+				code={meta.content}
+				language={meta.language}
+				maxHeight={MAX_HEIGHT_CODE_BLOCK}
+				streaming={ctx.isCodeStreaming}
+			/>
+			<div class="mt-1.5 text-xs text-muted-foreground/70 italic">
+				{#if meta.resultMessage}
+					{meta.resultMessage}{meta.bytesWritten != null ? RESULT_STAT_SEPARATOR : ''}{/if}
+				{#if meta.bytesWritten != null}
+					<span class="font-mono">{meta.bytesWritten}</span>
+					bytes
+				{/if}
+			</div>
+		{/if}
+	{/snippet}
+</ToolCallBlock>

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

@@ -0,0 +1,129 @@
+<script lang="ts" generics="TMeta">
+	// Generic chrome shell shared by every per-tool block under
+	// `ChatMessageToolCall/`. Owns:
+	//   - the collapsible wrapper (defaults to CollapsibleContentBlock;
+	//     `exec_shell_command` swaps in CollapsibleTerminalBlock via the
+	//     `wrapper` prop);
+	//   - the icon, spinner state, and MCP favicon fallback chain;
+	//   - the status subtitle pill.
+	// Components supply only their `meta`, a title snippet, and a body
+	// snippet - everything around them is this single source of truth.
+
+	import { Loader2, Wrench } from '@lucide/svelte';
+	import { CollapsibleContentBlock } from '$lib/components/app';
+	import { ICON_CLASS_DEFAULT, ICON_CLASS_SPIN } from '$lib/constants/css-classes';
+	import { AgenticSectionType } from '$lib/enums';
+	import { getBuiltinToolUi } from '$lib/constants/built-in-tools';
+	import { mcpStore } from '$lib/stores/mcp.svelte';
+	import type { Component, Snippet } from 'svelte';
+	import type { AgenticSection, BuiltinToolUiEntry } from '$lib/utils';
+
+	type ToolCallBlockMetaWithError = TMeta & { errorMessage?: string };
+
+	interface ToolCallCtx {
+		isStreaming: boolean;
+		isPending: boolean;
+		isStreamingCall: boolean;
+		isCodeStreaming: boolean;
+	}
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		/**
+		 * The per-tool meta, including any `errorMessage` field that the
+		 * shared chrome uses to compute the status pill subtitle.
+		 */
+		meta: ToolCallBlockMetaWithError | null | undefined;
+		/**
+		 * True while the tool's process is actively producing output
+		 * chunks after its args finished streaming (used by
+		 * `exec_shell_command`'s stdout feed).
+		 */
+		extraLiveStreaming?: boolean;
+		/**
+		 * Swap the title-row icon for a spinning `Loader2` while the
+		 * spinner is showing. Only meaningful for tools where "live"
+		 * is interesting (e.g. exec_shell_command showing the in-flight
+		 * process). Other tools leave it off and render the spinner
+		 * inline within the body.
+		 */
+		spinIconWhenActive?: boolean;
+		/**
+		 * Wrapper component that renders the title row and the body
+		 * children. Defaults to CollapsibleContentBlock;
+		 * `exec_shell_command` uses CollapsibleTerminalBlock for its
+		 * terminal-style frame.
+		 */
+		wrapper?: typeof CollapsibleContentBlock;
+		title?: string;
+		titleSnippet?: Snippet;
+		onToggle?: () => void;
+		children: Snippet<[TMeta | null | undefined, ToolCallCtx]>;
+	}
+
+	let {
+		section,
+		open,
+		isStreaming,
+		meta,
+		extraLiveStreaming = false,
+		spinIconWhenActive = false,
+		wrapper: Wrapper = CollapsibleContentBlock,
+		title,
+		titleSnippet,
+		onToggle,
+		children
+	}: Props = $props();
+
+	const isPending = $derived(section.type === AgenticSectionType.TOOL_CALL_PENDING);
+	const isStreamingCall = $derived(section.type === AgenticSectionType.TOOL_CALL_STREAMING);
+	const showSpinner = $derived(isPending || (isStreamingCall && isStreaming) || extraLiveStreaming);
+	const isCodeStreaming = $derived(isStreaming && (isPending || isStreamingCall));
+
+	const toolUi: BuiltinToolUiEntry | null = $derived(getBuiltinToolUi(section.toolName));
+	const toolIcon: Component = $derived(
+		spinIconWhenActive && showSpinner ? Loader2 : (toolUi?.icon ?? Wrench)
+	);
+	const toolIconClass = $derived(
+		spinIconWhenActive && showSpinner ? ICON_CLASS_SPIN : ICON_CLASS_DEFAULT
+	);
+	// Drop the MCP favicon while the spinner is on so the title row
+	// signals "in flight" without being overwritten by server branding.
+	const mcpServerFavicon = $derived(
+		showSpinner ? null : mcpStore.getServerFaviconForTool(section.toolName)
+	);
+	const iconUrl = $derived(
+		showSpinner || (toolUi?.icon ?? null) || !mcpServerFavicon ? null : mcpServerFavicon
+	);
+
+	function subtitleFor(errorMessage?: string): string | undefined {
+		if (extraLiveStreaming) return 'streaming...';
+		if (showSpinner) return 'executing...';
+		if (errorMessage) return 'failed';
+		if (isStreamingCall && !isStreaming) return 'incomplete';
+		return undefined;
+	}
+
+	const subtitle = $derived(subtitleFor(meta?.errorMessage));
+</script>
+
+<Wrapper
+	{open}
+	class="my-2"
+	icon={toolIcon}
+	iconClass={toolIconClass}
+	{iconUrl}
+	{title}
+	{titleSnippet}
+	{subtitle}
+	{onToggle}
+>
+	{@render children(meta, {
+		isStreaming,
+		isPending,
+		isStreamingCall,
+		isCodeStreaming
+	})}
+</Wrapper>

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

@@ -0,0 +1,49 @@
+// Helpers shared by the per-tool meta parsers under
+// `src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/`.
+// Each tool needs the same first three steps (tool-name check,
+// args-present check, JSON parse) - keeping them here lets each parser
+// stay focused on its own format quirks.
+
+import { BuiltInTool } from '$lib/enums';
+import { parsePartialJsonArgs } from '$lib/utils/parse-partial-json-args';
+import type { AgenticSection } from '$lib/utils/agentic';
+
+/**
+ * Strict (final-state) JSON parser for a tool-args blob. Mirrors the
+ * behaviour the per-tool components used before extraction: an
+ * invalid JSON blob, a JSON array, or a JSON primitive all map to
+ * `null` so callers don't have to guard against surprise shapes.
+ */
+function parseFinalToolArgs(blob: string): Record<string, unknown> | null {
+	try {
+		const parsed: unknown = JSON.parse(blob);
+		if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) {
+			return parsed as Record<string, unknown>;
+		}
+		return null;
+	} catch {
+		return null;
+	}
+}
+
+/**
+ * Parse a section's toolArgs against an expected tool name. Returns
+ * `null` when:
+ *   - the section's toolName doesn't match (component isn't for this
+ *     tool);
+ *   - the section has no args yet (call hasn't started streaming);
+ *   - or the args blob can't be parsed.
+ *
+ * Pass `{ partial: true }` for tools that need to render incrementally
+ * as each token lands (read_file, edit_file, write_file).
+ */
+export function parseToolArgs(
+	expected: BuiltInTool,
+	section: AgenticSection,
+	options: { partial?: boolean } = {}
+): Record<string, unknown> | null {
+	if (section.toolName !== expected || !section.toolArgs) return null;
+	return options.partial
+		? parsePartialJsonArgs(section.toolArgs)
+		: parseFinalToolArgs(section.toolArgs);
+}

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

@@ -0,0 +1,71 @@
+// Meta parser for `edit_file` tool calls. Reads the file path and the
+// array of edits from the streamed args (partial JSON for incremental
+// rendering), plus the result blob for `result` / `edits_applied` /
+// `error` fields.
+
+import { BuiltInTool } from '$lib/enums';
+import { FILE_PATH_SEPARATOR_REGEX } from '$lib/constants';
+import { tryParseToolResultObject, type AgenticSection } from '$lib/utils';
+import { parseToolArgs } from './_shared';
+
+export type EditFileEdit = {
+	oldText: string;
+	newText: string;
+};
+
+export type EditFileMeta = {
+	fileName: string;
+	filePath: string;
+	edits: EditFileEdit[];
+	resultMessage?: string;
+	editsApplied?: number;
+	errorMessage?: string;
+};
+
+export function parseEditFileMeta(section: AgenticSection): EditFileMeta | null {
+	const args = parseToolArgs(BuiltInTool.EDIT_FILE, section, { partial: true });
+	if (!args) return null;
+
+	const rawPath = args.path ?? args.file_path ?? args.filePath;
+	if (typeof rawPath !== 'string' || !rawPath) return null;
+
+	const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath;
+
+	// Filter the streamed edits array strictly: each entry must be an
+	// object with a non-empty `old_text`. Edits without an old_text
+	// would diff against empty and render as a full re-write.
+	const rawEdits = Array.isArray(args.edits) ? args.edits : [];
+	const edits: EditFileEdit[] = [];
+	for (const e of rawEdits) {
+		if (!e || typeof e !== 'object' || Array.isArray(e)) continue;
+		const obj = e as Record<string, unknown>;
+		const oldText = typeof obj.old_text === 'string' ? obj.old_text : '';
+		if (!oldText) continue;
+		const newText = typeof obj.new_text === 'string' ? obj.new_text : '';
+		edits.push({ oldText, newText });
+	}
+
+	const resultObj = tryParseToolResultObject(section.toolResult);
+	let resultMessage: string | undefined;
+	let editsApplied: number | undefined;
+	let errorMessage: string | undefined;
+	if (typeof resultObj?.error === 'string') {
+		errorMessage = resultObj.error;
+	} else if (resultObj) {
+		if (typeof resultObj.result === 'string') {
+			resultMessage = resultObj.result;
+		}
+		if (Number.isFinite(Number(resultObj.edits_applied))) {
+			editsApplied = Number(resultObj.edits_applied);
+		}
+	}
+
+	return {
+		fileName,
+		filePath: rawPath,
+		edits,
+		resultMessage,
+		editsApplied,
+		errorMessage
+	};
+}

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

@@ -0,0 +1,23 @@
+// Meta parser for `exec_shell_command` tool calls. Surfaces the
+// command text from args `command` / `cmd` / `shell_command` aliases.
+// The exit-status and error parsing live in their own utilities
+// (`parse-exec-shell-status.ts` / `parse-exec-shell-error.ts`) - this
+// file only deals with what's strictly about *calling* the tool, since
+// the error / exit status elide from call-section to result-section.
+
+import { BuiltInTool } from '$lib/enums';
+import type { AgenticSection } from '$lib/utils';
+import { parseToolArgs } from './_shared';
+
+export type ExecShellCommandMeta = {
+	command: string;
+};
+
+export function parseExecShellCommandMeta(section: AgenticSection): ExecShellCommandMeta | null {
+	const args = parseToolArgs(BuiltInTool.EXEC_SHELL_COMMAND, section);
+	if (!args) return null;
+
+	const commandRaw = args.command ?? args.cmd ?? args.shell_command;
+	if (typeof commandRaw !== 'string' || !commandRaw) return null;
+	return { command: commandRaw };
+}

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

@@ -0,0 +1,58 @@
+// Meta parser for `file_glob_search` tool calls. Reads the path,
+// include pattern, and optional exclude from the args (strict parsing)
+// and the matches from the result blob. Like grep_search, the result
+// parser keeps the original raw-text fallback for MCP servers that
+// emit unparseable output.
+
+import { BuiltInTool } from '$lib/enums';
+import { splitSearchSummaryList, type AgenticSection } from '$lib/utils';
+import { parseToolArgs } from './_shared';
+
+export type FileGlobSearchMeta = {
+	path: string;
+	include: string;
+	exclude?: string;
+	matches: string[];
+	totalMatches?: number;
+	errorMessage?: string;
+};
+
+export function parseFileGlobSearchMeta(section: AgenticSection): FileGlobSearchMeta | null {
+	const args = parseToolArgs(BuiltInTool.FILE_GLOB_SEARCH, section);
+	if (!args) return null;
+
+	const path = typeof args.path === 'string' ? args.path : '';
+	const include = typeof args.include === 'string' && args.include ? args.include : '**';
+	const exclude = typeof args.exclude === 'string' && args.exclude ? args.exclude : undefined;
+	if (!path) return null;
+
+	let matches: string[] = [];
+	let totalMatches: number | undefined;
+	let errorMessage: string | undefined;
+
+	const toolResultString = section.toolResult;
+	if (toolResultString) {
+		try {
+			const parsed: unknown = JSON.parse(toolResultString);
+			if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) {
+				const obj = parsed as Record<string, unknown>;
+				if (typeof obj.error === 'string') {
+					errorMessage = obj.error;
+				} else if (typeof obj.plain_text_response === 'string') {
+					const split = splitSearchSummaryList(obj.plain_text_response, (total) => {
+						totalMatches = total;
+					});
+					matches = split.lines;
+				}
+			}
+		} catch {
+			// See grep-search.ts: same fallback used there.
+			const split = splitSearchSummaryList(toolResultString, (total) => {
+				totalMatches = total;
+			});
+			matches = split.lines;
+		}
+	}
+
+	return { path, include, exclude, matches, totalMatches, errorMessage };
+}

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

@@ -0,0 +1,108 @@
+// Meta parser for `grep_search` tool calls. Reads the path/pattern
+// triplet from args (strict parsing - we wait for the args to
+// complete) and the matches from the result blob. The result parser
+// keeps the original "scan result as raw text on JSON.parse failure"
+// fallback so MCP servers that return unparseable output still get
+// surfaced.
+
+import { BuiltInTool } from '$lib/enums';
+import { splitSearchSummaryList, type AgenticSection } from '$lib/utils';
+import { parseToolArgs } from './_shared';
+
+export type GrepSearchMatch = {
+	file: string;
+	line?: number;
+	content: string;
+};
+
+export type GrepSearchMeta = {
+	path: string;
+	pattern: string;
+	include: string;
+	exclude?: string;
+	showLineNumbers: boolean;
+	matches: GrepSearchMatch[];
+	totalMatches?: number;
+	errorMessage?: string;
+};
+
+export function parseGrepSearchMeta(section: AgenticSection): GrepSearchMeta | null {
+	const args = parseToolArgs(BuiltInTool.GREP_SEARCH, section);
+	if (!args) return null;
+
+	const path = typeof args.path === 'string' ? args.path : '';
+	const pattern = typeof args.pattern === 'string' ? args.pattern : '';
+	if (!path || !pattern) return null;
+
+	const include = typeof args.include === 'string' && args.include ? args.include : '**';
+	const exclude = typeof args.exclude === 'string' && args.exclude ? args.exclude : undefined;
+	const showLineNumbers = args.return_line_numbers === true;
+
+	let matches: GrepSearchMatch[] = [];
+	let totalMatches: number | undefined;
+	let errorMessage: string | undefined;
+
+	const toolResultString = section.toolResult;
+	if (toolResultString) {
+		try {
+			const parsed: unknown = JSON.parse(toolResultString);
+			if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) {
+				const obj = parsed as Record<string, unknown>;
+				if (typeof obj.error === 'string') {
+					errorMessage = obj.error;
+				} else if (typeof obj.plain_text_response === 'string') {
+					const split = splitSearchSummaryList(obj.plain_text_response, (total) => {
+						totalMatches = total;
+					});
+					matches = split.lines.map((line) => parseGrepLine(line, showLineNumbers));
+				}
+			}
+		} catch {
+			// Result wasn't JSON: keep behaviour for MCP servers that
+			// emit raw text and treat each line as a `<file>:<content>`
+			// (or `<file>:<line>:<content>`) match.
+			const split = splitSearchSummaryList(toolResultString, (total) => {
+				totalMatches = total;
+			});
+			matches = split.lines.map((line) => parseGrepLine(line, showLineNumbers));
+		}
+	}
+
+	return {
+		path,
+		pattern,
+		include,
+		exclude,
+		showLineNumbers,
+		matches,
+		totalMatches,
+		errorMessage
+	};
+}
+
+function parseGrepLine(line: string, showLineNumbers: boolean): GrepSearchMatch {
+	// Server output:
+	//   <file>:<content>          when return_line_numbers=false
+	//   <file>:<lineno>:<content> when return_line_numbers=true
+	const firstColon = line.indexOf(':');
+	if (firstColon === -1) {
+		return { file: line, content: '' };
+	}
+	const file = line.slice(0, firstColon);
+	const tail = line.slice(firstColon + 1);
+
+	if (!showLineNumbers) {
+		return { file, content: tail };
+	}
+
+	const secondColon = tail.indexOf(':');
+	if (secondColon === -1) {
+		return { file, content: tail };
+	}
+	const lineNum = parseInt(tail.slice(0, secondColon), 10);
+	return {
+		file,
+		line: Number.isFinite(lineNum) ? lineNum : undefined,
+		content: tail.slice(secondColon + 1)
+	};
+}

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

@@ -0,0 +1,52 @@
+// Meta parser for `read_file` tool calls. Reads the file path and an
+// optional line range (either `start_line`+`end_line` or
+// `start_line`+`line_count`). Args are parsed partially so a header
+// can render incrementally as the file path streams in.
+
+import { BuiltInTool } from '$lib/enums';
+import {
+	DEFAULT_LANGUAGE,
+	FILE_PATH_SEPARATOR_REGEX,
+	TEXT_LANGUAGE_PREFIX_REGEX
+} from '$lib/constants';
+import { getFileTypeByExtension, type AgenticSection } from '$lib/utils';
+import { parseToolArgs } from './_shared';
+
+export type ReadFileMeta = {
+	fileName: string;
+	lineRange: { start: number; end: number } | null;
+	language: string;
+};
+
+export function parseReadFileMeta(section: AgenticSection): ReadFileMeta | null {
+	const args = parseToolArgs(BuiltInTool.READ_FILE, section, { partial: true });
+	if (!args) return null;
+
+	const rawPath = args.path ?? args.file_path ?? args.filePath;
+	if (typeof rawPath !== 'string' || !rawPath) return null;
+
+	const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath;
+
+	// Models emit range arguments under several aliases. Accept all to
+	// stay forgiving across prompt variations.
+	const startRaw = args.start_line ?? args.line_start ?? args.startLine ?? args.from_line;
+	const endRaw = args.end_line ?? args.line_end ?? args.endLine ?? args.to_line;
+	const countRaw = args.line_count ?? args.count ?? args.num_lines;
+
+	let lineRange: { start: number; end: number } | null = null;
+	const sNum = Number(startRaw);
+	const eNum = Number(endRaw);
+	if (startRaw != null && endRaw != null && Number.isFinite(sNum) && Number.isFinite(eNum)) {
+		lineRange = { start: sNum, end: eNum };
+	} else if (startRaw != null && countRaw != null) {
+		const cNum = Number(countRaw);
+		if (Number.isFinite(sNum) && Number.isFinite(cNum)) {
+			lineRange = { start: sNum, end: sNum + cNum - 1 };
+		}
+	}
+
+	const fileType = getFileTypeByExtension(fileName);
+	const language = fileType ? fileType.replace(TEXT_LANGUAGE_PREFIX_REGEX, '') : DEFAULT_LANGUAGE;
+
+	return { fileName, lineRange, language };
+}

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

@@ -0,0 +1,56 @@
+// Meta parser for `run_javascript` tool calls. Reads the JS code and
+// optional timeout from args (strict parsing) and surfaces any error
+// from the result blob. SandboxService.formatReply emits a JSON object
+// containing an `error` field on failure, but a partial/non-JSON
+// failure renders as a flat line beginning with `Error:`. Both shapes
+// are handled.
+
+import { BuiltInTool } from '$lib/enums';
+import type { AgenticSection } from '$lib/utils';
+import { parseToolArgs } from './_shared';
+
+export type RunJavascriptMeta = {
+	code: string;
+	timeoutMs?: number;
+	errorMessage?: string;
+};
+
+export function parseRunJavascriptMeta(section: AgenticSection): RunJavascriptMeta | null {
+	const args = parseToolArgs(BuiltInTool.RUN_JAVASCRIPT, section);
+	if (!args) return null;
+
+	const code = typeof args.code === 'string' ? args.code : '';
+	if (!code) return null;
+
+	const timeoutRaw = Number(args.timeout_ms);
+	const timeoutMs = Number.isFinite(timeoutRaw) && timeoutRaw > 0 ? timeoutRaw : undefined;
+
+	let errorMessage: string | undefined;
+	const toolResultString = section.toolResult;
+	if (toolResultString) {
+		// Branches matter here: a JSON object can carry `error`, but a
+		// JSON array always represents successful output (sandbox returns
+		// the array of values). Only when the result isn't a JSON object
+		// do we scan raw lines for the `Error:` prefix.
+		let parsedObject: Record<string, unknown> | null = null;
+		try {
+			const parsed: unknown = JSON.parse(toolResultString);
+			if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) {
+				parsedObject = parsed as Record<string, unknown>;
+			}
+		} catch {
+			parsedObject = null;
+		}
+		if (typeof parsedObject?.error === 'string') {
+			errorMessage = parsedObject.error;
+		} else if (!parsedObject) {
+			const errorLine = toolResultString
+				.split('\n')
+				.map((line) => line.trim())
+				.find((line) => line.startsWith('Error:'));
+			if (errorLine) errorMessage = errorLine.slice('Error:'.length).trim();
+		}
+	}
+
+	return { code, timeoutMs, errorMessage };
+}

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

@@ -0,0 +1,54 @@
+// Meta parser for `write_file` tool calls. Reads the path/content from
+// the streamed args (partial JSON so we can render before the call
+// finishes) and surfaces `bytes`, `result`, and `error` from the
+// result blob.
+
+import { BuiltInTool } from '$lib/enums';
+import {
+	DEFAULT_LANGUAGE,
+	FILE_PATH_SEPARATOR_REGEX,
+	TEXT_LANGUAGE_PREFIX_REGEX
+} from '$lib/constants';
+import { getFileTypeByExtension, tryParseToolResultObject, type AgenticSection } from '$lib/utils';
+import { parseToolArgs } from './_shared';
+
+export type WriteFileMeta = {
+	fileName: string;
+	filePath: string;
+	language: string;
+	content: string;
+	bytesWritten?: number;
+	resultMessage?: string;
+	errorMessage?: string;
+};
+
+export function parseWriteFileMeta(section: AgenticSection): WriteFileMeta | null {
+	const args = parseToolArgs(BuiltInTool.WRITE_FILE, section, { partial: true });
+	if (!args) return null;
+
+	// Tool contracts drifted over time: some models emit `path`,
+	// others `file_path` / `filePath`. Accept all three.
+	const rawPath = args.path ?? args.file_path ?? args.filePath;
+	if (typeof rawPath !== 'string' || !rawPath) return null;
+
+	const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath;
+	const content = typeof args.content === 'string' ? args.content : '';
+	const language =
+		getFileTypeByExtension(rawPath)?.replace(TEXT_LANGUAGE_PREFIX_REGEX, '') ?? DEFAULT_LANGUAGE;
+
+	const resultObj = tryParseToolResultObject(section.toolResult);
+	const bytesWritten =
+		resultObj && Number.isFinite(Number(resultObj.bytes)) ? Number(resultObj.bytes) : undefined;
+	const resultMessage = typeof resultObj?.result === 'string' ? resultObj.result : undefined;
+	const errorMessage = typeof resultObj?.error === 'string' ? resultObj.error : undefined;
+
+	return {
+		fileName,
+		filePath: rawPath,
+		language,
+		content,
+		bytesWritten,
+		resultMessage,
+		errorMessage
+	};
+}

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import type { Snippet, Component } from 'svelte';
 
 	interface Props {
@@ -12,7 +13,7 @@
 
 <div class="my-2 rounded-lg border border-border bg-card p-3">
 	<div class="mb-3 flex items-center gap-2 text-sm">
-		<IconComponent class="h-4 w-4 shrink-0 text-muted-foreground" />
+		<IconComponent class="{ICON_CLASS_DEFAULT} shrink-0 text-muted-foreground" />
 		<span>
 			{@render message()}
 		</span>

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

@@ -1,40 +1,29 @@
 <script lang="ts">
-	import { Wrench, Loader2, Brain } from '@lucide/svelte';
 	import {
 		ChatMessageStatistics,
-		CollapsibleContentBlock,
 		MarkdownContent,
-		SyntaxHighlightedCode,
 		ChatMessageActionCardPermissionRequest,
 		ChatMessageActionCardContinueRequest
 	} from '$lib/components/app';
 
-	import {
-		AgenticSectionType,
-		ChatMessageStatsView,
-		FileTypeText,
-		ToolPermissionDecision
-	} from '$lib/enums';
+	import { AgenticSectionType, ChatMessageStatsView, ToolPermissionDecision } from '$lib/enums';
 	import type {
 		ChatMessageAgenticTimings,
 		ChatMessageAgenticTurnStats,
 		DatabaseMessage
 	} from '$lib/types';
-	import {
-		deriveAgenticSections,
-		formatJsonPretty,
-		parseToolResultWithImages,
-		type AgenticSection,
-		type ToolResultLine
-	} from '$lib/utils';
+	import { deriveAgenticSections, type AgenticSection } from '$lib/utils';
 	import {
 		agenticPendingPermissionRequest,
 		agenticResolvePermission,
 		agenticPendingContinueRequest,
 		agenticResolveContinue,
-		agenticLastError
+		agenticLastError,
+		agenticExecutingToolCallId
 	} from '$lib/stores/agentic.svelte';
 	import { config } from '$lib/stores/settings.svelte';
+	import ChatMessageReasoningBlock from './ChatMessageReasoningBlock.svelte';
+	import ChatMessageToolCallBlock from './ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte';
 
 	interface Props {
 		message: DatabaseMessage;
@@ -52,10 +41,10 @@
 
 	let expandedStates: Record<number, boolean> = $state({});
 
-	const showToolCallInProgress = $derived(config().showToolCallInProgress as boolean);
-	const showThoughtInProgress = $derived(config().showThoughtInProgress as boolean);
 	const renderThinkingAsMarkdown = $derived(config().renderThinkingAsMarkdown as boolean);
-	const showMessageStats = $derived(config().showMessageStats as boolean);
+	const showThoughtInProgress = $derived(Boolean(config().showThoughtInProgress));
+	const showMessageStats = $derived(Boolean(config().showMessageStats));
+	const showAgenticTurnStats = $derived(showMessageStats && Boolean(config().showAgenticTurnStats));
 
 	const hasReasoningError = $derived(
 		isLastAssistantMessage ? !!agenticLastError(message.convId) : false
@@ -67,7 +56,6 @@
 		isStreaming && isLastAssistantMessage ? agenticPendingPermissionRequest(message.convId) : null
 	);
 
-	// Reset dismissed when pendingPermission changes (new request or cleared)
 	let prevPendingRef: typeof pendingPermission = null;
 	$effect(() => {
 		if (pendingPermission !== prevPendingRef) {
@@ -106,33 +94,43 @@
 
 	const sections = $derived(deriveAgenticSections(message, toolMessages, [], isStreaming));
 
-	// Parse tool results with images
-	const sectionsParsed = $derived(
-		sections.map((section) => ({
-			...section,
-			parsedLines: section.toolResult
-				? parseToolResultWithImages(section.toolResult, section.toolResultExtras || message?.extra)
-				: ([] as ToolResultLine[])
-		}))
+	const currentlyExecutingToolCallId = $derived(
+		isStreaming ? agenticExecutingToolCallId(message.convId) : null
 	);
 
-	// Group flat sections into agentic turns
-	// A new turn starts when a non-tool section follows a tool section
-	const turnGroups = $derived.by(() => {
-		const turns: { sections: (typeof sectionsParsed)[number][]; flatIndices: number[] }[] = [];
-		let currentTurn: (typeof sectionsParsed)[number][] = [];
+	// Skip sections the user manually collapsed - we never override an explicit false.
+	let lastSeenExecutingToolCallId: string | null = null;
+	$effect(() => {
+		const current = currentlyExecutingToolCallId;
+		const previous = lastSeenExecutingToolCallId;
+		lastSeenExecutingToolCallId = current;
+		if (!current || current === previous) return;
+		const idx = sections.findIndex((s) => s.toolCallId === current);
+		if (idx >= 0 && expandedStates[idx] === undefined) {
+			expandedStates[idx] = true;
+		}
+	});
+
+	type TurnGroup = {
+		sections: AgenticSection[];
+		flatIndices: number[];
+	};
+
+	const turnGroups: TurnGroup[] = $derived.by(() => {
+		const groups: TurnGroup[] = [];
+		let currentTurn: AgenticSection[] = [];
 		let currentIndices: number[] = [];
 		let prevWasTool = false;
 
-		for (let i = 0; i < sectionsParsed.length; i++) {
-			const section = sectionsParsed[i];
+		for (let i = 0; i < sections.length; i++) {
+			const section = sections[i];
 			const isTool =
 				section.type === AgenticSectionType.TOOL_CALL ||
 				section.type === AgenticSectionType.TOOL_CALL_PENDING ||
 				section.type === AgenticSectionType.TOOL_CALL_STREAMING;
 
 			if (!isTool && prevWasTool && currentTurn.length > 0) {
-				turns.push({ sections: currentTurn, flatIndices: currentIndices });
+				groups.push({ sections: currentTurn, flatIndices: currentIndices });
 				currentTurn = [];
 				currentIndices = [];
 			}
@@ -143,10 +141,10 @@
 		}
 
 		if (currentTurn.length > 0) {
-			turns.push({ sections: currentTurn, flatIndices: currentIndices });
+			groups.push({ sections: currentTurn, flatIndices: currentIndices });
 		}
 
-		return turns;
+		return groups;
 	});
 
 	function getDefaultExpanded(section: AgenticSection): boolean {
@@ -154,7 +152,7 @@
 			section.type === AgenticSectionType.TOOL_CALL_PENDING ||
 			section.type === AgenticSectionType.TOOL_CALL_STREAMING
 		) {
-			return showToolCallInProgress;
+			return false;
 		}
 
 		if (section.type === AgenticSectionType.REASONING_PENDING) {
@@ -189,181 +187,46 @@
 	}
 </script>
 
-{#snippet renderSection(section: (typeof sectionsParsed)[number], index: number)}
+{#snippet renderSection(section: AgenticSection, index: number)}
 	{#if section.type === AgenticSectionType.TEXT}
 		<div class="agentic-text">
 			<MarkdownContent content={section.content} attachments={message?.extra} />
 		</div>
-	{:else if section.type === AgenticSectionType.TOOL_CALL_STREAMING}
-		{@const streamingIcon = isStreaming ? Loader2 : Loader2}
-		{@const streamingIconClass = isStreaming ? 'h-4 w-4 animate-spin' : 'h-4 w-4'}
-
-		<CollapsibleContentBlock
+	{:else if section.type === AgenticSectionType.REASONING || section.type === AgenticSectionType.REASONING_PENDING}
+		<ChatMessageReasoningBlock
+			{section}
 			open={isExpanded(index, section)}
-			class="my-2"
-			icon={streamingIcon}
-			iconClass={streamingIconClass}
-			title={section.toolName || 'Tool call'}
-			subtitle={isStreaming ? '' : 'incomplete'}
 			{isStreaming}
+			{renderThinkingAsMarkdown}
+			{hasReasoningError}
+			attachments={message?.extra}
 			onToggle={() => toggleExpanded(index, section)}
-		>
-			<div class="pt-3">
-				<div class="my-3 flex items-center gap-2 text-xs text-muted-foreground">
-					<span>Arguments:</span>
-
-					{#if isStreaming}
-						<Loader2 class="h-3 w-3 animate-spin" />
-					{/if}
-				</div>
-				{#if section.toolArgs}
-					<SyntaxHighlightedCode
-						code={formatJsonPretty(section.toolArgs)}
-						language={FileTypeText.JSON}
-						maxHeight="20rem"
-						class="text-xs"
-					/>
-				{:else if isStreaming}
-					<div class="rounded bg-muted/30 p-2 text-xs text-muted-foreground italic">
-						Receiving arguments...
-					</div>
-				{:else}
-					<div
-						class="rounded bg-yellow-500/10 p-2 text-xs text-yellow-600 italic dark:text-yellow-400"
-					>
-						Response was truncated
-					</div>
-				{/if}
-			</div>
-		</CollapsibleContentBlock>
-	{:else if section.type === AgenticSectionType.TOOL_CALL || section.type === AgenticSectionType.TOOL_CALL_PENDING}
-		{@const isPending = section.type === AgenticSectionType.TOOL_CALL_PENDING}
-		{@const toolIcon = isPending ? Loader2 : Wrench}
-		{@const toolIconClass = isPending ? 'h-4 w-4 animate-spin' : 'h-4 w-4'}
-
-		<CollapsibleContentBlock
+		/>
+	{:else if section.type === AgenticSectionType.TOOL_CALL || section.type === AgenticSectionType.TOOL_CALL_PENDING || section.type === AgenticSectionType.TOOL_CALL_STREAMING}
+		<ChatMessageToolCallBlock
+			{section}
 			open={isExpanded(index, section)}
-			class="my-2"
-			icon={toolIcon}
-			iconClass={toolIconClass}
-			title={section.toolName || ''}
-			subtitle={isPending ? 'executing...' : undefined}
-			isStreaming={isPending}
-			onToggle={() => toggleExpanded(index, section)}
-		>
-			{#if section.toolArgs && section.toolArgs !== '{}'}
-				<div class="pt-3">
-					<div class="my-3 text-xs text-muted-foreground">Arguments:</div>
-
-					<SyntaxHighlightedCode
-						code={formatJsonPretty(section.toolArgs)}
-						language={FileTypeText.JSON}
-						maxHeight="20rem"
-						class="text-xs"
-					/>
-				</div>
-			{/if}
-
-			<div class="pt-3">
-				<div class="my-3 flex items-center gap-2 text-xs text-muted-foreground">
-					<span>Result:</span>
-
-					{#if isPending}
-						<Loader2 class="h-3 w-3 animate-spin" />
-					{/if}
-				</div>
-				{#if isPending}
-					<div class="rounded bg-muted/30 p-2 text-xs text-muted-foreground italic">
-						Waiting for result...
-					</div>
-				{:else if section.toolResult}
-					<div class="overflow-auto rounded-lg border border-border bg-muted p-4">
-						{#each section.parsedLines as line, i (i)}
-							<div class="font-mono text-xs leading-relaxed whitespace-pre-wrap">
-								{line.text}
-							</div>
-							{#if line.image}
-								<img
-									src={line.image.base64Url}
-									alt={line.image.name}
-									class="mt-2 mb-2 h-auto max-w-full rounded-lg"
-									loading="lazy"
-								/>
-							{/if}
-						{/each}
-					</div>
-				{:else}
-					<div class="rounded bg-muted/30 p-2 text-xs text-muted-foreground italic">No output</div>
-				{/if}
-			</div>
-		</CollapsibleContentBlock>
-	{:else if section.type === AgenticSectionType.REASONING}
-		{@const reasoningSubtitle = section.wasInterrupted
-			? hasReasoningError
-				? 'Error'
-				: 'Cancelled'
-			: isStreaming
-				? ''
-				: undefined}
-
-		<CollapsibleContentBlock
-			open={isExpanded(index, section)}
-			class="my-2"
-			icon={Brain}
-			title="Reasoning"
-			subtitle={reasoningSubtitle}
-			rawContent={section.content}
-			onToggle={() => toggleExpanded(index, section)}
-		>
-			<div class="pt-3">
-				{#if renderThinkingAsMarkdown}
-					<MarkdownContent content={section.content} attachments={message?.extra} />
-				{:else}
-					<div class="text-xs leading-relaxed break-words whitespace-pre-wrap">
-						{section.content}
-					</div>
-				{/if}
-			</div>
-		</CollapsibleContentBlock>
-	{:else if section.type === AgenticSectionType.REASONING_PENDING}
-		{@const reasoningTitle = isStreaming ? 'Reasoning...' : 'Reasoning'}
-		{@const reasoningSubtitle = isStreaming ? '' : hasReasoningError ? 'Error' : 'Cancelled'}
-
-		<CollapsibleContentBlock
-			open={isExpanded(index, section)}
-			class="my-2"
-			icon={Brain}
-			title={reasoningTitle}
-			subtitle={reasoningSubtitle}
-			rawContent={section.content}
 			{isStreaming}
+			isExecuting={section.toolCallId !== undefined &&
+				section.toolCallId === currentlyExecutingToolCallId}
+			attachments={message?.extra}
 			onToggle={() => toggleExpanded(index, section)}
-		>
-			<div class="pt-3">
-				{#if renderThinkingAsMarkdown}
-					<MarkdownContent content={section.content} attachments={message?.extra} />
-				{:else}
-					<div class="text-xs leading-relaxed break-words whitespace-pre-wrap">
-						{section.content}
-					</div>
-				{/if}
-			</div>
-		</CollapsibleContentBlock>
+		/>
 	{/if}
 {/snippet}
 
-<div class="agentic-content">
+<div class="agentic-content gap-2">
 	{#if turnGroups.length > 1}
 		{#each turnGroups as turn, turnIndex (turnIndex)}
 			{@const turnStats = message?.timings?.agentic?.perTurn?.[turnIndex]}
 
-			<div class="agentic-turn group/turn grid gap-3 mb-4">
+			<div class="agentic-turn group/turn grid gap-2">
 				{#each turn.sections as section, sIdx (turn.flatIndices[sIdx])}
 					{@render renderSection(section, turn.flatIndices[sIdx])}
 				{/each}
 
-				{#if turnStats && showMessageStats}
-					<div class="turn-stats transition-opacity duration-150">
+				{#if turnStats && showAgenticTurnStats}
+					<div class="turn-stats transition-opacity duration-150 mt-1 mb-4">
 						<ChatMessageStatistics
 							promptTokens={turnStats.llm.prompt_n}
 							promptMs={turnStats.llm.prompt_ms}
@@ -380,7 +243,7 @@
 			</div>
 		{/each}
 	{:else}
-		{#each sectionsParsed as section, index (index)}
+		{#each sections as section, index (index)}
 			{@render renderSection(section, index)}
 		{/each}
 	{/if}
@@ -404,7 +267,6 @@
 		flex-direction: column;
 		width: 100%;
 		max-width: 48rem;
-		gap: 1rem;
 	}
 
 	.agentic-content > :global(*),

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

@@ -0,0 +1,151 @@
+<script lang="ts">
+	import { Lightbulb } from '@lucide/svelte';
+	import { CollapsibleContentBlock, MarkdownContent } from '$lib/components/app';
+	import { AgenticSectionType } from '$lib/enums';
+	import { REASONING_SCROLL_AT_BOTTOM_THRESHOLD_PX } from '$lib/constants/auto-scroll';
+	import type { DatabaseMessageExtra } from '$lib/types';
+	import type { AgenticSection } from '$lib/utils';
+
+	interface Props {
+		section: AgenticSection;
+		open: boolean;
+		isStreaming: boolean;
+		renderThinkingAsMarkdown: boolean;
+		hasReasoningError?: boolean;
+		attachments?: DatabaseMessageExtra[];
+		onToggle?: () => void;
+	}
+
+	let {
+		section,
+		open,
+		isStreaming,
+		renderThinkingAsMarkdown,
+		hasReasoningError = false,
+		attachments,
+		onToggle
+	}: Props = $props();
+
+	const REASONING_HEADER = 'Reasoning';
+	const REASONING_HEADER_PENDING = 'Reasoning...';
+	const REASONING_SUBTITLE_ERROR = 'Error';
+	const REASONING_SUBTITLE_CANCELLED = 'Cancelled';
+
+	const isPending = $derived(section.type === AgenticSectionType.REASONING_PENDING);
+	const title = $derived(isPending && isStreaming ? REASONING_HEADER_PENDING : REASONING_HEADER);
+	const subtitle = $derived.by(() => {
+		if (isPending && !isStreaming) {
+			return hasReasoningError ? REASONING_SUBTITLE_ERROR : REASONING_SUBTITLE_CANCELLED;
+		}
+		if (section.wasInterrupted) {
+			return hasReasoningError ? REASONING_SUBTITLE_ERROR : REASONING_SUBTITLE_CANCELLED;
+		}
+		return isStreaming ? '' : undefined;
+	});
+	const shimmerTitle = $derived(isPending && isStreaming);
+
+	let scrollEl: HTMLDivElement | undefined = $state();
+
+	const SCROLL_BOTTOM_THRESHOLD_PX = REASONING_SCROLL_AT_BOTTOM_THRESHOLD_PX;
+
+	let userScrolledUp = $state(false);
+	let lastScrollTop = 0;
+	let pendingFrame: number | null = null;
+
+	function isAtBottom(): boolean {
+		if (!scrollEl) return false;
+		return (
+			scrollEl.scrollHeight - scrollEl.clientHeight - scrollEl.scrollTop <=
+			SCROLL_BOTTOM_THRESHOLD_PX
+		);
+	}
+
+	function scrollToBottomOnFrame() {
+		if (pendingFrame !== null || !scrollEl || userScrolledUp) return;
+		pendingFrame = requestAnimationFrame(() => {
+			pendingFrame = null;
+			// User may scroll between scheduling and paint.
+			if (scrollEl && !userScrolledUp) {
+				scrollEl.scrollTop = scrollEl.scrollHeight;
+			}
+		});
+	}
+
+	function handleScrollEvent() {
+		if (!scrollEl) return;
+		const isScrollingUp = scrollEl.scrollTop < lastScrollTop;
+		if (isScrollingUp && !isAtBottom()) {
+			userScrolledUp = true;
+		} else if (isAtBottom()) {
+			userScrolledUp = false;
+		}
+		lastScrollTop = scrollEl.scrollTop;
+	}
+
+	$effect(() => {
+		void section.content;
+		if (!scrollEl || !isPending || !isStreaming) return;
+		scrollToBottomOnFrame();
+	});
+
+	$effect(() => {
+		// Layout shifts that don't change section.content (markdown re-parse,
+		// syntax-highlight settle, image loads).
+		if (!scrollEl || !isPending || !isStreaming) return;
+
+		const observer = new MutationObserver(() => scrollToBottomOnFrame());
+		observer.observe(scrollEl, {
+			childList: true,
+			subtree: true,
+			characterData: true
+		});
+
+		return () => observer.disconnect();
+	});
+
+	$effect(() => {
+		// Pin to bottom at the start of each round.
+		if (!isPending) {
+			userScrolledUp = false;
+			lastScrollTop = 0;
+		}
+	});
+</script>
+
+<CollapsibleContentBlock
+	{open}
+	class="my-2"
+	icon={Lightbulb}
+	iconClass="h-3.5 w-3.5"
+	{title}
+	{subtitle}
+	{shimmerTitle}
+	{onToggle}
+>
+	<div
+		bind:this={scrollEl}
+		class="reasoning-content"
+		class:is-streaming={isPending}
+		onscroll={handleScrollEvent}
+	>
+		{#if renderThinkingAsMarkdown}
+			<MarkdownContent content={section.content} class="text-muted-foreground" {attachments} />
+		{:else}
+			<div
+				class="text-[13px] leading-relaxed wrap-break-word whitespace-pre-wrap text-muted-foreground"
+			>
+				{section.content}
+			</div>
+		{/if}
+	</div>
+</CollapsibleContentBlock>
+
+<style>
+	.reasoning-content.is-streaming {
+		max-height: 28rem;
+		overflow-y: auto;
+		overscroll-behavior: contain;
+		scrollbar-gutter: stable;
+		padding-right: 0.25rem;
+	}
+</style>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { ArrowDown } from '@lucide/svelte';
 	import ActionIcon from '$lib/components/app/actions/ActionIcon.svelte';
 
@@ -12,7 +13,7 @@
 		ariaLabel="Scroll to bottom"
 		tooltip="Scroll to bottom"
 		size="lg"
-		iconSize="h-4 w-4"
+		iconSize={ICON_CLASS_DEFAULT}
 		class="h-9 w-9 rounded-full bg-accent text-accent-foreground absolute bottom-4 shadow-md"
 	/>
 </div>

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { AlertTriangle, Loader2, RefreshCw } from '@lucide/svelte';
 	import { fadeInView } from '$lib/actions/fade-in-view.svelte';
 	import * as Alert from '$lib/components/ui/alert';
@@ -15,9 +16,9 @@
 	>
 		<Alert.Root variant={isLoadingModel ? 'default' : 'destructive'}>
 			{#if isLoadingModel}
-				<Loader2 class="h-4 w-4 animate-spin" />
+				<Loader2 class="{ICON_CLASS_DEFAULT} animate-spin" />
 			{:else}
-				<AlertTriangle class="h-4 w-4" />
+				<AlertTriangle class={ICON_CLASS_DEFAULT} />
 			{/if}
 
 			<Alert.Title class="flex items-center justify-between">

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

@@ -571,6 +571,10 @@ export { default as ChatMessageMcpPromptContent } from './ChatMessages/ChatMessa
  * Handles streaming state with real-time content updates.
  */
 export { default as ChatMessageAssistant } from './ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte';
+export { default as ChatMessageAssistantModel } from './ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantModel.svelte';
+export { default as ChatMessageAssistantProcessingInfo } from './ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte';
+export { default as ChatMessageAssistantRawOutput } from './ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte';
+export { default as ChatMessageAssistantStatistics } from './ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte';
 
 /**
  * Inline message editing form. Provides textarea for editing message content with

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

@@ -1,12 +1,8 @@
 <script lang="ts">
-	import ChevronsUpDownIcon from '@lucide/svelte/icons/chevrons-up-down';
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
+	import ChevronDown from '@lucide/svelte/icons/chevron-down';
 	import * as Collapsible from '$lib/components/ui/collapsible/index.js';
-	import { buttonVariants } from '$lib/components/ui/button/index.js';
-	import { Card } from '$lib/components/ui/card';
-	import { createAutoScrollController } from '$lib/hooks/use-auto-scroll.svelte';
-	import { useThrottle } from '$lib/hooks/use-throttle.svelte';
-	import { formatReasoningPreview } from '$lib/utils';
-	import { config } from '$lib/stores/settings.svelte';
+	import { cn } from '$lib/components/ui/utils';
 	import type { Snippet } from 'svelte';
 	import type { Component } from 'svelte';
 
@@ -15,11 +11,11 @@
 		class?: string;
 		icon?: Component;
 		iconClass?: string;
-		title: string;
+		iconUrl?: string | null;
+		title?: string;
+		titleSnippet?: Snippet;
 		subtitle?: string;
-		preview?: string;
-		rawContent?: string;
-		isStreaming?: boolean;
+		shimmerTitle?: boolean;
 		onToggle?: () => void;
 		children: Snippet;
 	}
@@ -28,45 +24,18 @@
 		open = $bindable(false),
 		class: className = '',
 		icon: IconComponent,
-		iconClass = 'h-4 w-4',
-		title,
+		iconClass = ICON_CLASS_DEFAULT,
+		iconUrl = null,
+		title = '',
+		titleSnippet,
 		subtitle,
-		preview,
-		rawContent,
-		isStreaming = false,
+		shimmerTitle = false,
 		onToggle,
 		children
 	}: Props = $props();
 
-	let contentContainer: HTMLDivElement | undefined = $state();
-
-	const showThoughtInProgress = $derived(config().showThoughtInProgress as boolean);
-
-	let previewKey = useThrottle(() => rawContent ?? preview ?? '', 500);
-	let displayedPreview = $state('');
-	let displayedOverflow = $state(0);
-
-	$effect(() => {
-		void previewKey.key;
-		const content = rawContent ?? preview ?? '';
-		const result = formatReasoningPreview(content);
-		displayedPreview = result.preview;
-		displayedOverflow = result.overflow;
-	});
-
-	const autoScroll = createAutoScrollController();
-
-	$effect(() => {
-		autoScroll.setContainer(contentContainer);
-	});
-
-	$effect(() => {
-		// Only auto-scroll when open and streaming
-		autoScroll.updateInterval(open && isStreaming);
-	});
-
-	function handleScroll() {
-		autoScroll.handleScroll();
+	function hideBrokenIcon(event: Event) {
+		(event.currentTarget as HTMLImageElement).style.display = 'none';
 	}
 </script>
 
@@ -76,59 +45,54 @@
 		open = value;
 		onToggle?.();
 	}}
-	class="{className} my-0!"
+	class={cn('group/collapsible', 'my-0!', className)}
 >
-	<Card class="gap-0 border-muted bg-muted/30 py-0">
-		<Collapsible.Trigger class="flex w-full cursor-pointer items-start justify-between gap-2 p-3">
-			<div class="flex min-w-0 items-center gap-2">
-				<div class="flex items-center gap-2 text-muted-foreground">
-					{#if IconComponent}
-						<IconComponent class={iconClass} />
-					{/if}
+	<Collapsible.Trigger
+		class={cn(
+			'flex w-full cursor-pointer items-start justify-between gap-2 text-left',
+			'py-1.5 pr-1'
+		)}
+	>
+		<div class="flex min-w-0 items-start gap-2 text-muted-foreground">
+			{#if iconUrl}
+				<img
+					src={iconUrl}
+					alt=""
+					class={cn('shrink-0 rounded-sm  mt-0.75', iconClass)}
+					onerror={hideBrokenIcon}
+				/>
+			{:else if IconComponent}
+				<IconComponent class={cn('shrink-0 text-muted-foreground/60 mt-0.75', iconClass)} />
+			{/if}
 
-					<span class="font-mono text-sm font-medium">{title}</span>
-
-					{#if subtitle}
-						<span class="text-xs italic">{subtitle}</span>
-					{/if}
-				</div>
-
-				{#if displayedPreview && !showThoughtInProgress}
-					<div class="flex min-w-0 items-baseline justify-between gap-2">
-						<div class="w-3/4 truncate text-xs text-muted-foreground/80">
-							{displayedPreview}
-						</div>
-						{#if displayedOverflow > 0}
-							<span class="shrink-0 text-xs text-muted-foreground/60"
-								>{displayedOverflow}+ chars</span
-							>
-						{/if}
-					</div>
+			<span class={cn('text-sm font-medium', shimmerTitle ? 'shimmer-text' : 'text-foreground/80')}>
+				{#if titleSnippet}
+					{@render titleSnippet()}
+				{:else}
+					{title}
 				{/if}
-			</div>
+			</span>
 
-			<div
-				class={buttonVariants({
-					variant: 'ghost',
-					size: 'sm',
-					class: 'h-6 w-6 p-0 text-muted-foreground hover:text-foreground'
-				})}
-			>
-				<ChevronsUpDownIcon class="h-4 w-4" />
+			{#if subtitle}
+				<span class="text-xs italic text-muted-foreground/70">{subtitle}</span>
+			{/if}
+		</div>
 
-				<span class="sr-only">Toggle content</span>
-			</div>
-		</Collapsible.Trigger>
+		<ChevronDown
+			class={cn(
+				'size-4 shrink-0 text-muted-foreground/60 transition-all duration-150 ease-out opacity-0 group-hover/collapsible:opacity-100 mt-0.75',
+				open && 'rotate-180'
+			)}
+		/>
 
-		<Collapsible.Content>
-			<div
-				bind:this={contentContainer}
-				class="overflow-y-auto border-t border-muted px-3 pb-3"
-				onscroll={handleScroll}
-				style="min-height: var(--min-message-height); max-height: var(--max-message-height);"
-			>
+		<span class="sr-only">Toggle content</span>
+	</Collapsible.Trigger>
+
+	<Collapsible.Content>
+		<div class="pl-1.5 grid min-w-0" style="min-height: var(--min-message-height);">
+			<div class="min-w-0 border-l border-muted-foreground/20 pl-4 pb-2 my-2">
 				{@render children()}
 			</div>
-		</Collapsible.Content>
-	</Card>
+		</div>
+	</Collapsible.Content>
 </Collapsible.Root>

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

@@ -0,0 +1,97 @@
+<script lang="ts">
+	import ChevronDown from '@lucide/svelte/icons/chevron-down';
+	import * as Collapsible from '$lib/components/ui/collapsible/index.js';
+	import { cn } from '$lib/components/ui/utils';
+	import { ICON_CLASS_DEFAULT } from '$lib/constants';
+	import type { Snippet } from 'svelte';
+	import type { Component } from 'svelte';
+
+	interface Props {
+		open?: boolean;
+		class?: string;
+		icon?: Component;
+		iconClass?: string;
+		iconUrl?: string | null;
+		title?: string;
+		titleSnippet?: Snippet;
+		subtitle?: string;
+		shimmerTitle?: boolean;
+		onToggle?: () => void;
+		children: Snippet;
+	}
+
+	let {
+		open = $bindable(false),
+		class: className = '',
+		icon: IconComponent,
+		iconClass = ICON_CLASS_DEFAULT,
+		iconUrl = null,
+		title = '',
+		titleSnippet,
+		subtitle,
+		shimmerTitle = false,
+		onToggle,
+		children
+	}: Props = $props();
+
+	function hideBrokenIcon(event: Event) {
+		(event.currentTarget as HTMLImageElement).style.display = 'none';
+	}
+</script>
+
+<Collapsible.Root
+	{open}
+	onOpenChange={(value) => {
+		open = value;
+		onToggle?.();
+	}}
+	class={cn('group/collapsible', 'overflow-hidden rounded-md', className)}
+	style="background: var(--code-background); border: 1px solid color-mix(in oklch, var(--border) 30%, transparent);"
+>
+	<Collapsible.Trigger
+		class={cn(
+			'flex w-full cursor-pointer items-start justify-between gap-2 text-left',
+			'px-3 py-2'
+		)}
+	>
+		<div class="flex min-w-0 items-start gap-2 text-muted-foreground">
+			{#if iconUrl}
+				<img
+					src={iconUrl}
+					alt=""
+					class={cn('shrink-0 rounded-sm mt-0.5', iconClass)}
+					onerror={hideBrokenIcon}
+				/>
+			{:else if IconComponent}
+				<IconComponent class={cn('shrink-0 text-muted-foreground/60 mt-0.5', iconClass)} />
+			{/if}
+
+			<span class={cn('text-sm font-medium', shimmerTitle ? 'shimmer-text' : 'text-foreground/80')}>
+				{#if titleSnippet}
+					{@render titleSnippet()}
+				{:else}
+					{title}
+				{/if}
+			</span>
+
+			{#if subtitle}
+				<span class="text-xs italic text-muted-foreground/70">{subtitle}</span>
+			{/if}
+		</div>
+
+		<ChevronDown
+			class={cn(
+				'size-4 shrink-0 text-muted-foreground/60 transition-all duration-150 ease-out opacity-0 group-hover/collapsible:opacity-100 mt-0.5',
+				open && 'rotate-180'
+			)}
+		/>
+
+		<span class="sr-only">Toggle content</span>
+	</Collapsible.Trigger>
+
+	<Collapsible.Content>
+		<div class="p-3 pt-1">
+			{@render children()}
+		</div>
+	</Collapsible.Content>
+</Collapsible.Root>

tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css

@@ -19,8 +19,16 @@
 	line-height: 1.75;
 }
 
+.markdown-content :global(.markdown-block:first-child p:first-child) {
+	margin-block-start: 0;
+}
+
+.markdown-content :global(.markdown-block:last-child p:last-child) {
+	margin-block-end: 0;
+}
+
 .markdown-content :global(:is(h1, h2, h3, h4, h5, h6):first-child) {
-	margin-top: 0;
+	margin-top: 0.5rem;
 }
 
 /* Headers with consistent spacing */

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Download } from '@lucide/svelte';
 	import ZoomInIcon from '@lucide/svelte/icons/zoom-in';
 	import ZoomOutIcon from '@lucide/svelte/icons/zoom-out';
@@ -36,7 +37,7 @@
 			title="Zoom out"
 			aria-label="Zoom out"
 		>
-			<ZoomOutIcon class="mermaid-preview-btn-icon h-4 w-4" />
+			<ZoomOutIcon class="mermaid-preview-btn-icon {ICON_CLASS_DEFAULT}" />
 		</button>
 		<span
 			class="mermaid-preview-zoom-label min-w-[3.5rem] px-0.5 text-center text-xs font-medium text-muted-foreground tabular-nums select-none"
@@ -48,7 +49,7 @@
 			title="Zoom in"
 			aria-label="Zoom in"
 		>
-			<ZoomInIcon class="mermaid-preview-btn-icon h-4 w-4" />
+			<ZoomInIcon class="mermaid-preview-btn-icon {ICON_CLASS_DEFAULT}" />
 		</button>
 		<div class="mermaid-preview-controls-separator mx-1 h-5 w-px bg-border/50"></div>
 
@@ -58,7 +59,7 @@
 			title="Reset view"
 			aria-label="Reset view"
 		>
-			<RotateCcwIcon class="mermaid-preview-btn-icon h-4 w-4" />
+			<RotateCcwIcon class="mermaid-preview-btn-icon {ICON_CLASS_DEFAULT}" />
 		</button>
 		<div class="mermaid-preview-controls-separator mx-1 h-5 w-px bg-border/50"></div>
 
@@ -68,7 +69,7 @@
 			title="Download SVG"
 			aria-label="Download SVG"
 		>
-			<Download class="mermaid-preview-btn-icon h-4 w-4" />
+			<Download class="mermaid-preview-btn-icon {ICON_CLASS_DEFAULT}" />
 		</button>
 	</div>
 </div>

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

@@ -1,11 +1,12 @@
 <script lang="ts">
-	import hljs from 'highlight.js';
 	import { browser } from '$app/environment';
 	import { mode } from 'mode-watcher';
 
 	import githubDarkCss from 'highlight.js/styles/github-dark.css?inline';
 	import githubLightCss from 'highlight.js/styles/github.css?inline';
 	import { ColorMode } from '$lib/enums';
+	import { SYNTAX_CODE_SCROLL_AT_BOTTOM_THRESHOLD_PX } from '$lib/constants/auto-scroll';
+	import { highlightCode } from '$lib/utils';
 
 	interface Props {
 		code: string;
@@ -13,6 +14,9 @@
 		class?: string;
 		maxHeight?: string;
 		maxWidth?: string;
+		/** Auto-scrolls to the bottom of new chunks; pauses on user scroll-up
+		 *  until the user returns to the bottom. */
+		streaming?: boolean;
 	}
 
 	let {
@@ -20,10 +24,17 @@
 		language = 'text',
 		class: className = '',
 		maxHeight = '60vh',
-		maxWidth = ''
+		maxWidth = '',
+		streaming = false
 	}: Props = $props();
 
-	let highlightedHtml = $state('');
+	const highlightedHtml = $derived(highlightCode(code, language));
+
+	let scrollEl = $state<HTMLDivElement>();
+	let userScrolledUp = $state(false);
+	let lastScrollTop = 0;
+	const SCROLL_BOTTOM_THRESHOLD_PX = SYNTAX_CODE_SCROLL_AT_BOTTOM_THRESHOLD_PX;
+	let pendingFrame: number | null = null;
 
 	function loadHighlightTheme(isDark: boolean) {
 		if (!browser) return;
@@ -38,6 +49,36 @@
 		document.head.appendChild(style);
 	}
 
+	function isAtBottom(): boolean {
+		if (!scrollEl) return false;
+		return (
+			scrollEl.scrollHeight - scrollEl.clientHeight - scrollEl.scrollTop <=
+			SCROLL_BOTTOM_THRESHOLD_PX
+		);
+	}
+
+	function scrollToBottomOnFrame() {
+		if (pendingFrame !== null || !scrollEl || userScrolledUp) return;
+		pendingFrame = requestAnimationFrame(() => {
+			pendingFrame = null;
+			// User may scroll between scheduling and paint.
+			if (scrollEl && !userScrolledUp) {
+				scrollEl.scrollTop = scrollEl.scrollHeight;
+			}
+		});
+	}
+
+	function handleScrollEvent() {
+		if (!scrollEl) return;
+		const isScrollingUp = scrollEl.scrollTop < lastScrollTop;
+		if (isScrollingUp && !isAtBottom()) {
+			userScrolledUp = true;
+		} else if (isAtBottom()) {
+			userScrolledUp = false;
+		}
+		lastScrollTop = scrollEl.scrollTop;
+	}
+
 	$effect(() => {
 		const currentMode = mode.current;
 		const isDark = currentMode === ColorMode.DARK;
@@ -45,46 +86,64 @@
 		loadHighlightTheme(isDark);
 	});
 
+	// Pin to bottom at the start of each streaming episode.
 	$effect(() => {
-		if (!code) {
-			highlightedHtml = '';
-			return;
+		if (streaming) {
+			userScrolledUp = false;
+			lastScrollTop = 0;
 		}
+	});
 
-		try {
-			// Check if the language is supported
-			const lang = language.toLowerCase();
-			const isSupported = hljs.getLanguage(lang);
+	$effect(() => {
+		void code;
+		if (!streaming || userScrolledUp) return;
+		scrollToBottomOnFrame();
+	});
 
-			if (isSupported) {
-				const result = hljs.highlight(code, { language: lang });
-				highlightedHtml = result.value;
-			} else {
-				// Try auto-detection or fallback to plain text
-				const result = hljs.highlightAuto(code);
-				highlightedHtml = result.value;
-			}
-		} catch {
-			// Fallback to escaped plain text
-			highlightedHtml = code.replace(/&/g, '&amp;').replace(/</g, '&lt;').replace(/>/g, '&gt;');
-		}
+	// Layout shifts that don't change `code` (highlight.js re-tokenize, line-wrap reflow).
+	$effect(() => {
+		if (!streaming || !scrollEl) return;
+
+		const observer = new MutationObserver(() => scrollToBottomOnFrame());
+		observer.observe(scrollEl, {
+			childList: true,
+			subtree: true,
+			characterData: true
+		});
+
+		return () => observer.disconnect();
 	});
 </script>
 
 <div
-	class="code-preview-wrapper min-w-0 max-w-full overflow-x-auto rounded-lg border border-border bg-muted {className}"
-	style="max-height: {maxHeight}; {maxWidth ? `max-width: ${maxWidth};` : ''}"
+	bind:this={scrollEl}
+	onscroll={handleScrollEvent}
+	class="code-preview-wrapper min-w-0 max-w-full overflow-auto rounded-xl border shadow-[0_1px_2px_0_rgb(0_0_0_/_0.05)] {className}"
+	style="border-color: color-mix(in oklch, var(--border) 30%, transparent); background: var(--code-background); max-height: {maxHeight}; {maxWidth
+		? `max-width: ${maxWidth};`
+		: ''}"
 >
-	<!-- Needs to be formatted as single line for proper rendering -->
+	<!-- Single line: hljs injection depends on a contiguous source string. -->
 	<pre class="m-0"><code class="hljs text-sm leading-relaxed">{@html highlightedHtml}</code></pre>
 </div>
 
 <style>
+	.code-preview-wrapper {
+		overscroll-behavior: contain;
+	}
+
 	.code-preview-wrapper pre {
 		background: transparent;
+		padding: 0;
 	}
 
 	.code-preview-wrapper code {
 		background: transparent;
+		display: block;
+		padding: 0.5rem;
+	}
+
+	:global(.dark) .code-preview-wrapper {
+		border-color: color-mix(in oklch, var(--border) 20%, transparent);
 	}
 </style>

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

@@ -68,7 +68,6 @@ export { default as SyntaxHighlightedCode } from './SyntaxHighlightedCode.svelte
  * ```svelte
  * <CollapsibleContentBlock
  *   bind:open
- *   icon={BrainIcon}
  *   title="Thinking..."
  *   isStreaming
  * >
@@ -78,6 +77,22 @@ export { default as SyntaxHighlightedCode } from './SyntaxHighlightedCode.svelte
  */
 export { default as CollapsibleContentBlock } from './CollapsibleContentBlock.svelte';
 
+/**
+ * **CollapsibleTerminalBlock** - Expandable content card with a terminal-style frame
+ *
+ * Same shape as CollapsibleContentBlock, but with a `code-background`
+ * fill, subtle border, and tightened padding suited for shell command
+ * output and similar dense / monospace content.
+ *
+ * @example
+ * ```svelte
+ * <CollapsibleTerminalBlock bind:open title="Run command">
+ *   <pre>{output}</pre>
+ * </CollapsibleTerminalBlock>
+ * ```
+ */
+export { default as CollapsibleTerminalBlock } from './CollapsibleTerminalBlock.svelte';
+
 /**
  * **MermaidPreview** - Interactive Mermaid diagram viewer
  *

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { FolderOpen, Plus, Loader2, Braces } from '@lucide/svelte';
 	import { toast } from 'svelte-sonner';
 	import * as Dialog from '$lib/components/ui/dialog';
@@ -289,7 +290,7 @@
 				{#if selectedTemplate && !templatePreviewContent}
 					<div class="flex h-full flex-col">
 						<div class="mb-3 flex items-center gap-2">
-							<Braces class="h-4 w-4 text-muted-foreground" />
+							<Braces class="{ICON_CLASS_DEFAULT} text-muted-foreground" />
 
 							<span class="text-sm font-medium">
 								{selectedTemplate.title || selectedTemplate.name}
@@ -371,9 +372,9 @@
 			{#if hasTemplateResult}
 				<Button onclick={handleAttachTemplateResource} disabled={isAttaching}>
 					{#if isAttaching}
-						<Loader2 class="mr-2 h-4 w-4 animate-spin" />
+						<Loader2 class="mr-2 {ICON_CLASS_DEFAULT} animate-spin" />
 					{:else}
-						<Plus class="mr-2 h-4 w-4" />
+						<Plus class="mr-2 {ICON_CLASS_DEFAULT}" />
 					{/if}
 
 					Attach Resource
@@ -381,9 +382,9 @@
 			{:else}
 				<Button onclick={handleAttach} disabled={selectedResources.size === 0 || isAttaching}>
 					{#if isAttaching}
-						<Loader2 class="mr-2 h-4 w-4 animate-spin" />
+						<Loader2 class="mr-2 {ICON_CLASS_DEFAULT} animate-spin" />
 					{:else}
-						<Plus class="mr-2 h-4 w-4" />
+						<Plus class="mr-2 {ICON_CLASS_DEFAULT}" />
 					{/if}
 
 					Attach {selectedResources.size > 0 ? `(${selectedResources.size})` : 'Resource'}

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

@@ -1,11 +1,20 @@
 <script lang="ts">
 	import { Button } from '$lib/components/ui/button';
 	import * as Dialog from '$lib/components/ui/dialog';
-	import { McpServerForm } from '$lib/components/app/mcp';
+	import { McpServerCardCompact, McpServerForm } from '$lib/components/app/mcp';
 	import { mcpStore } from '$lib/stores/mcp.svelte';
 	import { conversationsStore } from '$lib/stores/conversations.svelte';
-	import { parseHeadersToArray, uuid } from '$lib/utils';
-	import { MCP_SERVER_ID_PREFIX } from '$lib/constants';
+	import { parseHeadersToArray, uuid, canonicalizeServerUrl } from '$lib/utils';
+	import {
+		BEARER_PREFIX,
+		BOOL_FALSE_STRING,
+		BOOL_TRUE_STRING,
+		DISMISSED_RECOMMENDED_MCP_SERVERS_LOCALSTORAGE_KEY,
+		MCP_SERVER_ID_PREFIX,
+		RECOMMENDED_MCP_SERVERS,
+		REDACTED_HEADERS
+	} from '$lib/constants';
+	import { browser } from '$app/environment';
 
 	interface Props {
 		open: boolean;
@@ -17,6 +26,39 @@
 	let newServerUrl = $state('');
 	let newServerHeaders = $state('');
 	let newServerUseProxy = $state(false);
+
+	let newServerWantsAuthorization = $state(false);
+
+	let selectedRecommendationId = $derived.by(() => {
+		const url = newServerUrl.trim();
+		if (!url) return null;
+		const targetCanonical = canonicalizeServerUrl(url);
+		return (
+			RECOMMENDED_MCP_SERVERS.find((rec) => canonicalizeServerUrl(rec.url) === targetCanonical)
+				?.id ?? null
+		);
+	});
+	let selectedRecommendation = $derived(
+		selectedRecommendationId
+			? (RECOMMENDED_MCP_SERVERS.find((rec) => rec.id === selectedRecommendationId) ?? null)
+			: null
+	);
+	let authRequired = $derived(selectedRecommendation?.needsAuthorization ?? false);
+
+	let bearerTokenFilled = $derived.by(() => {
+		const pairs = parseHeadersToArray(newServerHeaders);
+		const bearerPrefix = BEARER_PREFIX.toLowerCase();
+		const bearer = pairs.find(
+			(p) =>
+				REDACTED_HEADERS.has(p.key.trim().toLowerCase()) &&
+				p.value.trim().toLowerCase().startsWith(bearerPrefix)
+		);
+
+		if (!bearer) return false;
+
+		return bearer.value.trim().slice(bearerPrefix.length).trim().length > 0;
+	});
+
 	let newServerUrlError = $derived.by(() => {
 		if (!newServerUrl.trim()) return 'URL is required';
 		try {
@@ -30,13 +72,83 @@
 	let newServerHeaderPairsValid = $derived(
 		parseHeadersToArray(newServerHeaders).every((p) => p.key.trim() && p.value.trim())
 	);
-	let canSave = $derived(!newServerUrlError && newServerHeaderPairsValid);
+	let canSave = $derived(
+		!newServerUrlError && newServerHeaderPairsValid && (!authRequired || bearerTokenFilled)
+	);
+
+	// Backward-compatible read: older versions stored a JSON array of dismissed ids.
+	function readRecommendationsDismissed(): boolean {
+		if (!browser) return false;
+		const raw = localStorage.getItem(DISMISSED_RECOMMENDED_MCP_SERVERS_LOCALSTORAGE_KEY);
+
+		if (!raw) return false;
+
+		if (raw === BOOL_TRUE_STRING) return true;
+		if (raw === BOOL_FALSE_STRING) return false;
+
+		try {
+			const parsed = JSON.parse(raw);
+			return Array.isArray(parsed) && parsed.length > 0;
+		} catch {
+			return false;
+		}
+	}
+
+	function writeRecommendationsDismissed(dismissed: boolean) {
+		recommendationsDismissed = dismissed;
+
+		if (browser) {
+			localStorage.setItem(
+				DISMISSED_RECOMMENDED_MCP_SERVERS_LOCALSTORAGE_KEY,
+				dismissed ? BOOL_TRUE_STRING : BOOL_FALSE_STRING
+			);
+		}
+	}
+
+	let recommendationsDismissed = $state<boolean>(readRecommendationsDismissed());
+
+	// Read-only once a recommendation is picked: switch is disabled, so we keep
+	// the Authorization field in sync with the requirement.
+	$effect(() => {
+		if (authRequired) {
+			newServerWantsAuthorization = true;
+		}
+	});
+
+	let hasSelection = $derived(selectedRecommendationId !== null);
+
+	let unconfiguredRecommendations = $derived.by(() => {
+		const configuredCanonicals = new Set(
+			mcpStore.getServers().map((s) => canonicalizeServerUrl(s.url))
+		);
+
+		return RECOMMENDED_MCP_SERVERS.filter(
+			(rec) => !configuredCanonicals.has(canonicalizeServerUrl(rec.url))
+		);
+	});
+
+	let recommendationsToShow = $derived(recommendationsDismissed ? [] : unconfiguredRecommendations);
+
+	function handleRecommendationClick(recommendedId: string) {
+		const recommendation = RECOMMENDED_MCP_SERVERS.find((rec) => rec.id === recommendedId);
+
+		if (!recommendation) return;
+
+		newServerUrl = recommendation.url;
+		newServerHeaders = '';
+		newServerWantsAuthorization = recommendation.needsAuthorization ?? false;
+	}
+
+	function handleDismissAll() {
+		writeRecommendationsDismissed(true);
+	}
 
 	function handleOpenChange(value: boolean) {
 		if (!value) {
 			newServerUrl = '';
 			newServerHeaders = '';
 			newServerUseProxy = false;
+			newServerWantsAuthorization = false;
 		}
 		open = value;
 		onOpenChange?.(value);
@@ -67,11 +179,33 @@
 </script>
 
 <Dialog.Root {open} onOpenChange={handleOpenChange}>
-	<Dialog.Content class="sm:max-w-md">
+	<Dialog.Content class="sm:max-w-2xl">
 		<Dialog.Header>
-			<Dialog.Title>Add New Server</Dialog.Title>
+			<Dialog.Title class="select-none">Add New MCP Server</Dialog.Title>
 		</Dialog.Header>
 
+		{#if recommendationsToShow.length > 0}
+			<div class="space-y-3 pt-2">
+				<div class="flex items-center justify-between gap-3">
+					<h3 class="text-sm font-medium">Recommended Servers</h3>
+					<Button class="text-muted-foreground" variant="ghost" size="sm" onclick={handleDismissAll}
+						>Dismiss</Button
+					>
+				</div>
+
+				<div class="grid grid-cols-1 gap-3 sm:grid-cols-2">
+					{#each recommendationsToShow as recommendation (recommendation.id)}
+						<McpServerCardCompact
+							server={recommendation}
+							onClick={() => handleRecommendationClick(recommendation.id)}
+							selected={selectedRecommendationId === recommendation.id}
+							dimmed={hasSelection && selectedRecommendationId !== recommendation.id}
+						/>
+					{/each}
+				</div>
+			</div>
+		{/if}
+
 		<form onsubmit={handleSubmit} class="contents">
 			<div class="space-y-4 py-4">
 				<McpServerForm
@@ -83,6 +217,8 @@
 					onUseProxyChange={(v) => (newServerUseProxy = v)}
 					urlError={newServerUrl ? newServerUrlError : null}
 					id="new-server"
+					bind:wantsAuthorization={newServerWantsAuthorization}
+					required={authRequired}
 				/>
 			</div>
 

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import * as AlertDialog from '$lib/components/ui/alert-dialog';
 	import { AlertTriangle, ArrowRight } from '@lucide/svelte';
 	import { goto } from '$app/navigation';
@@ -60,7 +61,7 @@
 							>
 								<span class="min-w-0 truncate font-mono text-xs">{model}</span>
 								<ArrowRight
-									class="h-4 w-4 shrink-0 text-muted-foreground opacity-0 transition-opacity group-hover:opacity-100"
+									class="{ICON_CLASS_DEFAULT} shrink-0 text-muted-foreground opacity-0 transition-opacity group-hover:opacity-100"
 								/>
 							</button>
 						{/each}

tools/ui/src/lib/components/app/forms/KeyValuePairs.svelte

@@ -101,7 +101,7 @@
 <div class={className}>
 	<div class="mb-2 flex items-center justify-between">
 		{#if sectionLabel}
-			<span class="text-xs font-medium">
+			<span class="text-xs font-medium select-none">
 				{sectionLabel}
 				{#if sectionLabelOptional}
 					<span class="text-muted-foreground">(optional)</span>
@@ -118,6 +118,7 @@
 			{addButtonLabel}
 		</button>
 	</div>
+
 	{#if pairs.length > 0}
 		<div class="space-y-3">
 			{#each pairs as pair, index (index)}
@@ -159,6 +160,6 @@
 			{/each}
 		</div>
 	{:else}
-		<p class="text-xs text-muted-foreground">{emptyMessage}</p>
+		<p class="select-none text-xs text-muted-foreground">{emptyMessage}</p>
 	{/if}
 </div>

tools/ui/src/lib/components/app/forms/SearchInput.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { Input } from '$lib/components/ui/input';
 	import { Search, X } from '@lucide/svelte';
 
@@ -50,7 +51,7 @@
 
 <div class="relative {className}">
 	<Search
-		class="absolute top-1/2 left-3 z-10 h-4 w-4 -translate-y-1/2 transform text-muted-foreground"
+		class="absolute top-1/2 left-3 z-10 {ICON_CLASS_DEFAULT} -translate-y-1/2 transform text-muted-foreground"
 	/>
 
 	<Input
@@ -72,7 +73,7 @@
 			onclick={handleClear}
 			aria-label={value ? 'Clear search' : 'Close'}
 		>
-			<X class="h-4 w-4" />
+			<X class={ICON_CLASS_DEFAULT} />
 		</button>
 	{/if}
 </div>

tools/ui/src/lib/components/app/mcp/McpActiveServersAvatars.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import * as Tooltip from '$lib/components/ui/tooltip';
 	import { conversationsStore } from '$lib/stores/conversations.svelte';
 	import { mcpStore } from '$lib/stores/mcp.svelte';
@@ -50,7 +51,7 @@
 	>
 		<Tooltip.Root>
 			<Tooltip.Trigger>
-				<McpLogo class="h-4 w-4" />
+				<McpLogo class={ICON_CLASS_DEFAULT} />
 			</Tooltip.Trigger>
 
 			<Tooltip.Content>
@@ -68,7 +69,7 @@
 							<img
 								src={favicon.url}
 								alt=""
-								class="h-4 w-4"
+								class={ICON_CLASS_DEFAULT}
 								onerror={(e) => {
 									(e.currentTarget as HTMLImageElement).style.display = 'none';
 								}}

tools/ui/src/lib/components/app/mcp/McpResourcePreview.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { FileText, Loader2, AlertCircle, Download } from '@lucide/svelte';
 	import { Button } from '$lib/components/ui/button';
 	import { mcpStore } from '$lib/stores/mcp.svelte';
@@ -140,7 +141,7 @@
 						/>
 					{:else}
 						<div class="flex items-center gap-2 rounded bg-muted p-2 text-sm text-muted-foreground">
-							<FileText class="h-4 w-4" />
+							<FileText class={ICON_CLASS_DEFAULT} />
 
 							<span>Binary content ({blob.mimeType || 'unknown type'})</span>
 						</div>

tools/ui/src/lib/components/app/mcp/McpResourcesBrowser/McpResourcesBrowserHeader.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { RefreshCw, Loader2 } from '@lucide/svelte';
 	import { Button } from '$lib/components/ui/button';
 	import { SearchInput } from '$lib/components/app/forms';
@@ -30,9 +31,9 @@
 			title="Refresh resources"
 		>
 			{#if isLoading}
-				<Loader2 class="h-4 w-4 animate-spin" />
+				<Loader2 class="{ICON_CLASS_DEFAULT} animate-spin" />
 			{:else}
-				<RefreshCw class="h-4 w-4" />
+				<RefreshCw class={ICON_CLASS_DEFAULT} />
 			{/if}
 		</Button>
 	</div>

tools/ui/src/lib/components/app/mcp/McpResourcesBrowser/McpResourcesBrowserServerItem.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { FolderOpen, ChevronDown, ChevronRight, Loader2, Braces } from '@lucide/svelte';
 	import { Checkbox } from '$lib/components/ui/checkbox';
 	import * as Collapsible from '$lib/components/ui/collapsible';
@@ -122,7 +123,7 @@
 					checked={isSelected}
 					onCheckedChange={(checked: boolean | 'indeterminate') =>
 						handleCheckboxChange(resource, checked === true)}
-					class="h-4 w-4"
+					class={ICON_CLASS_DEFAULT}
 				/>
 			{/if}
 
@@ -160,7 +161,7 @@
 				<McpServerIdentity
 					displayName={serverDisplayName}
 					faviconUrl={serverFaviconUrl}
-					iconClass="h-4 w-4"
+					iconClass={ICON_CLASS_DEFAULT}
 					showVersion={false}
 				/>
 			</div>

tools/ui/src/lib/components/app/mcp/McpServerCard/McpServerCard.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { tick } from 'svelte';
 	import * as Card from '$lib/components/ui/card';
 	import { Skeleton } from '$lib/components/ui/skeleton';
@@ -134,7 +135,7 @@
 			{#if showSkeleton}
 				<div class="space-y-2">
 					<div class="flex items-center gap-2">
-						<Skeleton class="h-4 w-4 rounded" />
+						<Skeleton class="{ICON_CLASS_DEFAULT} rounded" />
 						<Skeleton class="h-3 w-24" />
 					</div>
 					<div class="flex flex-wrap gap-1.5">
@@ -146,7 +147,7 @@
 
 				<div class="space-y-1.5">
 					<div class="flex items-center gap-2">
-						<Skeleton class="h-4 w-4 rounded" />
+						<Skeleton class="{ICON_CLASS_DEFAULT} rounded" />
 						<Skeleton class="h-3 w-32" />
 					</div>
 				</div>

tools/ui/src/lib/components/app/mcp/McpServerCard/McpServerCardCompact.svelte

@@ -0,0 +1,44 @@
+<script lang="ts">
+	import * as Card from '$lib/components/ui/card';
+	import { mode } from 'mode-watcher';
+	import type { RecommendedMCPServer } from '$lib/types';
+
+	interface Props {
+		server: RecommendedMCPServer;
+		onClick?: () => void;
+		selected?: boolean;
+		dimmed?: boolean;
+	}
+
+	let { server, onClick, selected = false, dimmed = false }: Props = $props();
+
+	let activeIconUrl = $derived.by(() => {
+		const isDark = mode.current === 'dark';
+
+		if (isDark && server.iconUrlDark) return server.iconUrlDark;
+		if (!isDark && server.iconUrlLight) return server.iconUrlLight;
+
+		return server.iconUrl;
+	});
+</script>
+
+<Card.Root
+	class={`relative gap-3! select-none bg-muted/30 p-4 transition-all ${onClick ? 'cursor-pointer hover:bg-muted/50 hover:opacity-100' : ''} ${selected ? 'bg-muted/30 ring-1 ring-primary/40' : ''} ${dimmed ? 'opacity-50' : ''}`}
+	onclick={onClick}
+>
+	<div class="flex min-w-0 items-center gap-2">
+		{#if activeIconUrl}
+			<img
+				src={activeIconUrl}
+				alt=""
+				class="h-5 w-5 shrink-0 rounded"
+				loading="lazy"
+				decoding="async"
+			/>
+		{/if}
+
+		<h4 class="min-w-0 flex-1 truncate font-medium">{server.name}</h4>
+	</div>
+
+	<p class="text-xs text-muted-foreground">{server.description}</p>
+</Card.Root>

tools/ui/src/lib/components/app/mcp/McpServerForm.svelte

@@ -5,9 +5,14 @@
 	import type { KeyValuePair } from '$lib/types';
 	import { parseHeadersToArray, serializeHeaders } from '$lib/utils';
 	import { UrlProtocol } from '$lib/enums';
-	import { MCP_SERVER_URL_PLACEHOLDER } from '$lib/constants';
+	import {
+		AUTHORIZATION_HEADER,
+		BEARER_PREFIX,
+		CLI_FLAGS,
+		MCP_SERVER_URL_PLACEHOLDER,
+		REDACTED_HEADERS
+	} from '$lib/constants';
 	import { mcpStore } from '$lib/stores/mcp.svelte';
-	import { CLI_FLAGS } from '$lib/constants';
 
 	interface Props {
 		url: string;
@@ -18,6 +23,22 @@
 		onUseProxyChange?: (useProxy: boolean) => void;
 		urlError?: string | null;
 		id?: string;
+		/**
+		 * "Wants Authorization" is the user's *intent* to add a Bearer token
+		 * (separate from `hasAuthorization` which reflects what's already in
+		 * the headers). Bindable so a parent - e.g. the recommendation cards
+		 * on the "Add New Server" dialog - can flip the switch on when the
+		 * picked server ships a `needsAuthorization: true` flag.
+		 */
+		wantsAuthorization?: boolean;
+		/**
+		 * Marks the "Authorization" field as required. Locks the toggle so the
+		 * user can't dismiss it, and visually marks the field with a red
+		 * asterisk. The parent is expected to gate its submit affordance on
+		 * the bearer token actually being filled. Used by the "Add New Server"
+		 * dialog for recommendations whose `needsAuthorization` flag is true.
+		 */
+		required?: boolean;
 	}
 
 	let {
@@ -28,7 +49,9 @@
 		onHeadersChange,
 		onUseProxyChange,
 		urlError = null,
-		id = 'server'
+		id = 'server',
+		wantsAuthorization = $bindable(false),
+		required = false
 	}: Props = $props();
 
 	let isWebSocket = $derived(
@@ -38,14 +61,11 @@
 
 	let headerPairs = $derived<KeyValuePair[]>(parseHeadersToArray(headers));
 
-	const AUTHORIZATION_HEADER = 'Authorization';
-	const BEARER_PREFIX = 'Bearer ';
-
 	// Heuristic: this dedicated UI only owns Authorization headers that already
 	// carry a Bearer scheme. Anything else (e.g. Basic, raw tokens) stays in the
 	// KV section so the user can still edit those values verbatim.
 	const matchesAuthorizationKey = (key: string): boolean =>
-		key.trim().toLowerCase() === AUTHORIZATION_HEADER.toLowerCase();
+		REDACTED_HEADERS.has(key.trim().toLowerCase());
 
 	const isBearerScheme = (value: string): boolean =>
 		value.trim().toLowerCase().startsWith(BEARER_PREFIX.toLowerCase());
@@ -55,8 +75,6 @@
 
 	let hasAuthorization = $derived(headerPairs.some(ownedByBearerUi));
 
-	let wantsAuthorization = $state(false);
-
 	let showAuthorization = $derived(hasAuthorization || wantsAuthorization);
 
 	let urlInput: HTMLInputElement | null = $state(null);
@@ -119,7 +137,7 @@
 
 <div class="grid gap-2">
 	<div class="mb-4">
-		<label for="server-url-{id}" class="mb-2 block text-xs font-medium">
+		<label for="server-url-{id}" class="mb-2 block text-xs font-medium select-none">
 			Server URL <span class="text-destructive">*</span>
 		</label>
 
@@ -138,14 +156,18 @@
 		{/if}
 	</div>
 
-	<label class="flex items-center gap-2 cursor-pointer">
+	<label class="flex items-center gap-2 cursor-pointer select-none">
 		<Switch
 			id="use-authorization-{id}"
 			checked={showAuthorization}
 			onCheckedChange={setUseAuthorization}
+			disabled={required}
 		/>
 
-		<span class="text-xs text-muted-foreground">Authorization</span>
+		<span class="text-xs text-muted-foreground">
+			Authorization{#if required}
+				<span class="text-destructive">*</span>{/if}
+		</span>
 	</label>
 
 	{#if showAuthorization}

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

@@ -180,6 +180,17 @@ export { default as McpServerCardDeleteDialog } from './McpServerCard/McpServerC
 /** Skeleton loading state for server card during health checks. */
 export { default as McpServerCardSkeleton } from './McpServerCardSkeleton.svelte';
 
+/**
+ * **McpServerCardCompact** - Condensed MCP server card
+ *
+ * Static card for picker-style UIs (e.g. recommended MCP servers in the
+ * Add New Server dialog). Shows an optional favicon, the server name, and
+ * a short description. Performs no network requests - safe to render
+ * without contacting any upstream server until the user explicitly adds
+ * the server.
+ */
+export { default as McpServerCardCompact } from './McpServerCard/McpServerCardCompact.svelte';
+
 /**
  * **McpServerIdentity** - Server identity display (icon, name, version)
  *

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { ChevronLeft, ChevronRight } from '@lucide/svelte';
 	import type { Snippet } from 'svelte';
 
@@ -71,7 +72,7 @@
 		disabled={!canScrollLeft}
 		aria-label="Scroll left"
 	>
-		<ChevronLeft class="h-4 w-4" />
+		<ChevronLeft class={ICON_CLASS_DEFAULT} />
 	</button>
 
 	<div
@@ -88,6 +89,6 @@
 		disabled={!canScrollRight}
 		aria-label="Scroll right"
 	>
-		<ChevronRight class="h-4 w-4" />
+		<ChevronRight class={ICON_CLASS_DEFAULT} />
 	</button>
 </div>

tools/ui/src/lib/components/app/models/ModelsSelectorOption.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import {
 		CircleAlert,
 		Heart,
@@ -121,7 +122,7 @@
 
 		{#if isLoading}
 			<div class="flex w-4 items-center justify-center [@media(pointer:coarse)]:w-5">
-				<Loader2 class="h-4 w-4 animate-spin text-muted-foreground" />
+				<Loader2 class="{ICON_CLASS_DEFAULT} animate-spin text-muted-foreground" />
 			</div>
 		{:else if isFailed}
 			<div class="flex w-4 items-center justify-center [@media(pointer:coarse)]:w-auto">

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { goto } from '$app/navigation';
 	import { page } from '$app/state';
 	import { Search } from '@lucide/svelte';
@@ -85,7 +86,7 @@
 </script>
 
 {#snippet itemIcon(IconComponent: Component)}
-	<IconComponent class="h-4 w-4" />
+	<IconComponent class={ICON_CLASS_DEFAULT} />
 {/snippet}
 
 {#if isSearchModeActive}
@@ -183,7 +184,7 @@
 						tooltip={item.tooltip}
 						tooltipSide={TooltipSide.RIGHT}
 						size="lg"
-						iconSize="h-4 w-4"
+						iconSize={ICON_CLASS_DEFAULT}
 						class="h-9 w-9 rounded-full hover:bg-accent! {isActive
 							? 'bg-accent text-accent-foreground'
 							: ''}"

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

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import {
 		Trash2,
 		Pencil,
@@ -147,7 +148,7 @@
 			<Tooltip.Root>
 				<Tooltip.Trigger>
 					<div
-						class="stop-button flex h-4 w-4 shrink-0 cursor-pointer items-center justify-center rounded text-muted-foreground transition-colors hover:text-foreground"
+						class="stop-button flex {ICON_CLASS_DEFAULT} shrink-0 cursor-pointer items-center justify-center rounded text-muted-foreground transition-colors hover:text-foreground"
 						onclick={handleStop}
 						onkeydown={(e) => e.key === 'Enter' && handleStop(e)}
 						role="button"

tools/ui/src/lib/components/app/server/ServerErrorSplash.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { base } from '$app/paths';
 	import { AlertTriangle, RefreshCw, Key, CheckCircle, XCircle } from '@lucide/svelte';
 	import { goto } from '$app/navigation';
@@ -7,7 +8,7 @@
 	import Label from '$lib/components/ui/label/label.svelte';
 	import { serverStore, serverLoading } from '$lib/stores/server.svelte';
 	import { config, settingsStore } from '$lib/stores/settings.svelte';
-	import { SETTINGS_KEYS } from '$lib/constants';
+	import { AUTHORIZATION_HEADER, BEARER_PREFIX, SETTINGS_KEYS } from '$lib/constants';
 	import { ROUTES } from '$lib/constants/routes';
 	import { fade, fly, scale } from 'svelte/transition';
 	import { KeyboardKey } from '$lib/enums';
@@ -71,7 +72,7 @@
 			const response = await fetch(`${base}/props`, {
 				headers: {
 					'Content-Type': 'application/json',
-					Authorization: `Bearer ${apiKeyInput.trim()}`
+					[AUTHORIZATION_HEADER]: `${BEARER_PREFIX}${apiKeyInput.trim()}`
 				}
 			});
 
@@ -145,7 +146,7 @@
 		{#if isAccessDeniedError && !showApiKeyInput}
 			<div in:fly={{ y: 10, duration: 300, delay: 200 }} class="mb-4">
 				<Button onclick={handleShowApiKeyInput} variant="outline" class="w-full">
-					<Key class="h-4 w-4" />
+					<Key class={ICON_CLASS_DEFAULT} />
 					Enter API Key
 				</Button>
 			</div>
@@ -171,21 +172,21 @@
 						/>
 						{#if apiKeyState === 'validating'}
 							<div class="absolute top-1/2 right-3 -translate-y-1/2">
-								<RefreshCw class="h-4 w-4 animate-spin text-muted-foreground" />
+								<RefreshCw class="{ICON_CLASS_DEFAULT} animate-spin text-muted-foreground" />
 							</div>
 						{:else if apiKeyState === 'success'}
 							<div
 								class="absolute top-1/2 right-3 -translate-y-1/2"
 								in:scale={{ duration: 200, start: 0.8 }}
 							>
-								<CheckCircle class="h-4 w-4 text-green-500" />
+								<CheckCircle class="{ICON_CLASS_DEFAULT} text-green-500" />
 							</div>
 						{:else if apiKeyState === 'error'}
 							<div
 								class="absolute top-1/2 right-3 -translate-y-1/2"
 								in:scale={{ duration: 200, start: 0.8 }}
 							>
-								<XCircle class="h-4 w-4 text-destructive" />
+								<XCircle class="{ICON_CLASS_DEFAULT} text-destructive" />
 							</div>
 						{/if}
 					</div>
@@ -209,7 +210,7 @@
 						class="flex-1"
 					>
 						{#if apiKeyState === 'validating'}
-							<RefreshCw class="h-4 w-4 animate-spin" />
+							<RefreshCw class="{ICON_CLASS_DEFAULT} animate-spin" />
 							Validating...
 						{:else if apiKeyState === 'success'}
 							Success!
@@ -237,11 +238,11 @@
 			<div in:fly={{ y: 10, duration: 300, delay: 200 }}>
 				<Button onclick={handleRetryConnection} disabled={isServerLoading} class="w-full">
 					{#if isServerLoading}
-						<RefreshCw class="h-4 w-4 animate-spin" />
+						<RefreshCw class="{ICON_CLASS_DEFAULT} animate-spin" />
 
 						Connecting...
 					{:else}
-						<RefreshCw class="h-4 w-4" />
+						<RefreshCw class={ICON_CLASS_DEFAULT} />
 
 						Retry Connection
 					{/if}

tools/ui/src/lib/components/app/server/ServerStatus.svelte

@@ -1,4 +1,5 @@
 <script lang="ts">
+	import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes';
 	import { AlertTriangle, Server } from '@lucide/svelte';
 	import { Badge } from '$lib/components/ui/badge';
 	import { Button } from '$lib/components/ui/button';
@@ -57,7 +58,7 @@
 
 	{#if showActions && error}
 		<Button variant="outline" size="sm" class="text-destructive">
-			<AlertTriangle class="h-4 w-4" />
+			<AlertTriangle class={ICON_CLASS_DEFAULT} />
 
 			{error}
 		</Button>