From de7fa0a3c6a2e1b4cd9f22eb8d6bf5b12dbdb63b Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Thu, 8 Oct 2026 12:56:54 -0700 Subject: [PATCH 001/400] Musa FWHT fix (#30167) --- ggml/src/ggml-cuda/fwht.cu | 3 +++ 1 file changed, 3 insertions(+) diff --git a/ggml/src/ggml-cuda/fwht.cu b/ggml/src/ggml-cuda/fwht.cu index d69e2ed81c7b..78768eace3d0 100644 --- a/ggml/src/ggml-cuda/fwht.cu +++ b/ggml/src/ggml-cuda/fwht.cu @@ -200,9 +200,12 @@ static bool ggml_cuda_op_fwht_impl(ggml_backend_cuda_context & ctx, const ggml_t case 4096: ggml_cuda_kernel_launch(fwht_cuda_block<4096, nt, T>, launch_params_w, src_d, dst_d, rows, scale); return true; +#if !defined(GGML_USE_MUSA) + // 32 KB of shared memory, above the MUSA limit; falls back there case 8192: ggml_cuda_kernel_launch(fwht_cuda_block<8192, nt, T>, launch_params_w, src_d, dst_d, rows, scale); return true; +#endif // !defined(GGML_USE_MUSA) default: return false; } From 3d65c90d04d337e88f2b1f7f0061f40a5324e662 Mon Sep 17 00:00:00 2001 From: Nicolas Mowen Date: Thu, 8 Oct 2026 20:19:32 -0600 Subject: [PATCH 002/400] sycl : Q5_K reorder-layout MMVQ and fused GLU (#29375) --- ggml/src/ggml-sycl/ggml-sycl.cpp | 13 +++- ggml/src/ggml-sycl/mmvq.cpp | 108 +++++++++++++++++-------------- ggml/src/ggml-sycl/mmvq.hpp | 6 ++ ggml/src/ggml-sycl/vecdotq.hpp | 90 ++++++++++++++++---------- 4 files changed, 133 insertions(+), 84 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 2a5184e01283..d2c1a1a44839 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5102,12 +5102,19 @@ static bool ggml_sycl_mul_mat_glu_mmvq_fused(ggml_backend_sycl_context & ctx, gg return false; } - // quant pairs the reorder kernel cannot serve (mixed gate/up types) take the - // standard-layout fused path instead; q4_K keeps the reorder path below - if (wg->type != GGML_TYPE_Q4_K || wu->type != GGML_TYPE_Q4_K) { + // quant pairs the reorder kernel does not serve (mixed gate/up types, q5_K off BMG) take the + // standard-layout fused path instead; same-type q4_K / q5_K keep the reorder path below + const bool reorder_pair = wg->type == wu->type && + (wu->type == GGML_TYPE_Q4_K || (wu->type == GGML_TYPE_Q5_K && ggml_sycl_q5_k_mmvq_reuse(ctx.device))); + if (!reorder_pair) { return ggml_sycl_mul_mat_glu_mmvq_plain(ctx, glu, gate, up, wu, wg, act); } + // past 5 columns the two unfused q5_K GEMVs are faster than the fused kernel + if (wu->type == GGML_TYPE_Q5_K && act->ne[1] > 5) { + return false; + } + // install the reorder (SoA) layout the fused kernel needs, as the unfused mmvq path would; // a no-op once done. after the bail checks so a declined op does not pay for it. opt_for_reorder(&ctx, wu, act, up, mul_mat_algo::MMVQ); diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index ceaacb59ece5..e8b293a97d75 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -110,7 +110,8 @@ static void mul_mat_vec_q_reorder(const void * __restrict__ vx, const void * __r // With has_fusion, `vgate` is a second weight matrix sharing vx's shape, stride and reorder // layout: one pass computes both row dot products and the epilogue writes glu(gate, up). -template +template ::value> static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void * __restrict__ vgate, const void * __restrict__ vy, float * __restrict__ dst, const int ncols, const int nrows, const int stride_col_y_bytes, const int stride_col_dst, @@ -181,7 +182,7 @@ static void mul_mat_vec_q_reorder_ncols(const void * __restrict__ vx, const void } } } - } else if constexpr (reorder_vec_dot_shared_weights::value) { + } else if constexpr (shared_weights) { const int ibx = row0 * blocks_per_row + i; const auto bx_offset = block_type::get_block_offset(ibx, nblocks); const auto d_offset = block_type::get_d_offset(nrows, ncols, ibx); @@ -1945,8 +1946,8 @@ static void reorder_mul_mat_vec_q5_k_q8_1_sycl(const void * vx, const void * vy, }); } -template -static void reorder_mul_mat_vec_q5_k_q8_1_sycl_ncols( +template +static void reorder_mul_mat_vec_q5_k_q8_1_sycl_ncols_impl( const void * vx, const void * vy, float * dst, const int ncols, const int nrows, const int stride_col_y_bytes, const int stride_col_dst, @@ -1954,20 +1955,35 @@ static void reorder_mul_mat_vec_q5_k_q8_1_sycl_ncols( GGML_ASSERT(ncols % QK_K == 0); constexpr size_t num_subgroups = WARP_SIZE; - const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups); + const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups * rows_per_sg); const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); stream->submit([&](sycl::handler & cgh) { cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { - mul_mat_vec_q_reorder_ncols, ncols_dst>( + mul_mat_vec_q_reorder_ncols, ncols_dst, + /*has_fusion=*/ false, rows_per_sg, shared_weights>( vx, /*vgate=*/ nullptr, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, /*glu_op=*/ GGML_GLU_OP_SWIGLU, nd_item); }); }); } +template +static void reorder_mul_mat_vec_q5_k_q8_1_sycl_ncols( + const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, + const int stride_col_y_bytes, const int stride_col_dst, + dpct::queue_ptr stream) { + if (ggml_sycl_q5_k_mmvq_reuse(ggml_sycl_get_device())) { + constexpr int rows_per_sg = ncols_dst >= 3 ? 2 : 1; + reorder_mul_mat_vec_q5_k_q8_1_sycl_ncols_impl(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); + } else { + reorder_mul_mat_vec_q5_k_q8_1_sycl_ncols_impl(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); + } +} + static void reorder_mul_mat_vec_q5_k_q8_1_sycl_switch_ncols( const void * vx, const void * vy, float * dst, const int ncols, const int nrows, const int ncols_dst, @@ -3129,8 +3145,11 @@ static void launch_mul_mat_vec_q_reorder_glu(const void * vx, const void * vgate const int ncols, const int nrows, const int stride_col_y_bytes, const int stride_col_dst, const ggml_glu_op glu_op, dpct::queue_ptr stream) { + // q4_K pairs rows for 3..4 columns, q5_K for 3..5 + constexpr int row_pair_max = reorder_vec_dot_q_sycl::gtype == GGML_TYPE_Q5_K ? 5 : 4; constexpr int rows_per_sg = - reorder_vec_dot_shared_activations::value && ncols_dst >= 3 && ncols_dst <= 4 + reorder_vec_dot_shared_activations::value && ncols_dst >= 3 && + ncols_dst <= row_pair_max ? 2 : 1; launch_mul_mat_vec_q_reorder_glu_impl(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream); @@ -3321,55 +3340,48 @@ bool ggml_sycl_mul_mat_vec_q_glu_plain(enum ggml_type gate_type, enum ggml_type return false; } +template +static bool mul_mat_vec_q_glu_reorder_ncols(enum ggml_glu_op glu_op, const void * vx, const void * vgate, + const void * vy, float * dst, int ncols, int nrows, int ncols_dst, + int stride_col_y_bytes, int stride_col_dst, dpct::queue_ptr stream) { + using vec_dot = reorder_vec_dot_q_sycl; + + auto launch = [&](auto I) -> bool { + constexpr int n = decltype(I)::value; + if (ncols_dst != n) { + return false; + } + if constexpr (type == GGML_TYPE_Q4_K && n == 2) { + if (nrows >= Q4_K_MMVQ_ROW_PAIR_MIN_NROWS) { + launch_mul_mat_vec_q_reorder_glu_impl(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream); + return true; + } + } + launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, + stride_col_dst, glu_op, stream); + return true; + }; + + // unary fold over launch + return (launch(std::integral_constant{}) || ...); +} + bool ggml_sycl_mul_mat_vec_q_glu_reorder(enum ggml_type src0_type, enum ggml_glu_op glu_op, const void * vx, const void * vgate, const void * vy, float * dst, int ncols, int nrows, int ncols_dst, int stride_col_y_bytes, int stride_col_dst, dpct::queue_ptr stream) { - if (src0_type != GGML_TYPE_Q4_K) { - return false; - } if (glu_op != GGML_GLU_OP_SWIGLU && glu_op != GGML_GLU_OP_GEGLU) { return false; } - using vec_dot = reorder_vec_dot_q_sycl; - - switch (ncols_dst) { - case 1: - launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, - stride_col_dst, glu_op, stream); - return true; - case 2: - if (nrows >= Q4_K_MMVQ_ROW_PAIR_MIN_NROWS) { - launch_mul_mat_vec_q_reorder_glu_impl(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream); - } else { - launch_mul_mat_vec_q_reorder_glu_impl(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, glu_op, stream); - } - return true; - case 3: - launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, - stride_col_dst, glu_op, stream); - return true; - case 4: - launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, - stride_col_dst, glu_op, stream); - return true; - case 5: - launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, - stride_col_dst, glu_op, stream); - return true; - case 6: - launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, - stride_col_dst, glu_op, stream); - return true; - case 7: - launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, - stride_col_dst, glu_op, stream); - return true; - case 8: - launch_mul_mat_vec_q_reorder_glu(vx, vgate, vy, dst, ncols, nrows, stride_col_y_bytes, - stride_col_dst, glu_op, stream); - return true; + switch (src0_type) { + case GGML_TYPE_Q4_K: + return mul_mat_vec_q_glu_reorder_ncols( + glu_op, vx, vgate, vy, dst, ncols, nrows, ncols_dst, stride_col_y_bytes, stride_col_dst, stream); + case GGML_TYPE_Q5_K: + // fusion declines q5_K past 5 columns + return mul_mat_vec_q_glu_reorder_ncols( + glu_op, vx, vgate, vy, dst, ncols, nrows, ncols_dst, stride_col_y_bytes, stride_col_dst, stream); default: return false; } diff --git a/ggml/src/ggml-sycl/mmvq.hpp b/ggml/src/ggml-sycl/mmvq.hpp index 7fb9cf6f82e6..3a0ef07d67e4 100644 --- a/ggml/src/ggml-sycl/mmvq.hpp +++ b/ggml/src/ggml-sycl/mmvq.hpp @@ -15,6 +15,12 @@ #include "common.hpp" +// q5_K multi-column MMVQ shares weights across columns, pairs rows and fuses gate/up in the reorder +// layout: faster on Xe2 (BMG), so untested archs keep the per-column kernel +inline bool ggml_sycl_q5_k_mmvq_reuse(int device) { + const gpu_arch arch = ggml_sycl_info().devices[device].hw_info.arch; + return arch == gpu_arch::intel_gpu_bmg_g21 || arch == gpu_arch::intel_gpu_bmg_g31; +} void ggml_sycl_op_mul_mat_vec_q( ggml_backend_sycl_context & ctx, diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index cc6ae6a8d3df..6ae951525e87 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -362,6 +362,10 @@ template <> struct reorder_vec_dot_shared_weights { static constexpr bool value = true; }; +template <> struct reorder_vec_dot_shared_weights { + static constexpr bool value = true; +}; + template struct reorder_vec_dot_shared_activations { static constexpr bool value = false; }; @@ -370,6 +374,10 @@ template <> struct reorder_vec_dot_shared_activations { static constexpr bool value = true; }; +template <> struct reorder_vec_dot_shared_activations { + static constexpr bool value = true; +}; + template <> struct reorder_vec_dot_q_sycl { static constexpr ggml_type gtype = GGML_TYPE_Q4_0; @@ -676,56 +684,72 @@ template <> struct reorder_vec_dot_q_sycl { using q5_k_block = ggml_sycl_reordered::block_q_t; using q5_k_traits = typename q5_k_block::traits; - __dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair ibx_offset, - const std::pair d_offset, const int8_t * q8_1_quant_ptr, - const sycl::half2 * q8_1_ds, const int & iqs) { - const uint8_t * base = static_cast(vbq); - const uint8_t * qs = base + ibx_offset.first; // low 4 bits - const uint8_t * qh_base = base + ibx_offset.second; // high bit - const uint8_t * scs = base + d_offset.first; - const ggml_half2 * dms = reinterpret_cast(base + d_offset.second); + struct weights { + int vl[2]; + int vh[2]; + uint16_t aux[2]; + ggml_half2 dm; + }; + + // same activation layout as Q4_K + static_assert(QR5_K == QR4_K); + using activations = reorder_vec_dot_q_sycl::activations; + + __dpct_inline__ static weights load(const void * __restrict__ vbq, const std::pair ibx_offset, + const std::pair d_offset, const int & iqs) { + const uint8_t * base = static_cast(vbq); + const uint8_t * qs = base + ibx_offset.first; // low 4 bits + const uint8_t * qh_base = base + ibx_offset.second; // high bit + const uint8_t * scs = base + d_offset.first; + const ggml_half2 * dms = reinterpret_cast(base + d_offset.second); const int bq8_offset = QR5_K * ((iqs / 2) / (QI8_1 / 2)); const int * ql_ptr = (const int *) (qs + 16 * bq8_offset + 4 * ((iqs / 2) % 4)); const int * qh_ptr = (const int *) (qh_base + 4 * ((iqs / 2) % 4)); const uint16_t * scales = (const uint16_t *) scs; - int vl[2]; - int vh[2]; - int u[2 * QR5_K]; - float d8[QR5_K]; - - vl[0] = ql_ptr[0]; - vl[1] = ql_ptr[4]; + weights w; + w.vl[0] = ql_ptr[0]; + w.vl[1] = ql_ptr[4]; - vh[0] = qh_ptr[0] >> bq8_offset; - vh[1] = qh_ptr[4] >> bq8_offset; + w.vh[0] = qh_ptr[0] >> bq8_offset; + w.vh[1] = qh_ptr[4] >> bq8_offset; - uint16_t aux[2]; const int j = (QR5_K * ((iqs / 2) / (QI8_1 / 2))) / 2; if (j < 2) { - aux[0] = scales[j + 0] & 0x3f3f; - aux[1] = scales[j + 2] & 0x3f3f; + w.aux[0] = scales[j + 0] & 0x3f3f; + w.aux[1] = scales[j + 2] & 0x3f3f; } else { - aux[0] = ((scales[j + 2] >> 0) & 0x0f0f) | ((scales[j - 2] & 0xc0c0) >> 2); - aux[1] = ((scales[j + 2] >> 4) & 0x0f0f) | ((scales[j - 0] & 0xc0c0) >> 2); + w.aux[0] = ((scales[j + 2] >> 0) & 0x0f0f) | ((scales[j - 2] & 0xc0c0) >> 2); + w.aux[1] = ((scales[j + 2] >> 4) & 0x0f0f) | ((scales[j - 0] & 0xc0c0) >> 2); } - const uint8_t * sc = (const uint8_t *) aux; - const uint8_t * m = sc + 2; + w.dm = *dms; - for (int i = 0; i < QR5_K; ++i) { - const int8_t* quant_base_ptr = q8_1_quant_ptr + (bq8_offset + i) * QK8_1; - sycl::half2 ds_values = *(q8_1_ds + bq8_offset + i); + return w; + } - d8[i] = ds_values[0]; + __dpct_inline__ static activations load_activations(const int8_t * q8_1_quant_ptr, + const sycl::half2 * q8_1_ds, const int & iqs) { + return reorder_vec_dot_q_sycl::load_activations(q8_1_quant_ptr, q8_1_ds, iqs); + } - const int * q8 = (const int *) quant_base_ptr + ((iqs / 2) % 4); - u[2 * i + 0] = q8[0]; - u[2 * i + 1] = q8[4]; - } + __dpct_inline__ static float apply(const weights & w, const activations & a) { + const uint8_t * sc = (const uint8_t *) w.aux; + const uint8_t * m = sc + 2; + + return vec_dot_q5_K_q8_1_impl_vmmq(w.vl, w.vh, a.u, sc, m, w.dm, a.d8); + } - return vec_dot_q5_K_q8_1_impl_vmmq(vl, vh, u, sc, m, *dms, d8); + __dpct_inline__ static float dot(const weights & w, const int8_t * q8_1_quant_ptr, + const sycl::half2 * q8_1_ds, const int & iqs) { + return apply(w, load_activations(q8_1_quant_ptr, q8_1_ds, iqs)); + } + + __dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair ibx_offset, + const std::pair d_offset, const int8_t * q8_1_quant_ptr, + const sycl::half2 * q8_1_ds, const int & iqs) { + return dot(load(vbq, ibx_offset, d_offset, iqs), q8_1_quant_ptr, q8_1_ds, iqs); } }; From d4d82d67f48bb8f51e9a3518063e49d9ee54f58b Mon Sep 17 00:00:00 2001 From: Mark Todorovich Date: Thu, 8 Oct 2026 23:04:09 -0700 Subject: [PATCH 003/400] hexagon: fix IM2COL patch-embed DMA ring overflow (#30189) * hexagon: fix IM2COL patch-embed DMA ring overflow The exact-tiling (stride == kernel, no pad/dilation) IM2COL DMA kernel issues IC*KH DDR->VTCM descriptors per output row without checking the return value of dma_queue_push(), and then pops IC*KH times. The per-thread DMA ring holds 256 entries and a push into a full ring returns false and drops the transfer, so for IC*KH > 255 the remaining rows of the VTCM staging buffer were never written and stale data (often NaN/inf) leaked into the output. Solution is to retire the oldest descriptor when the ring is full, just as the blocked kernel in the same file already does, and wait with dma_queue_flush(). For testing, added exact-tiling test cases with IC*KH > 256 (2D 1x1, 2D 2x2 patch embed and 1D, F16 and F32 dst), which fail on HTP without this fix. * Apply suggestion from @max-krasnyansky --------- Co-authored-by: Max Krasnyansky --- ggml/src/ggml-hexagon/htp/im2col-ops.c | 11 +++++++---- tests/test-backend-ops.cpp | 8 ++++++++ 2 files changed, 15 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-hexagon/htp/im2col-ops.c b/ggml/src/ggml-hexagon/htp/im2col-ops.c index 2e05cf3e1002..b99d711c964d 100644 --- a/ggml/src/ggml-hexagon/htp/im2col-ops.c +++ b/ggml/src/ggml-hexagon/htp/im2col-ops.c @@ -351,12 +351,15 @@ IM2COL_BLOCKED_DMA_BODY(im2col_blocked_dma_f32_thread, float, hvx_copy_f32_uu, const dma_addr_t vsrc = ok \ ? (src_data + (size_t) ((in * IC + iic) * IH + iih) * IW * sizeof(float)) \ : src_data; \ - dma_queue_push(dma_q, dma_make_data(vdst, vsrc), \ - IW * sizeof(float), IW * sizeof(float), IW * sizeof(float), ok ? 1 : 0); \ + /* IC*KH descriptors per row can exceed the ring capacity: retire the oldest when full */ \ + while (!dma_queue_push(dma_q, dma_make_data(vdst, vsrc), \ + IW * sizeof(float), IW * sizeof(float), IW * sizeof(float), \ + ok ? 1 : 0)) { \ + dma_queue_pop(dma_q); \ + } \ } \ } \ - for (uint32_t i = 0; i < IC * KH; i++) \ - dma_queue_pop(dma_q); \ + dma_queue_flush(dma_q); \ htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_COMP, r); \ for (uint32_t iow = 0; iow < OW; iow++) { \ DST_CTYPE * dst_patch = dstb + (uint64_t) iow * patch_stride; \ diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 7b7e6cae10da..9a8d46bfdb4d 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9789,6 +9789,14 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F32, GGML_TYPE_F32, {2, 2, 1536, 729}, {2, 2, 1536, 4096}, 1, 1, 0, 0, 1, 1, true)); test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_F16, {128, 128, 1, 2}, {32, 33, 1, 2}, 1, 1, 1, 1, 1, 1, true)); test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_F16, {128, 128, 2, 1}, {33, 34, 2, 1}, 1, 1, 1, 1, 1, 1, true)); + // non-overlapping (stride == kernel, no pad/dilation) with IC*KH well above 256, e.g. SD1.5 1x1 convs + for (ggml_type dst_type : {GGML_TYPE_F16, GGML_TYPE_F32}) { + test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, dst_type, {64, 64, 320, 1}, {1, 1, 320, 1}, 1, 1, 0, 0, 1, 1, true)); + test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, dst_type, {16, 16, 1280, 2}, {1, 1, 1280, 1}, 1, 1, 0, 0, 1, 1, true)); + test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, dst_type, {8, 8, 2560, 1}, {1, 1, 2560, 1}, 1, 1, 0, 0, 1, 1, true)); + test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, dst_type, {32, 32, 192, 1}, {2, 2, 192, 1}, 2, 2, 0, 0, 1, 1, true)); + test_cases.emplace_back(new test_im2col(GGML_TYPE_F32, GGML_TYPE_F16, dst_type, {64, 512, 1, 1}, {2, 512, 1, 1}, 2, 0, 0, 0, 1, 0, false)); + } // im2col 3D test_cases.emplace_back(new test_im2col_3d(GGML_TYPE_F32, GGML_TYPE_F32, GGML_TYPE_F32)); From 8a1a9b5126126e5228b95fa909d4b08fac65e8b3 Mon Sep 17 00:00:00 2001 From: ynankani Date: Fri, 9 Oct 2026 06:32:52 +0000 Subject: [PATCH 004/400] CUDA: pass src1 precision to host MMQ config helpers (#30168) * CUDA: pass src1 precision to host MMQ config helpers Signed-off-by: ynankani * make src1 prec explicit Signed-off-by: ynankani --------- Signed-off-by: ynankani --- ggml/src/ggml-cuda/mmq-load-tiles.cuh | 128 ++++++++--------- ggml/src/ggml-cuda/mmq-vec-dot.cuh | 80 +++++------ ggml/src/ggml-cuda/mmq.cu | 44 +++--- ggml/src/ggml-cuda/mmq.cuh | 196 +++++++++++++------------- 4 files changed, 223 insertions(+), 225 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index d5654868bc08..0e59c0347b72 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -7,9 +7,9 @@ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q1_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -98,9 +98,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -187,9 +187,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -250,9 +250,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_1( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -313,9 +313,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -393,9 +393,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_1( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -471,9 +471,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q8_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -537,9 +537,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_K( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -598,9 +598,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q3_K( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -711,9 +711,9 @@ static __device__ __forceinline__ int unpack_scales_q45_K(const int * scales, co template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_K( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -822,9 +822,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q5_K( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -946,9 +946,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q6_K( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1036,9 +1036,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq1_s( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1098,9 +1098,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_xxs( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1162,9 +1162,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_xs( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1227,9 +1227,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq2_s( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1295,9 +1295,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq3_xxs( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1359,9 +1359,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq3_s( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1428,9 +1428,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_xs( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1495,9 +1495,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_iq4_nl( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1564,9 +1564,9 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_mxfp4( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) int * x_qs = (int *) x_tile; @@ -1670,7 +1670,7 @@ template static __device__ __forceinline_ } } -template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4( +template static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4( const char * __restrict__ x, int * __restrict__ x_tile, const int kb0, const int i_max, const int stride) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size; diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index b39e4d579a8b..a75dd046e5c3 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh @@ -10,8 +10,8 @@ using namespace ggml_cuda_mma; template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q4_0_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_0, I); const int * x_qs = (const int *) x; @@ -60,8 +60,8 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q4_1_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_1, I); const int * x_qs = (const int *) x; @@ -110,8 +110,8 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I); const int * x_qs = (const int *) x; @@ -148,8 +148,8 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma( typedef tile<16, 8, int, input_layout> tile_B; typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -203,8 +203,8 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma( typedef tile< 8, 8, int> tile_B; typedef tile<16, 8, int> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -281,8 +281,8 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma( template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_1_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q5_1, I); const int * x_qs = (const int *) x; @@ -318,8 +318,8 @@ template static __device__ __forceinline_ typedef tile<16, 8, int, input_layout> tile_B; typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -368,8 +368,8 @@ template static __device__ __forceinline_ typedef tile< 8, 8, int> tile_B; typedef tile<16, 8, int> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -442,8 +442,8 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(type, I); const int * x_qs = (const int *) x; @@ -474,7 +474,7 @@ template static __device__ __forceinline_ } // Used for Q3_K, IQ2_S, and IQ2_XS: -template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma( +template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) constexpr data_layout input_layout = get_input_data_layout(); @@ -483,7 +483,7 @@ template tile_C; constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, prec_src1); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, prec_src1); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -533,7 +533,7 @@ template tile_C; constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, prec_src1); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, prec_src1); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -610,8 +610,8 @@ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q2_K_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q2_K, I); const int * x_qs = (const int *) x; @@ -680,8 +680,8 @@ template static __device__ __forceinline_ typedef tile<16, 4, int, input_layout> tile_B; typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -749,8 +749,8 @@ template static __device__ __forceinline_ typedef tile< 8, 4, int> tile_B; typedef tile<16, 8, int> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -870,8 +870,8 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q3_K_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q3_K, I); const int * x_qs = (const int *) x; @@ -905,8 +905,8 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q4_K_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_K, I); const int * x_qs = (const int *) x; @@ -940,8 +940,8 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q5_K_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q5_K, I); const int * x_qs = (const int *) x; @@ -975,8 +975,8 @@ template static __device__ __forceinline_ template static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q6_K_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q8) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q8); constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q6_K, I); const int * x_qs = (const int *) x; @@ -1015,8 +1015,8 @@ template static __device__ __forceinline_ typedef tile<16, 4, int, input_layout> tile_B; typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -1066,8 +1066,8 @@ template static __device__ __forceinline_ typedef tile< 8, 4, int> tile_B; typedef tile<16, 8, int> tile_C; - constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q8); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q8); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); @@ -1181,7 +1181,7 @@ template static __device__ __forceinline_ typedef tile<16, 8, float> tile_C; constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q4); - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, GGML_PREC_Q4); constexpr int ntx = rows_per_warp / tile_C::I; constexpr int nfrags = MMQ_TILE_NE_K / tile_A::J; diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 027a7e60da23..91b89999ef8e 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -8,66 +8,66 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream, const ggml_prec prec_src1) { switch (args.type_x) { case GGML_TYPE_Q1_0: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q2_0: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q4_0: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q4_1: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q5_0: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q5_1: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q8_0: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; // ----------------------------------------------------------------------- case GGML_TYPE_Q2_K: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q3_K: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q4_K: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q5_K: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_Q6_K: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; // ----------------------------------------------------------------------- case GGML_TYPE_IQ1_S: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_IQ2_XXS: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_IQ2_XS: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_IQ2_S: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_IQ3_XXS: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_IQ3_S: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_IQ4_XS: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_IQ4_NL: - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; // ----------------------------------------------------------------------- case GGML_TYPE_MXFP4: @@ -76,14 +76,14 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con mul_mat_q_case(ctx, args, stream); break; } - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; case GGML_TYPE_NVFP4: if (prec_src1 == GGML_PREC_Q4) { mul_mat_q_case(ctx, args, stream); break; } - mul_mat_q_case(ctx, args, stream); + mul_mat_q_case(ctx, args, stream); break; default: GGML_ABORT("fatal error"); diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 198aab4b5044..d3e7b75ea7c8 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -227,7 +227,7 @@ struct ggml_cuda_mmq_config { #undef CASE -static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1 = GGML_PREC_Q8) { +static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { if (GGML_CUDA_CC_IS_AMD(cc)) { if (GGML_CUDA_CC_IS_GCN(cc)) { return ggml_cuda_mmq_get_config_gcn(type, J, fallback); @@ -262,7 +262,7 @@ static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type ty return ggml_cuda_mmq_get_config_pascal_older(type, J, fallback); } -static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { #ifdef GGML_USE_HIP #ifdef GCN return ggml_cuda_mmq_get_config_gcn(type, J, fallback); @@ -295,79 +295,77 @@ static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_t GGML_UNUSED_VARS(type, J, fallback, prec_src1); } -// FIXME all of the host functions are missing prec_src1, this can lead to inconsitent behavior. - -static __host__ int ggml_cuda_mmq_get_type(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_config(type, J, fallback, cc).type; +static __host__ int ggml_cuda_mmq_get_type(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1).type; } -static constexpr __device__ int ggml_cuda_mmq_get_type(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_type(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).type; } -static constexpr __device__ int ggml_cuda_mmq_get_nthreads(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_nthreads(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).nthreads; } -static constexpr __device__ int ggml_cuda_mmq_get_occupancy(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_occupancy(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).occupancy; } -static __host__ int ggml_cuda_mmq_get_I(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_config(type, J, fallback, cc).I; +static __host__ int ggml_cuda_mmq_get_I(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1).I; } -static constexpr __device__ int ggml_cuda_mmq_get_I(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_I(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).I; } -static __host__ int ggml_cuda_mmq_get_J(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_config(type, J, fallback, cc).J; +static __host__ int ggml_cuda_mmq_get_J(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1).J; } -static constexpr __device__ int ggml_cuda_mmq_get_J(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_J(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).J; } -static __host__ ggml_cuda_mmq_sram_layout ggml_cuda_mmq_get_sram_layout(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_config(type, J, fallback, cc).sram_layout; +static __host__ ggml_cuda_mmq_sram_layout ggml_cuda_mmq_get_sram_layout(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1).sram_layout; } -static constexpr __device__ ggml_cuda_mmq_sram_layout ggml_cuda_mmq_get_sram_layout(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ ggml_cuda_mmq_sram_layout ggml_cuda_mmq_get_sram_layout(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).sram_layout; } -static __host__ int ggml_cuda_mmq_get_K_vram(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_config(type, J, fallback, cc).K_vram; +static __host__ int ggml_cuda_mmq_get_K_vram(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1).K_vram; } -static constexpr __device__ int ggml_cuda_mmq_get_K_vram(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_K_vram(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).K_vram; } -static __host__ bool ggml_cuda_mmq_get_stream_k(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_config(type, J, fallback, cc).stream_k; +static __host__ bool ggml_cuda_mmq_get_stream_k(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1).stream_k; } -static constexpr __device__ bool ggml_cuda_mmq_get_stream_k(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ bool ggml_cuda_mmq_get_stream_k(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).stream_k; } -static __host__ int ggml_cuda_mmq_get_fallback(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_config(type, J, fallback, cc).fallback; +static __host__ int ggml_cuda_mmq_get_fallback(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, cc, prec_src1).fallback; } -static constexpr __device__ int ggml_cuda_mmq_get_fallback(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_fallback(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).fallback; } // --------------------------------------------------------------------------------------------- -static __host__ int ggml_cuda_mmq_get_sram_stride(const ggml_type type, const int J, const bool fallback, const int cc) { - return ggml_cuda_mmq_get_sram_stride(ggml_cuda_mmq_get_sram_layout(type, J, fallback, cc)); +static __host__ int ggml_cuda_mmq_get_sram_stride(const ggml_type type, const int J, const bool fallback, const int cc, const ggml_prec prec_src1) { + return ggml_cuda_mmq_get_sram_stride(ggml_cuda_mmq_get_sram_layout(type, J, fallback, cc, prec_src1)); } -static constexpr __device__ int ggml_cuda_mmq_get_sram_stride(ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8) { +static constexpr __device__ int ggml_cuda_mmq_get_sram_stride(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { return ggml_cuda_mmq_get_sram_stride(ggml_cuda_mmq_get_sram_layout(type, J, fallback, prec_src1)); } @@ -375,8 +373,8 @@ static __host__ bool ggml_cuda_mmq_needs_fallback(const int64_t nrows_x) { return nrows_x % 128 != 0; } -static constexpr __device__ int ggml_cuda_mmq_get_rows_per_warp(ggml_type type, int J, bool fallback) { - return ggml_cuda_mmq_get_config(type, J, fallback).rows_per_warp(); +static constexpr __device__ int ggml_cuda_mmq_get_rows_per_warp(ggml_type type, int J, bool fallback, ggml_prec prec_src1) { + return ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).rows_per_warp(); } #define MMQ_DP4A_TXS_Q4_0 tile_x_sizes{I*MMQ_TILE_NE_K + I, I*MMQ_TILE_NE_K/QI4_0 + I/QI4_0, 0} @@ -432,12 +430,12 @@ static __host__ int ggml_cuda_mmq_get_nbytes_shared_x(const ggml_cuda_mmq_config #include "mmq-load-tiles.cuh" #include "mmq-vec-dot.cuh" -template static __device__ __forceinline__ void ggml_cuda_mmq_write_back_dp4a( +template static __device__ __forceinline__ void ggml_cuda_mmq_write_back_dp4a( const float * __restrict__ sum, const int32_t * __restrict__ ids_dst, float * __restrict__ dst, const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max) { constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size; - constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback); + constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size; + constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1); const bool y_scale_used = y_scale != nullptr; @@ -471,7 +469,7 @@ template static __device__ __forceinline_ } } -template +template static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( const float * __restrict__ sum, const int * __restrict__ ids_dst, float * __restrict__ dst, const float * __restrict__ y_scale, const int stride, const int i_max, const int j_max) { @@ -482,7 +480,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma( typedef tile<16, 8, int> tile_C; #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback); + constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback, prec_src1); constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I); @@ -536,7 +534,7 @@ struct ggml_cuda_mmq_util_funcs { vdr(vdr), load_tiles(load_tiles), vec_dot(vec_dot), write_back(write_back) {} }; -template +template static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_funcs() { if (!ggml_cuda_mmq_get_config(type, J, fallback, prec_src1).use_mma_data_layout()) { switch (type) { @@ -545,136 +543,136 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func VDR_Q1_0_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q1_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q2_0: return ggml_cuda_mmq_util_funcs( VDR_Q2_0_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q2_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q4_0: return ggml_cuda_mmq_util_funcs( VDR_Q4_0_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q4_0, ggml_cuda_mmq_vec_dot_q4_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q4_1: return ggml_cuda_mmq_util_funcs( VDR_Q4_1_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q4_1, ggml_cuda_mmq_vec_dot_q4_1_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q5_0: return ggml_cuda_mmq_util_funcs( VDR_Q5_0_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q5_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q5_1: return ggml_cuda_mmq_util_funcs( VDR_Q5_1_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q5_1, ggml_cuda_mmq_vec_dot_q8_1_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q8_0: return ggml_cuda_mmq_util_funcs( VDR_Q8_0_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q8_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); // --------------------------------------------------------------------------------------------- case GGML_TYPE_Q2_K: return ggml_cuda_mmq_util_funcs( VDR_Q2_K_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q2_K, ggml_cuda_mmq_vec_dot_q2_K_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q3_K: return ggml_cuda_mmq_util_funcs( VDR_Q3_K_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q3_K, ggml_cuda_mmq_vec_dot_q3_K_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q4_K: return ggml_cuda_mmq_util_funcs( VDR_Q4_K_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q4_K, ggml_cuda_mmq_vec_dot_q4_K_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q5_K: return ggml_cuda_mmq_util_funcs( VDR_Q5_K_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q5_K, ggml_cuda_mmq_vec_dot_q5_K_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_Q6_K: return ggml_cuda_mmq_util_funcs( VDR_Q6_K_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_q6_K, ggml_cuda_mmq_vec_dot_q6_K_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); // --------------------------------------------------------------------------------------------- case GGML_TYPE_IQ1_S: return ggml_cuda_mmq_util_funcs( VDR_IQ1_S_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq1_s, ggml_cuda_mmq_vec_dot_q8_1_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ2_XXS: return ggml_cuda_mmq_util_funcs( VDR_IQ2_XXS_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq2_xxs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ2_XS: return ggml_cuda_mmq_util_funcs( VDR_IQ2_XS_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq2_xs, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ2_S: return ggml_cuda_mmq_util_funcs( VDR_IQ2_S_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq2_s, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ3_XXS: return ggml_cuda_mmq_util_funcs( VDR_IQ3_XXS_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq3_xxs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ3_S: return ggml_cuda_mmq_util_funcs( VDR_IQ3_S_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq3_s, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ4_XS: return ggml_cuda_mmq_util_funcs( VDR_IQ4_XS_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq4_xs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_IQ4_NL: return ggml_cuda_mmq_util_funcs( VDR_IQ4_NL_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_iq4_nl, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); // --------------------------------------------------------------------------------------------- case GGML_TYPE_MXFP4: return ggml_cuda_mmq_util_funcs( VDR_MXFP4_Q8_1_MMQ, ggml_cuda_mmq_load_tiles_mxfp4, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); case GGML_TYPE_NVFP4: return ggml_cuda_mmq_util_funcs( VDR_NVFP4_Q8_1_MMQ, - ggml_cuda_mmq_load_tiles_nvfp4, + ggml_cuda_mmq_load_tiles_nvfp4, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a, - ggml_cuda_mmq_write_back_dp4a); + ggml_cuda_mmq_write_back_dp4a); default: return ggml_cuda_mmq_util_funcs(1, nullptr, nullptr, nullptr); } @@ -690,7 +688,7 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func -1, ggml_cuda_mmq_load_tiles_mxfp4_fp4, ggml_cuda_mmq_vec_dot_fp4_fp4_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); } break; case GGML_TYPE_NVFP4: @@ -699,7 +697,7 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func -1, ggml_cuda_mmq_load_tiles_nvfp4_nvfp4, ggml_cuda_mmq_vec_dot_fp4_fp4_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); } break; default: @@ -715,164 +713,164 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func -1, ggml_cuda_mmq_load_tiles_q1_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q2_0: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q2_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q4_0: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q4_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q4_1: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q4_1, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q5_0: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q5_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q5_1: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q5_1, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q8_0: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q8_0, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); // --------------------------------------------------------------------------------------------- case GGML_TYPE_Q2_K: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q2_K, ggml_cuda_mmq_vec_dot_q2_K_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q3_K: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q3_K, - ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q4_K: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q4_K, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q5_K: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q5_K, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_Q6_K: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_q6_K, ggml_cuda_mmq_vec_dot_q6_K_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); // --------------------------------------------------------------------------------------------- case GGML_TYPE_IQ1_S: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq1_s, ggml_cuda_mmq_vec_dot_q8_1_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ2_XXS: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq2_xxs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ2_XS: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq2_xs, - ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ2_S: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq2_s, - ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ3_XXS: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq3_xxs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ3_S: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq3_s, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ4_XS: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq4_xs, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_IQ4_NL: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_iq4_nl, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); // --------------------------------------------------------------------------------------------- case GGML_TYPE_MXFP4: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_mxfp4, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); case GGML_TYPE_NVFP4: return ggml_cuda_mmq_util_funcs( -1, ggml_cuda_mmq_load_tiles_nvfp4, ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma, - ggml_cuda_mmq_write_back_mma); + ggml_cuda_mmq_write_back_mma); default: return ggml_cuda_mmq_util_funcs(1, nullptr, nullptr, nullptr); } } -template +template static constexpr __device__ int ggml_cuda_mmq_get_vdr() { return ggml_cuda_mmq_get_util_funcs().vdr; } -template +template static constexpr __device__ ggml_cuda_mmq_load_tiles_t ggml_cuda_mmq_get_load_tiles() { return ggml_cuda_mmq_get_util_funcs().load_tiles; } -template +template static constexpr __device__ ggml_cuda_mmq_vec_dot_t ggml_cuda_mmq_get_vec_dot() { return ggml_cuda_mmq_get_util_funcs().vec_dot; } -template +template static constexpr __device__ ggml_cuda_mmq_write_back_t ggml_cuda_mmq_get_write_back() { return ggml_cuda_mmq_get_util_funcs().write_back; } // --------------------------------------------------------------------------------------------- -template +template static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const int offset_x, const int * __restrict__ y, const int * __restrict__ ids_dst, float * __restrict__ dst, float * __restrict__ tmp_fixup, @@ -953,7 +951,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( // The mul_mat_q kernel implements "stream-k" work partitioning as described in https://arxiv.org/abs/2301.03598 -template +template __launch_bounds__(ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1), ggml_cuda_mmq_get_occupancy(type, J, fallback, prec_src1)) static __global__ void mul_mat_q( const char * __restrict__ x, const int * __restrict__ y, const int32_t * __restrict__ ids_dst, @@ -1240,7 +1238,7 @@ static __global__ void mul_mat_q( tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); } -template +template __launch_bounds__(ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1)/2, 1) static __global__ void mul_mat_q_stream_k_fixup( const int32_t * __restrict__ ids_dst, const int32_t * __restrict__ expert_bounds, float * __restrict__ dst, @@ -1395,7 +1393,7 @@ static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const i return nbs_ids + nbs_x + GGML_PAD(nbs_y, config.nthreads*sizeof(int)); } -template +template static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { const int id = ggml_cuda_get_device(); const int cc = ggml_cuda_info().devices[id].cc; @@ -1477,7 +1475,7 @@ static void launch_mul_mat_q(ggml_backend_cuda_context & ctx, const mmq_args & a ntx_fd); } -template +template void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { switch (args.J_best) { case 8: @@ -1535,7 +1533,7 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args, } } -template +template void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { if (ggml_cuda_mmq_needs_fallback(args.nrows_x)) { constexpr bool fallback = true; @@ -1547,7 +1545,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda } #define DECL_MMQ_CASE(type) \ - template void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) \ + template void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) \ // FP4 variant: uses native FP4 MMA instead of keeping src1 at Q8_1. #define DECL_MMQ_CASE_W4A4(type) \ From 1f8fa52318b594b373e7447a3c2bb36cdba9515f Mon Sep 17 00:00:00 2001 From: Pascal Date: Fri, 9 Oct 2026 10:15:32 +0200 Subject: [PATCH 005/400] metal: add the 128/96 flash attention kernels, accept every compiled pair (#30209) --- ggml/src/ggml-metal/ggml-metal-device.m | 11 +++++++---- ggml/src/ggml-metal/ggml-metal-ops.cpp | 1 + ggml/src/ggml-metal/ggml-metal-tuning.cpp | 3 +++ ggml/src/ggml-metal/kernels/fa_f16.metal | 2 ++ ggml/src/ggml-metal/kernels/fa_f32.metal | 1 + ggml/src/ggml-metal/kernels/fa_q4_0.metal | 1 + ggml/src/ggml-metal/kernels/fa_q4_1.metal | 1 + ggml/src/ggml-metal/kernels/fa_q5_0.metal | 1 + ggml/src/ggml-metal/kernels/fa_q5_1.metal | 1 + ggml/src/ggml-metal/kernels/fa_q8_0.metal | 1 + ggml/src/ggml-metal/kernels/fa_vec_f16.metal | 4 ++++ ggml/src/ggml-metal/kernels/fa_vec_f32.metal | 1 + ggml/src/ggml-metal/kernels/fa_vec_q4_0.metal | 3 +++ ggml/src/ggml-metal/kernels/fa_vec_q4_1.metal | 3 +++ ggml/src/ggml-metal/kernels/fa_vec_q5_0.metal | 3 +++ ggml/src/ggml-metal/kernels/fa_vec_q5_1.metal | 3 +++ ggml/src/ggml-metal/kernels/fa_vec_q8_0.metal | 3 +++ 17 files changed, 39 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 827d92622d7c..13aecbdd320a 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1741,10 +1741,13 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te op->src[0]->ne[0] != 576) { return false; } - if (op->src[1]->ne[0] == 72 && op->src[1]->ne[0] != op->src[2]->ne[0]) { - return false; - } - if (op->src[1]->ne[0] < op->src[2]->ne[0]) { + // the kernels exist for K == V and for these K > V pairs only + if (op->src[1]->ne[0] != op->src[2]->ne[0] && + !(op->src[1]->ne[0] == 96 && op->src[2]->ne[0] == 64) && + !(op->src[1]->ne[0] == 128 && op->src[2]->ne[0] == 96) && + !(op->src[1]->ne[0] == 192 && op->src[2]->ne[0] == 128) && + !(op->src[1]->ne[0] == 320 && op->src[2]->ne[0] == 256) && + !(op->src[1]->ne[0] == 576 && op->src[2]->ne[0] == 512)) { return false; } if (op->src[1]->type != op->src[2]->type) { diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 597d5300c071..86b9fe4ef360 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3180,6 +3180,7 @@ static int ggml_metal_op_flash_attn_ext_n_kv_max_sparse(const ggml_tensor * op) (dk == 96 && dv == 96) || (dk == 96 && dv == 64) || (dk == 128 && dv == 128) || + (dk == 128 && dv == 96) || (dk == 192 && dv == 128) || (dk == 192 && dv == 192) || (dk == 256 && dv == 256) || diff --git a/ggml/src/ggml-metal/ggml-metal-tuning.cpp b/ggml/src/ggml-metal/ggml-metal-tuning.cpp index 217ad43b311c..0fe50277ea78 100644 --- a/ggml/src/ggml-metal/ggml-metal-tuning.cpp +++ b/ggml/src/ggml-metal/ggml-metal-tuning.cpp @@ -40,6 +40,9 @@ int fa_vec_baseline_ne(int dk, int dv) { if (dk == 128 && dv == 128) { return 1; } + if (dk == 128 && dv == 96) { + return 4; + } if (dk == 192 && dv == 192) { return 2; } diff --git a/ggml/src/ggml-metal/kernels/fa_f16.metal b/ggml/src/ggml-metal/kernels/fa_f16.metal index 2dfd3799d137..0f44b419bdea 100644 --- a/ggml/src/ggml-metal/kernels/fa_f16.metal +++ b/ggml/src/ggml-metal/kernels/fa_f16.metal @@ -44,6 +44,7 @@ template [[host_name("kernel_flash_attn_ext_f16_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_f16_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f16_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f16_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_f16_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f16_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f16_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f16_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; @@ -62,6 +63,7 @@ template [[host_name("kernel_flash_attn_ext_bf16_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_bf16_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_bf16_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_bf16_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_bf16_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_bf16_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_bf16_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_bf16_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; diff --git a/ggml/src/ggml-metal/kernels/fa_f32.metal b/ggml/src/ggml-metal/kernels/fa_f32.metal index 8d38d775229e..449079e8f6e5 100644 --- a/ggml/src/ggml-metal/kernels/fa_f32.metal +++ b/ggml/src/ggml-metal/kernels/fa_f32.metal @@ -41,6 +41,7 @@ template [[host_name("kernel_flash_attn_ext_f32_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_f32_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f32_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f32_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_f32_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f32_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f32_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_f32_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; diff --git a/ggml/src/ggml-metal/kernels/fa_q4_0.metal b/ggml/src/ggml-metal/kernels/fa_q4_0.metal index 49bb4c13cf7b..6b05b5df6bdf 100644 --- a/ggml/src/ggml-metal/kernels/fa_q4_0.metal +++ b/ggml/src/ggml-metal/kernels/fa_q4_0.metal @@ -41,6 +41,7 @@ template [[host_name("kernel_flash_attn_ext_q4_0_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_q4_0_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_0_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_0_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_q4_0_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_0_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_0_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_0_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; diff --git a/ggml/src/ggml-metal/kernels/fa_q4_1.metal b/ggml/src/ggml-metal/kernels/fa_q4_1.metal index 79797ffced9d..87fac701ea92 100644 --- a/ggml/src/ggml-metal/kernels/fa_q4_1.metal +++ b/ggml/src/ggml-metal/kernels/fa_q4_1.metal @@ -41,6 +41,7 @@ template [[host_name("kernel_flash_attn_ext_q4_1_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_q4_1_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_1_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_1_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_q4_1_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_1_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_1_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q4_1_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; diff --git a/ggml/src/ggml-metal/kernels/fa_q5_0.metal b/ggml/src/ggml-metal/kernels/fa_q5_0.metal index 93de514b425b..cd0e31f715a7 100644 --- a/ggml/src/ggml-metal/kernels/fa_q5_0.metal +++ b/ggml/src/ggml-metal/kernels/fa_q5_0.metal @@ -41,6 +41,7 @@ template [[host_name("kernel_flash_attn_ext_q5_0_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_q5_0_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_0_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_0_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_q5_0_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_0_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_0_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_0_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; diff --git a/ggml/src/ggml-metal/kernels/fa_q5_1.metal b/ggml/src/ggml-metal/kernels/fa_q5_1.metal index 3918c571e20c..4d74951acc5f 100644 --- a/ggml/src/ggml-metal/kernels/fa_q5_1.metal +++ b/ggml/src/ggml-metal/kernels/fa_q5_1.metal @@ -41,6 +41,7 @@ template [[host_name("kernel_flash_attn_ext_q5_1_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_q5_1_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_1_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_1_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_q5_1_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_1_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_1_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q5_1_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; diff --git a/ggml/src/ggml-metal/kernels/fa_q8_0.metal b/ggml/src/ggml-metal/kernels/fa_q8_0.metal index 146bbb31c526..bb29a262c1ea 100644 --- a/ggml/src/ggml-metal/kernels/fa_q8_0.metal +++ b/ggml/src/ggml-metal/kernels/fa_q8_0.metal @@ -41,6 +41,7 @@ template [[host_name("kernel_flash_attn_ext_q8_0_dk96_dv96" )]] kernel flash_at template [[host_name("kernel_flash_attn_ext_q8_0_dk96_dv64" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q8_0_dk112_dv112")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q8_0_dk128_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; +template [[host_name("kernel_flash_attn_ext_q8_0_dk128_dv96" )]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q8_0_dk192_dv192")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q8_0_dk192_dv128")]] kernel flash_attn_ext_t kernel_flash_attn_ext; template [[host_name("kernel_flash_attn_ext_q8_0_dk256_dv256")]] kernel flash_attn_ext_t kernel_flash_attn_ext; diff --git a/ggml/src/ggml-metal/kernels/fa_vec_f16.metal b/ggml/src/ggml-metal/kernels/fa_vec_f16.metal index 56b5adb840cd..bea54f9f73fd 100644 --- a/ggml/src/ggml-metal/kernels/fa_vec_f16.metal +++ b/ggml/src/ggml-metal/kernels/fa_vec_f16.metal @@ -56,8 +56,12 @@ template [[host_name("kernel_flash_attn_ext_vec_f16_dk128_dv128_q2_ne4")]] kerne template [[host_name("kernel_flash_attn_ext_vec_f16_dk128_dv128_q4_ne1")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f16_dk128_dv128_q4_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f16_dk128_dv128_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_f16_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_f16_dk128_dv96_q2_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_f16_dk128_dv96_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; #if defined(GGML_METAL_HAS_BF16) template [[host_name("kernel_flash_attn_ext_vec_bf16_dk128_dv128")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_bf16_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; #endif template [[host_name("kernel_flash_attn_ext_vec_f16_dk192_dv192")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f16_dk192_dv192_q1_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; diff --git a/ggml/src/ggml-metal/kernels/fa_vec_f32.metal b/ggml/src/ggml-metal/kernels/fa_vec_f32.metal index 1c6bed0c61cf..18f862966a78 100644 --- a/ggml/src/ggml-metal/kernels/fa_vec_f32.metal +++ b/ggml/src/ggml-metal/kernels/fa_vec_f32.metal @@ -25,6 +25,7 @@ template [[host_name("kernel_flash_attn_ext_vec_f32_dk64_dv64")]] kernel flas template [[host_name("kernel_flash_attn_ext_vec_f32_dk96_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f32_dk96_dv64")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f32_dk128_dv128")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_f32_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f32_dk192_dv192")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f32_dk192_dv128")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_f32_dk256_dv256")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; diff --git a/ggml/src/ggml-metal/kernels/fa_vec_q4_0.metal b/ggml/src/ggml-metal/kernels/fa_vec_q4_0.metal index 7bb99ffd724c..6bf5fad75cf1 100644 --- a/ggml/src/ggml-metal/kernels/fa_vec_q4_0.metal +++ b/ggml/src/ggml-metal/kernels/fa_vec_q4_0.metal @@ -44,6 +44,9 @@ template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk128_dv128_q2_ne4")]] kern template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk128_dv128_q4_ne1")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk128_dv128_q4_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk128_dv128_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk128_dv96_q2_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk128_dv96_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk192_dv192")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk192_dv192_q1_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_0_dk192_dv192_q2_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; diff --git a/ggml/src/ggml-metal/kernels/fa_vec_q4_1.metal b/ggml/src/ggml-metal/kernels/fa_vec_q4_1.metal index 8a6b08bca3fb..37277ee10203 100644 --- a/ggml/src/ggml-metal/kernels/fa_vec_q4_1.metal +++ b/ggml/src/ggml-metal/kernels/fa_vec_q4_1.metal @@ -44,6 +44,9 @@ template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk128_dv128_q2_ne4")]] kern template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk128_dv128_q4_ne1")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk128_dv128_q4_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk128_dv128_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk128_dv96_q2_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk128_dv96_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk192_dv192")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk192_dv192_q1_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q4_1_dk192_dv192_q2_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; diff --git a/ggml/src/ggml-metal/kernels/fa_vec_q5_0.metal b/ggml/src/ggml-metal/kernels/fa_vec_q5_0.metal index 1633c4764a8a..26868a228a40 100644 --- a/ggml/src/ggml-metal/kernels/fa_vec_q5_0.metal +++ b/ggml/src/ggml-metal/kernels/fa_vec_q5_0.metal @@ -44,6 +44,9 @@ template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk128_dv128_q2_ne4")]] kern template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk128_dv128_q4_ne1")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk128_dv128_q4_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk128_dv128_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk128_dv96_q2_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk128_dv96_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk192_dv192")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk192_dv192_q1_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_0_dk192_dv192_q2_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; diff --git a/ggml/src/ggml-metal/kernels/fa_vec_q5_1.metal b/ggml/src/ggml-metal/kernels/fa_vec_q5_1.metal index 2e7db017e621..9d39dc58a8d5 100644 --- a/ggml/src/ggml-metal/kernels/fa_vec_q5_1.metal +++ b/ggml/src/ggml-metal/kernels/fa_vec_q5_1.metal @@ -44,6 +44,9 @@ template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk128_dv128_q2_ne4")]] kern template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk128_dv128_q4_ne1")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk128_dv128_q4_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk128_dv128_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk128_dv96_q2_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk128_dv96_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk192_dv192")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk192_dv192_q1_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q5_1_dk192_dv192_q2_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; diff --git a/ggml/src/ggml-metal/kernels/fa_vec_q8_0.metal b/ggml/src/ggml-metal/kernels/fa_vec_q8_0.metal index e9c6b152209b..7686ccc00126 100644 --- a/ggml/src/ggml-metal/kernels/fa_vec_q8_0.metal +++ b/ggml/src/ggml-metal/kernels/fa_vec_q8_0.metal @@ -44,6 +44,9 @@ template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk128_dv128_q2_ne4")]] kern template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk128_dv128_q4_ne1")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk128_dv128_q4_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk128_dv128_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk128_dv96")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk128_dv96_q2_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; +template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk128_dv96_q4_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk192_dv192")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk192_dv192_q1_ne4")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; template [[host_name("kernel_flash_attn_ext_vec_q8_0_dk192_dv192_q2_ne2")]] kernel flash_attn_ext_vec_t kernel_flash_attn_ext_vec; From 2c564f43dfdb0565745dc71cc7748d2a1305246a Mon Sep 17 00:00:00 2001 From: Pascal Date: Fri, 9 Oct 2026 10:29:22 +0200 Subject: [PATCH 006/400] ci: fix HIP quality check by ignoring the 320/256 FA kernel spill (#30212) --- scripts/hip/gcn-cdna-vgpr-check.py | 1 + 1 file changed, 1 insertion(+) diff --git a/scripts/hip/gcn-cdna-vgpr-check.py b/scripts/hip/gcn-cdna-vgpr-check.py index 6461fccd510b..1f1c2c903bc7 100644 --- a/scripts/hip/gcn-cdna-vgpr-check.py +++ b/scripts/hip/gcn-cdna-vgpr-check.py @@ -83,6 +83,7 @@ def main(): '_ZL18flash_attn_ext_f16ILi576ELi512ELi2ELi16ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', '_ZL18flash_attn_ext_f16ILi576ELi512ELi8ELi4ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', '_ZL18flash_attn_ext_vecILi256ELi1EL9ggml_type2ELS0_2ELb0EEvPKcS2_S2_S2_S2_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS6_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi320ELi256ELi2ELi32ELb0ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', } functions = parse_log_file(log_file) From b013e56a7147dbff5ec03175828442a94086507c Mon Sep 17 00:00:00 2001 From: Aleksander Grygier Date: Fri, 9 Oct 2026 10:34:45 +0200 Subject: [PATCH 007/400] ui : add the models manager (#29583) * ui : add the drawer, sheet and grouped list primitives Add the drawer and sheet overlay components, the toggle and toggle group, the shared searchable input, the collapsible sections and the grouped list, and the near viewport helper they position with. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : add the shared model components and data layer Add the model row pieces shared by the selector and the manager, the models store with its download status feed, the huggingface and migration services, and the model utils, enums and constants they read through. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : rework the models selector around its providers Group the selector components under their own folder, rework the list around the provider grouping, add the mobile trigger and the download item, and derive the selection state in one hook. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : add the models manager Add the manager table with its repo, quant and status rows, the filters and the toolbar, the row actions in a drawer, the downloads section, the manage models dialog and the stories and tests. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : rework the chat form actions and the mobile experience Move the add actions into a drawer and a sheet, float the model actions in a bar on a phone, open the context panel in a drawer, and derive the attachment, reasoning and tools menus in hooks. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : move the mcp servers to the sidebar rail and tidy the dialogs Move the mcp servers dialog to the sidebar rail with its menu entries, even out the dialogs on a phone, and group the mcp and settings stores behind their own modules. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : polish the shell The secondary button gets its own look back and the chat add button its own light surface. Pressable elements get a pointer cursor again, the font rendering smooths on the app shell, and the agent skills stay out of prettier's way. The root layout props probe rework that came with this polish reads the providers api url and stays with the providers change instead. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : wrap the model id classes Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * ui : load the model the new chat CTA picks Start a new chat selected the model but left it unloaded, so the chat opened against a server with nothing in memory. Both CTAs share a load helper now. Assisted-by: pi:llama.cpp/DeepSeek-V4.1-Flash * refactor: Post-review fixes --- tools/ui/.prettierignore | 4 + tools/ui/README.md | 160 ++++--- tools/ui/package-lock.json | 64 ++- tools/ui/package.json | 3 +- tools/ui/src/app.css | 29 +- .../app/chat/ChatForm/ChatForm.svelte | 23 +- .../ChatFormActionAddButton.svelte | 4 +- .../ChatFormActionAddDrawer.svelte | 206 +++++++++ .../ChatFormActionAddDropdown.svelte | 49 +- .../ChatFormActionAddReasoningSubmenu.svelte | 75 ---- .../ChatFormActionAddSheet.svelte | 284 ------------ .../ChatFormActionAddToolsSubmenu.svelte | 4 +- .../ChatFormActionsAdd.svelte | 6 +- .../ChatFormActionModels.svelte | 170 ------- .../ChatFormActionReasoning.svelte | 147 ++++++ .../ChatFormActionRecord.svelte | 2 +- .../ChatFormActions/ChatFormActions.svelte | 47 +- .../ChatFormContextGauge.svelte | 11 +- .../ContextGaugeDetailRow.svelte | 4 +- .../ContextGaugeDetails.svelte | 29 +- .../ContextGaugeDial.svelte | 5 +- .../ContextGaugePopup.svelte | 129 +++--- .../gauge-popup.svelte.ts | 2 +- .../ChatFormPickerCommand.svelte | 12 +- .../ChatMessageAssistantModel.svelte | 18 +- ...atMessageToolCallBlockSearchResults.svelte | 3 +- .../ChatMessageStatistics.svelte | 2 +- .../app/chat/ChatScreen/ChatScreenForm.svelte | 1 + .../ChatScreen/ChatScreenServerError.svelte | 2 +- tools/ui/src/lib/components/app/chat/index.ts | 17 +- .../app/dialogs/DialogManageModels.svelte | 40 ++ .../dialogs/DialogMcpResourcePreview.svelte | 7 +- .../dialogs/DialogMcpResourcesBrowser.svelte | 22 +- .../app/dialogs/DialogMcpServerAddNew.svelte | 7 +- .../app/dialogs/DialogMcpServers.svelte | 4 +- .../app/dialogs/DialogModelInformation.svelte | 408 +++-------------- .../app/dialogs/DialogSettingsChat.svelte | 18 +- .../src/lib/components/app/dialogs/index.ts | 43 +- .../components/app/forms/SearchInput.svelte | 5 +- .../app/misc/CollapsibleRegion.svelte | 83 ++++ .../app/misc/CollapsibleSection.svelte | 111 +++++ .../components/app/misc/GroupedList.svelte | 188 ++++++++ tools/ui/src/lib/components/app/misc/index.ts | 9 + .../components/app/models/ModelAvatar.svelte | 106 +++++ .../components/app/models/ModelBadge.svelte | 4 +- .../app/models/ModelCapabilities.svelte | 63 +++ .../app/models/ModelCapabilityIcons.svelte | 104 ++--- .../components/app/models/ModelContext.svelte | 58 +++ ...svelte => ModelDownloadProgressBar.svelte} | 4 +- .../app/models/ModelDraftSidecars.svelte | 42 ++ .../lib/components/app/models/ModelId.svelte | 47 +- .../app/models/ModelLoadControl.svelte | 107 +++++ .../app/models/ModelOrgAvatar.svelte | 130 ++++++ .../app/models/ModelRowActions.svelte | 75 ++++ .../models/ModelsManager/ModelsManager.svelte | 357 +++++++++++++++ .../ModelsManagerDownloadControl.svelte | 50 +++ .../ModelsManager/ModelsManagerFilters.svelte | 135 ++++++ .../ModelsManagerModelConfiguration.svelte | 110 +++++ ...elsManagerModelConfigurationActions.svelte | 85 ++++ ...delsManagerModelConfigurationHeader.svelte | 111 +++++ ...anagerModelConfigurationInformation.svelte | 197 ++++++++ .../ModelsManagerModelRow.svelte | 104 +++++ .../ModelsManagerModelsTable.svelte | 422 ++++++++++++++++++ .../ModelsManagerQuantRow.svelte | 71 +++ .../ModelsManager/ModelsManagerRepoRow.svelte | 95 ++++ .../ModelsManagerStatusCell.svelte | 43 ++ .../ModelsManagerTableToolbar.svelte | 88 ++++ .../app/models/ModelsManager/row-actions.ts | 61 +++ .../app/models/ModelsManager/utils.ts | 202 +++++++++ .../app/models/ModelsSection.svelte | 97 ++++ .../ModelsSelector/ModelsSelector.svelte | 405 +++++++++++++++++ .../ModelsSelectorDownloadItem.svelte | 130 ++++++ .../ModelsSelectorDrawer.svelte | 53 +++ .../ModelsSelectorDropdown.svelte | 248 ++++++++++ .../ModelsSelector/ModelsSelectorList.svelte | 235 ++++++++++ .../ModelsSelectorOption.svelte | 105 +++++ .../ModelsSelectorTriggerIcon.svelte | 23 + .../app/models/ModelsSelector/index.ts | 10 + .../app/models/ModelsSelectorDropdown.svelte | 417 ----------------- .../app/models/ModelsSelectorList.svelte | 76 ---- .../app/models/ModelsSelectorOption.svelte | 216 --------- .../app/models/ModelsSelectorSheet.svelte | 210 --------- .../discover/ModelsDiscoverAvatar.svelte | 109 ----- .../ui/src/lib/components/app/models/index.ts | 84 +--- .../app/navigation/DropdownMenuActions.svelte | 159 +++++-- .../navigation/DropdownMenuSearchable.svelte | 47 +- .../SidebarNavigation.svelte | 27 +- .../SidebarNavigationActions.svelte | 82 ++-- .../lib/components/ui/button/button.svelte | 4 +- .../ui/dialog/dialog-content.svelte | 4 +- .../components/ui/dialog/dialog-header.svelte | 3 +- .../components/ui/drawer/drawer-close.svelte | 7 + .../ui/drawer/drawer-content.svelte | 38 ++ .../ui/drawer/drawer-description.svelte | 17 + .../drawer-footer.svelte} | 4 +- .../drawer-header.svelte} | 7 +- .../components/ui/drawer/drawer-nested.svelte | 12 + .../ui/drawer/drawer-overlay.svelte | 20 + .../components/ui/drawer/drawer-portal.svelte | 7 + .../components/ui/drawer/drawer-title.svelte | 17 + .../ui/drawer/drawer-trigger.svelte | 7 + .../lib/components/ui/drawer/drawer.svelte | 12 + .../ui/src/lib/components/ui/drawer/index.ts | 38 ++ .../dropdown-menu-content.svelte | 2 +- tools/ui/src/lib/components/ui/sheet/index.ts | 36 -- .../components/ui/sheet/sheet-close.svelte | 7 - .../components/ui/sheet/sheet-content.svelte | 64 --- .../ui/sheet/sheet-description.svelte | 17 - .../components/ui/sheet/sheet-overlay.svelte | 20 - .../components/ui/sheet/sheet-title.svelte | 17 - .../components/ui/sheet/sheet-trigger.svelte | 7 - .../lib/components/ui/toggle-group/index.ts | 10 + .../ui/toggle-group/toggle-group-item.svelte | 35 ++ .../ui/toggle-group/toggle-group.svelte | 75 ++++ .../ui/src/lib/components/ui/toggle/index.ts | 13 + .../lib/components/ui/toggle/toggle.svelte | 54 +++ .../constants/attachment-menu.constants.ts | 17 + .../ui/src/lib/constants/backend.constants.ts | 2 + .../lib/constants/css-classes.constants.ts | 54 +++ tools/ui/src/lib/constants/icons.constants.ts | 59 ++- tools/ui/src/lib/constants/index.ts | 2 + .../lib/constants/model-loading.constants.ts | 22 + .../lib/constants/models-manager.constants.ts | 18 + .../lib/constants/settings-keys.constants.ts | 7 +- .../src/lib/constants/settings.constants.ts | 29 +- .../ui/src/lib/constants/storage.constants.ts | 9 + tools/ui/src/lib/constants/ui.constants.ts | 27 +- tools/ui/src/lib/enums/index.ts | 9 +- tools/ui/src/lib/enums/model.enums.ts | 28 ++ tools/ui/src/lib/enums/ui.enums.ts | 2 + .../lib/hooks/use-attachment-menu.svelte.ts | 2 +- .../lib/hooks/use-chat-form-model.svelte.ts | 121 +++++ .../lib/hooks/use-models-selector.svelte.ts | 140 +++--- .../lib/hooks/use-reasoning-menu.svelte.ts | 4 +- .../src/lib/hooks/use-tools-panel.svelte.ts | 2 +- .../src/lib/services/huggingface.service.ts | 131 +++++- .../ui/src/lib/services/migration.service.ts | 34 ++ tools/ui/src/lib/stores/mcp/index.svelte.ts | 14 +- .../ui/src/lib/stores/models/index.svelte.ts | 302 ++++++++++++- .../ui/src/lib/stores/models/status.svelte.ts | 73 ++- .../src/lib/stores/settings/index.svelte.ts | 2 +- tools/ui/src/lib/stores/ui.svelte.ts | 27 ++ tools/ui/src/lib/types/chat.d.ts | 5 +- tools/ui/src/lib/types/huggingface.d.ts | 7 + tools/ui/src/lib/types/index.ts | 1 + tools/ui/src/lib/types/models.d.ts | 32 ++ tools/ui/src/lib/utils/cors-proxy.ts | 15 + tools/ui/src/lib/utils/formatters.ts | 19 + tools/ui/src/lib/utils/index.ts | 19 +- tools/ui/src/lib/utils/model-families.ts | 47 ++ .../models/utils.ts => utils/model-list.ts} | 78 ++-- tools/ui/src/lib/utils/model-names.ts | 16 + tools/ui/src/lib/utils/near-viewport.ts | 34 ++ .../components/DialogPanelWrapper.svelte | 12 + .../components/ModelsManagerRowWrapper.svelte | 24 + .../components/ModelsManagerWrapper.svelte | 9 + .../ModelsSelectorHookHarness.svelte | 36 ++ .../context-panel-mobile.svelte.test.ts | 42 ++ .../client/dialog-focus-mobile.svelte.test.ts | 30 ++ .../client/mcp-icon-proxy-url.svelte.test.ts | 66 +++ .../client/model-avatar-orgs.svelte.test.ts | 60 +++ .../client/models-manager-rows.svelte.test.ts | 91 ++++ .../models-manager-utils.svelte.test.ts | 164 +++++++ .../models-selector-hidden.svelte.test.ts | 77 ++++ .../models-selector-mobile.svelte.test.ts | 56 +++ .../reasoning-control-mobile.svelte.test.ts | 41 ++ ...odel-org-name-key-migration.svelte.test.ts | 59 +++ .../stories/ModelsManager.stories.svelte | 110 +++++ .../stories/ModelsSelector.stories.svelte | 45 +- .../a11y/ChatScreenForm.a11y.stories.svelte | 21 +- .../a11y/ModelsManager.a11y.stories.svelte | 52 +++ tools/ui/tests/unit/formatters.test.ts | 24 + tools/ui/tests/unit/model-families.test.ts | 54 +++ tools/ui/tests/unit/model-list.test.ts | 104 +++++ 174 files changed, 8361 insertions(+), 2841 deletions(-) create mode 100644 tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDrawer.svelte delete mode 100644 tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddReasoningSubmenu.svelte delete mode 100644 tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte delete mode 100644 tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionReasoning.svelte create mode 100644 tools/ui/src/lib/components/app/dialogs/DialogManageModels.svelte create mode 100644 tools/ui/src/lib/components/app/misc/CollapsibleRegion.svelte create mode 100644 tools/ui/src/lib/components/app/misc/CollapsibleSection.svelte create mode 100644 tools/ui/src/lib/components/app/misc/GroupedList.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelAvatar.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelCapabilities.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelContext.svelte rename tools/ui/src/lib/components/app/models/{discover/DownloadProgressBar.svelte => ModelDownloadProgressBar.svelte} (80%) create mode 100644 tools/ui/src/lib/components/app/models/ModelDraftSidecars.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelLoadControl.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelOrgAvatar.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelRowActions.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManager.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerDownloadControl.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerFilters.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfiguration.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationActions.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationHeader.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationInformation.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelRow.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelsTable.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerQuantRow.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerRepoRow.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerStatusCell.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerTableToolbar.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/row-actions.ts create mode 100644 tools/ui/src/lib/components/app/models/ModelsManager/utils.ts create mode 100644 tools/ui/src/lib/components/app/models/ModelsSection.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelector.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDownloadItem.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDrawer.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDropdown.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorList.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorOption.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorTriggerIcon.svelte create mode 100644 tools/ui/src/lib/components/app/models/ModelsSelector/index.ts delete mode 100644 tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte delete mode 100644 tools/ui/src/lib/components/app/models/ModelsSelectorList.svelte delete mode 100644 tools/ui/src/lib/components/app/models/ModelsSelectorOption.svelte delete mode 100644 tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte delete mode 100644 tools/ui/src/lib/components/app/models/discover/ModelsDiscoverAvatar.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-close.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-content.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-description.svelte rename tools/ui/src/lib/components/ui/{sheet/sheet-header.svelte => drawer/drawer-footer.svelte} (81%) rename tools/ui/src/lib/components/ui/{sheet/sheet-footer.svelte => drawer/drawer-header.svelte} (60%) create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-nested.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-overlay.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-portal.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-title.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer-trigger.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/drawer.svelte create mode 100644 tools/ui/src/lib/components/ui/drawer/index.ts delete mode 100644 tools/ui/src/lib/components/ui/sheet/index.ts delete mode 100644 tools/ui/src/lib/components/ui/sheet/sheet-close.svelte delete mode 100644 tools/ui/src/lib/components/ui/sheet/sheet-content.svelte delete mode 100644 tools/ui/src/lib/components/ui/sheet/sheet-description.svelte delete mode 100644 tools/ui/src/lib/components/ui/sheet/sheet-overlay.svelte delete mode 100644 tools/ui/src/lib/components/ui/sheet/sheet-title.svelte delete mode 100644 tools/ui/src/lib/components/ui/sheet/sheet-trigger.svelte create mode 100644 tools/ui/src/lib/components/ui/toggle-group/index.ts create mode 100644 tools/ui/src/lib/components/ui/toggle-group/toggle-group-item.svelte create mode 100644 tools/ui/src/lib/components/ui/toggle-group/toggle-group.svelte create mode 100644 tools/ui/src/lib/components/ui/toggle/index.ts create mode 100644 tools/ui/src/lib/components/ui/toggle/toggle.svelte create mode 100644 tools/ui/src/lib/constants/backend.constants.ts create mode 100644 tools/ui/src/lib/constants/models-manager.constants.ts create mode 100644 tools/ui/src/lib/hooks/use-chat-form-model.svelte.ts create mode 100644 tools/ui/src/lib/utils/model-families.ts rename tools/ui/src/lib/{components/app/models/utils.ts => utils/model-list.ts} (54%) create mode 100644 tools/ui/src/lib/utils/near-viewport.ts create mode 100644 tools/ui/tests/client/components/DialogPanelWrapper.svelte create mode 100644 tools/ui/tests/client/components/ModelsManagerRowWrapper.svelte create mode 100644 tools/ui/tests/client/components/ModelsManagerWrapper.svelte create mode 100644 tools/ui/tests/client/components/ModelsSelectorHookHarness.svelte create mode 100644 tools/ui/tests/client/context-panel-mobile.svelte.test.ts create mode 100644 tools/ui/tests/client/dialog-focus-mobile.svelte.test.ts create mode 100644 tools/ui/tests/client/mcp-icon-proxy-url.svelte.test.ts create mode 100644 tools/ui/tests/client/model-avatar-orgs.svelte.test.ts create mode 100644 tools/ui/tests/client/models-manager-rows.svelte.test.ts create mode 100644 tools/ui/tests/client/models-manager-utils.svelte.test.ts create mode 100644 tools/ui/tests/client/models-selector-hidden.svelte.test.ts create mode 100644 tools/ui/tests/client/models-selector-mobile.svelte.test.ts create mode 100644 tools/ui/tests/client/reasoning-control-mobile.svelte.test.ts create mode 100644 tools/ui/tests/client/settings-model-org-name-key-migration.svelte.test.ts create mode 100644 tools/ui/tests/stories/ModelsManager.stories.svelte create mode 100644 tools/ui/tests/stories/a11y/ModelsManager.a11y.stories.svelte create mode 100644 tools/ui/tests/unit/formatters.test.ts create mode 100644 tools/ui/tests/unit/model-families.test.ts create mode 100644 tools/ui/tests/unit/model-list.test.ts diff --git a/tools/ui/.prettierignore b/tools/ui/.prettierignore index 635cf99c7e9c..ba47fbe165fa 100644 --- a/tools/ui/.prettierignore +++ b/tools/ui/.prettierignore @@ -19,3 +19,7 @@ test-results # Vendored third party sources, kept byte identical to upstream src/lib/vendors/ + +# Agent skills installed from upstream, kept as installed +.agents/ +skills-lock.json diff --git a/tools/ui/README.md b/tools/ui/README.md index 99abfaa41fcf..04a6ea01657b 100644 --- a/tools/ui/README.md +++ b/tools/ui/README.md @@ -54,10 +54,13 @@ Llama UI supports two server operation modes: ### Multi-Model Support (ROUTER mode) +- **Models manager** - one table for every model the server can serve, split into loaded, downloading, favorites, local and hidden sections, with search, context/capability/modality filters, sorting, load/unload, delete from disk and hide. It opens from the sidebar, from a model row or from the selector. - **Model selector** with Loaded/Available groups - **Automatic loading** - Models load on selection - **Modality validation** - Prevents sending images to non-vision models - **LRU unloading** - Server auto-manages model cache +- **Downloads** - In-flight downloads stay visible above the selector and can be paused, resumed or cancelled from the manager +- **Hub metadata** - Avatars, context length and chat template come from the Hugging Face Hub only while the `Use Hugging Face Hub API for models metadata` setting is on. It is off by default, and the UI then shows what the server reports for `/v1/models` ### Keyboard Shortcuts @@ -144,6 +147,7 @@ proxy: { | ----------------- | ------------------------------- | -------------------------------------------------------- | | **Framework** | SvelteKit + Svelte 5 | Reactive UI with runes (`$state`, `$derived`, `$effect`) | | **UI Components** | shadcn-svelte + bits-ui | Accessible, customizable component library | +| **Drawers** | vaul-svelte | Draggable drawer overlays (mobile pickers, model pane) | | **Styling** | TailwindCSS 4 | Utility-first CSS with design tokens | | **Database** | IndexedDB (Dexie) | Client-side storage for conversations and messages | | **Build** | Vite | Fast bundling with static adapter | @@ -350,94 +354,103 @@ Components are organized in `app/` (application-specific) and `ui/` (shadcn-svel **Dialog Components** (`app/dialogs/`): -| Component | Responsibility | -| ------------------------------- | -------------------------------------------------------- | -| `DialogChatSettings` | Full-screen settings configuration | -| `DialogModelInformation` | Model details (context size, modalities, parallel slots) | -| `DialogChatAttachmentPreview` | Full preview for images, PDFs (text or page view), code | -| `DialogConfirmation` | Generic confirmation for destructive actions | -| `DialogConversationTitleUpdate` | Edit conversation title | +| Component | Responsibility | +| ------------------------------ | -------------------------------------------------------- | +| `DialogChatSettings` | Full-screen settings configuration | +| `DialogManageModels` | Models manager: browse, load, download and delete models | +| `DialogModelInformation` | Model details (context size, modalities, parallel slots) | +| `DialogChatAttachmentsPreview` | Full preview for images, PDFs (text or page view), code | +| `DialogConfirmation` | Generic confirmation for destructive actions | +| `DialogConversationRename` | Edit conversation title | **Server/Model Components** (`app/server/`, `app/models/`): -| Component | Responsibility | -| ------------------- | --------------------------------------------------------- | -| `ServerErrorSplash` | Error display when server is unreachable | -| `ModelsSelector` | Model dropdown with Loaded/Available groups (ROUTER mode) | +| Component | Responsibility | +| ----------------------------------- | ------------------------------------------------------------------------------------------------------------- | +| `ServerErrorSplash` | Error display when server is unreachable | +| `ModelsSelector/` | Model dropdown and mobile sheet: favorites, loaded and local sections, rows windowed and folded into families | +| `ModelsManager/` | Models manager table and model pane: sections, filters, sorting, quants folded per repo, downloads | +| `ModelAvatar` | Org avatar of a model, with the quantizer badge; follows the family grouping of the list | +| `ModelId`, `ModelBadge` | Model name, aliases, tags, quantization and draft sidecar badges | +| `ModelContext`, `ModelCapabilities` | Context window and capability icons, filled from the Hub when it is enabled | **Shared UI Components** (`app/misc/`): -| Component | Responsibility | -| -------------------------------- | ---------------------------------------------------------------- | -| `MarkdownContent` | Markdown rendering with KaTeX, syntax highlighting, copy buttons | -| `SyntaxHighlightedCode` | Code blocks with language detection and highlighting | -| `ActionButton`, `ActionDropdown` | Reusable action buttons and menus | -| `BadgeModality`, `BadgeInfo` | Status and capability badges | +| Component | Responsibility | +| ----------------------------------------- | ----------------------------------------------------------------------- | +| `MarkdownContent` | Markdown rendering with KaTeX, syntax highlighting, copy buttons | +| `SyntaxHighlightedCode` | Code blocks with language detection and highlighting | +| `ActionIcon`, `DropdownMenuActions` | Reusable action button and action menu | +| `BadgesModality`, `BadgeInfo` | Status and capability badges | +| `CollapsibleSection`, `CollapsibleRegion` | Section header with a trigger that expands in place | +| `GroupedList` | Grouped rows with foldable groups, per-group windows and show-more rows | +| `ScrollCarousel`, `TruncatedText` | Horizontal scroll strip and single-line truncation with tooltip | #### Hooks (`src/lib/hooks/`) Hooks are the thin view-layer between components and stores: they own UI concerns (scroll, drag-and-drop, keyboard shortcuts, pickers, selection) and translate store state into view state. -| Hook | Responsibility | -| ------------------------------- | -------------------------------------------------------------- | -| `use-chat-screen-active-model` | Active model resolution + modality capability detection | -| `use-processing-state` | View over `chatStore.processing` for streaming progress/tokens | -| `use-context-gauge` | View over `contextStatsStore` for the context usage gauge | -| `use-models-selector` | Model selector dropdown state (loaded/available groups) | -| `use-tools-panel` | Tools panel state | -| `use-reasoning-menu` | Reasoning-effort menu state | -| `use-attachment-menu` | Attachment menu + modality flags | -| `use-draft-messages` | Per-chat draft message/files persistence | -| `use-chat-form-pickers` | Chat form pickers (commands, mentions) | -| `use-debounced-search` | Shared debounced async search for pickers | -| `use-picker-navigation` | Picker keyboard navigation | -| `use-chat-message-edit-context` | Message edit context (content + extras) | -| `use-chat-screen-drag-and-drop` | Drag-and-drop state machine | -| `use-chat-screen-file-upload` | File upload queue + capability validation | -| `use-chat-screen-scroll` | Scroll container binding + navigation guard | -| `use-auto-scroll` | Auto-scroll controller for streaming | -| `use-marquee-selection` | Shift+click / marquee range selection | -| `use-keyboard-shortcuts` | Global keyboard shortcuts | -| `use-settings-navigation` | Settings section navigation | -| `use-pwa` | PWA install/update + version mismatch detection | +| Hook | Responsibility | +| ------------------------------- | ----------------------------------------------------------------------------------------------- | +| `use-chat-screen-active-model` | Active model resolution + modality capability detection | +| `use-processing-state` | View over `chatStore.processing` for streaming progress/tokens | +| `use-context-gauge` | View over `contextStatsStore` for the context usage gauge | +| `use-models-selector` | Model selector state: hidden models, family grouping, windowed rows, selection and load on pick | +| `use-tools-panel` | Tools panel state | +| `use-reasoning-menu` | Reasoning-effort menu state | +| `use-attachment-menu` | Attachment menu + modality flags | +| `use-draft-messages` | Per-chat draft message/files persistence | +| `use-chat-form-pickers` | Chat form pickers (commands, mentions) | +| `use-debounced-search` | Shared debounced async search for pickers | +| `use-picker-navigation` | Picker keyboard navigation | +| `use-chat-message-edit-context` | Message edit context (content + extras) | +| `use-chat-screen-drag-and-drop` | Drag-and-drop state machine | +| `use-chat-screen-file-upload` | File upload queue + capability validation | +| `use-chat-screen-scroll` | Scroll container binding + navigation guard | +| `use-auto-scroll` | Auto-scroll controller for streaming | +| `use-marquee-selection` | Shift+click / marquee range selection | +| `use-keyboard-shortcuts` | Global keyboard shortcuts | +| `use-settings-navigation` | Settings section navigation | +| `use-pwa` | PWA install/update + version mismatch detection | #### Stores (`src/lib/stores/`) Stores own reactive application state as Svelte 5 runes. Larger stores are split into directories and compose focused sub-stores behind a narrow host interface (see Architectural Patterns). -| Store | Responsibility | -| -------------------- | --------------------------------------------------------------------------------------------------------------- | -| `chatStore` | Chat lifecycle, streaming, abort control, error handling; composes `processing`, `activity`, `streams`, `flows` | -| `conversationsStore` | Conversation CRUD, message branching, navigation, import/export; composes `preferences` | -| `modelsStore` | Model list, selection, loading/unloading (ROUTER); composes `props`, `status` | -| `mcpStore` | MCP host role: multi-server lifecycle, tool routing; composes `health`, `resources` | -| `agenticStore` | Multi-turn agentic loop orchestration, tool execution; composes `gates` | -| `serverStore` | Server connection state, `/props`, role detection, modalities | -| `settingsStore` | User preferences, theme, parameter sync with server defaults | -| `toolsStore` | Tool registry: server + MCP tools, enabled set for the LLM | -| `permissionsStore` | Persisted tool permission grants | -| `contextStatsStore` | Context window usage for the active conversation | -| `draftMessagesStore` | Per-chat draft message/files | -| `deviceStore` | Browser environment signals (mobile, OS, theme) | -| `versionStore` | Build version information | +| Store | Responsibility | +| -------------------- | ----------------------------------------------------------------------------------------------------------------------------------------- | +| `chatStore` | Chat lifecycle, streaming, abort control, error handling; composes `processing`, `activity`, `streams`, `flows` | +| `conversationsStore` | Conversation CRUD, message branching, navigation, import/export; composes `preferences` | +| `modelsStore` | Model list, selection, loading/unloading (ROUTER), recent picks, favorites, hidden models and list open state; composes `props`, `status` | +| `mcpStore` | MCP host role: multi-server lifecycle, tool routing, server icons; composes `health`, `resources` | +| `agenticStore` | Multi-turn agentic loop orchestration, tool execution; composes `gates` | +| `serverStore` | Server connection state, `/props`, role detection, modalities | +| `settingsStore` | User preferences, theme, parameter sync with server defaults | +| `toolsStore` | Tool registry: server + MCP tools, enabled set for the LLM | +| `permissionsStore` | Persisted tool permission grants | +| `contextStatsStore` | Context window usage for the active conversation | +| `draftMessagesStore` | Per-chat draft message/files | +| `deviceStore` | Browser environment signals (mobile, OS, theme) | +| `versionStore` | Build version information | #### Services (`src/lib/services/`) Services are a stateless protocol layer: static methods, pure I/O, no reactive state. Stores consume them for all API and storage access. -| Service | Responsibility | -| ----------------------------- | ------------------------------------------------------------------------- | -| `ChatService` | `/v1/chat/completions` streaming + SSE parsing, message format conversion | -| `ModelsService` | `/models`, `/models/load`, `/models/unload` | -| `PropsService` | `/props`, `/props?model=` | -| `DatabaseService` | IndexedDB operations via Dexie | -| `MCPService` | MCP protocol: transports, connect, list/execute tools, prompts, resources | -| `ToolsService` | Server tool list/execute/stream (`/tools`) | -| `SandboxService` | Browser JS execution in a sandboxed worker | -| `ParameterSyncService` | Syncs settings with server defaults | -| `ConversationTransferService` | Conversation import/export JSONL + ZIP format | -| `MigrationService` | Non-destructive localStorage/IndexedDB migrations | -| `RouterService` | Dynamic route URL construction | +| Service | Responsibility | +| ----------------------------- | ----------------------------------------------------------------------------------------------------------------------------- | +| `ChatService` | `/v1/chat/completions` streaming + SSE parsing, message format conversion | +| `ModelsService` | `/models`, `/models/load`, `/models/unload` | +| `PropsService` | `/props`, `/props?model=` | +| `HuggingFaceService` | Hugging Face Hub metadata: model details, avatars, README and file tree, cached per repo and off unless the Hub setting is on | +| `DatabaseService` | IndexedDB operations via Dexie | +| `MCPService` | MCP protocol: transports, connect, list/execute tools, prompts, resources | +| `ToolsService` | Server tool list/execute/stream (`/tools`) | +| `SandboxService` | Browser JS execution in a sandboxed worker | +| `ParameterSyncService` | Syncs settings with server defaults | +| `ConversationTransferService` | Conversation import/export JSONL + ZIP format | +| `MigrationService` | Non-destructive localStorage/IndexedDB migrations | +| `RouterService` | Dynamic route URL construction | --- @@ -660,8 +673,17 @@ flowchart TB ``` - **IndexedDB**: Conversations and messages (large, structured data) -- **LocalStorage**: Settings, user parameter overrides, theme (small key-value data) -- **Memory only**: Server props, model list (fetched fresh on each session) +- **LocalStorage**: Settings, user parameter overrides, theme, model favorites, recent picks, hidden models and model list open state (small key-value data) +- **Memory only**: Server props, model list (fetched fresh on each session), Hugging Face Hub details (cached per repo for the session) + +### 9. Remote Assets Under Cross-Origin Isolation + +llama-server serves the UI with `Cross-Origin-Embedder-Policy: require-corp` and `Cross-Origin-Opener-Policy: same-origin`, so a cross-origin subresource has to be allowed by the target origin. The UI handles that per asset kind: + +- **Hub avatars** are requested in CORS mode (`crossorigin="anonymous"`), which the Hub allows +- **Hub metadata** is fetched with `fetch()`, which is a CORS request by default, and only while the Hub setting is on +- **MCP server icons and web-search result favicons** have no CORS guarantee, so they go through the server's `/cors-proxy` when the user started llama-server with `--ui-mcp-proxy`; without it they fall back to the bundled glyph +- **Per-row Hub lookups** wait until the row is near the viewport (the `nearViewport` action), so a long model list does not fire one request per row on mount --- diff --git a/tools/ui/package-lock.json b/tools/ui/package-lock.json index 86d599e671c5..1c0e4ec98abf 100644 --- a/tools/ui/package-lock.json +++ b/tools/ui/package-lock.json @@ -12,7 +12,7 @@ "@eslint/compat": "1.4.1", "@eslint/js": "9.39.2", "@internationalized/date": "3.12.2", - "@lucide/svelte": "1.25.0", + "@lucide/svelte": "1.46.0", "@modelcontextprotocol/sdk": "1.30.0", "@playwright/test": "1.56.1", "@storybook/addon-a11y": "10.5.6", @@ -79,6 +79,7 @@ "unified": "11.0.5", "unist-util-visit": "5.1.0", "uuid": "13.0.2", + "vaul-svelte": "1.0.0-next.7", "vite": "7.3.6", "vite-plugin-devtools-json": "0.2.1", "vitest": "4.1.10", @@ -3225,9 +3226,9 @@ } }, "node_modules/@lucide/svelte": { - "version": "1.25.0", - "resolved": "https://registry.npmjs.org/@lucide/svelte/-/svelte-1.25.0.tgz", - "integrity": "sha512-v9m+dD68jxVnqkU3K59mG/RSRFlPGzmKCGSyMfnXcaGv9jODDQMyQkcp1CGvk3Y/cUj9v7f8rw1n//K0B53xGQ==", + "version": "1.46.0", + "resolved": "https://registry.npmjs.org/@lucide/svelte/-/svelte-1.46.0.tgz", + "integrity": "sha512-Ru8jDjoCVLJ7J7pPRUdY+odSBV2BI6rj68JTmT+ENq/ZwYOPGPAiKVvKz+XLyAoK4HAwlRTzZHpsMvrHfsetqg==", "dev": true, "license": "ISC", "peerDependencies": { @@ -16998,6 +16999,61 @@ "node": ">= 0.8" } }, + "node_modules/vaul-svelte": { + "version": "1.0.0-next.7", + "resolved": "https://registry.npmjs.org/vaul-svelte/-/vaul-svelte-1.0.0-next.7.tgz", + "integrity": "sha512-7zN7Bi3dFQixvvbUJY9uGDe7Ws/dGZeBQR2pXdXmzQiakjrxBvWo0QrmsX3HK+VH+SZOltz378cmgmCS9f9rSg==", + "dev": true, + "license": "MIT", + "dependencies": { + "runed": "^0.23.2", + "svelte-toolbelt": "^0.7.1" + }, + "engines": { + "node": ">=18", + "pnpm": ">=8.7.0" + }, + "peerDependencies": { + "svelte": "^5.0.0" + } + }, + "node_modules/vaul-svelte/node_modules/runed": { + "version": "0.23.4", + "resolved": "https://registry.npmjs.org/runed/-/runed-0.23.4.tgz", + "integrity": "sha512-9q8oUiBYeXIDLWNK5DfCWlkL0EW3oGbk845VdKlPeia28l751VpfesaB/+7pI6rnbx1I6rqoZ2fZxptOJLxILA==", + "dev": true, + "funding": [ + "https://github.com/sponsors/huntabyte", + "https://github.com/sponsors/tglide" + ], + "dependencies": { + "esm-env": "^1.0.0" + }, + "peerDependencies": { + "svelte": "^5.7.0" + } + }, + "node_modules/vaul-svelte/node_modules/svelte-toolbelt": { + "version": "0.7.1", + "resolved": "https://registry.npmjs.org/svelte-toolbelt/-/svelte-toolbelt-0.7.1.tgz", + "integrity": "sha512-HcBOcR17Vx9bjaOceUvxkY3nGmbBmCBBbuWLLEWO6jtmWH8f/QoWmbyUfQZrpDINH39en1b8mptfPQT9VKQ1xQ==", + "dev": true, + "funding": [ + "https://github.com/sponsors/huntabyte" + ], + "dependencies": { + "clsx": "^2.1.1", + "runed": "^0.23.2", + "style-to-object": "^1.0.8" + }, + "engines": { + "node": ">=18", + "pnpm": ">=8.7.0" + }, + "peerDependencies": { + "svelte": "^5.0.0" + } + }, "node_modules/vfile": { "version": "6.0.3", "resolved": "https://registry.npmjs.org/vfile/-/vfile-6.0.3.tgz", diff --git a/tools/ui/package.json b/tools/ui/package.json index 40176047417f..7b15f3343ff4 100644 --- a/tools/ui/package.json +++ b/tools/ui/package.json @@ -32,7 +32,7 @@ "@eslint/compat": "1.4.1", "@eslint/js": "9.39.2", "@internationalized/date": "3.12.2", - "@lucide/svelte": "1.25.0", + "@lucide/svelte": "1.46.0", "@modelcontextprotocol/sdk": "1.30.0", "@playwright/test": "1.56.1", "@storybook/addon-a11y": "10.5.6", @@ -99,6 +99,7 @@ "unified": "11.0.5", "unist-util-visit": "5.1.0", "uuid": "13.0.2", + "vaul-svelte": "1.0.0-next.7", "vite": "7.3.6", "vite-plugin-devtools-json": "0.2.1", "vitest": "4.1.10", diff --git a/tools/ui/src/app.css b/tools/ui/src/app.css index f9b544bebc56..ccf13fbec3fc 100644 --- a/tools/ui/src/app.css +++ b/tools/ui/src/app.css @@ -6,6 +6,33 @@ @custom-variant dark (&:is(.dark *)); +/* shadcn components key their orientation styles off these, toggle-group among them */ +@custom-variant data-horizontal (&[data-orientation='horizontal']); +@custom-variant data-vertical (&[data-orientation='vertical']); + +/* + * Tailwind v4 dropped the pointer cursor from buttons. Every control a user can + * press gets it back, and anything disabled keeps the cursor its own class asks for. + */ +@layer base { + button:not(:disabled), + [role='button']:not([data-disabled]), + [role='checkbox']:not([data-disabled]), + [role='combobox']:not([data-disabled]), + [role='menuitem']:not([data-disabled]), + [role='menuitemcheckbox']:not([data-disabled]), + [role='menuitemradio']:not([data-disabled]), + [role='option']:not([data-disabled]), + [role='radio']:not([data-disabled]), + [role='switch']:not([data-disabled]), + [role='tab']:not([data-disabled]), + label[for], + select, + summary { + cursor: pointer; + } +} + :root { --radius: 0.625rem; --background: oklch(1 0 0); @@ -140,7 +167,7 @@ } body { - @apply bg-background text-foreground; + @apply bg-background text-foreground antialiased; scrollbar-width: thin; scrollbar-gutter: stable; } diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte index 893f8077dd6a..397ddc322b1a 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte @@ -8,8 +8,7 @@ ChatFormInputFileInputInvisible, ChatFormMcpResourcesList, ChatFormPickers, - DialogMcpResourcesBrowser, - DialogMcpServers + DialogMcpResourcesBrowser } from '$lib/components/app'; import { CLIPBOARD_CONTENT_QUOTE_PREFIX, @@ -34,7 +33,8 @@ modelsStore, serverStore, settingsStore, - toolsStore + toolsStore, + uiStore } from '$lib/stores'; import type { FileMentionEntry, @@ -193,7 +193,6 @@ let preSelectedResourceUri = $state(undefined); // MCP Servers Dialog State - let isMcpServersDialogOpen = $state(false); let currentConfig = $derived(settingsStore.config); @@ -478,6 +477,19 @@ queueMicrotask(() => inputRef?.focus()); } + // A flow outside the chat, e.g. the models manager closing onto it, asks for the + // composer back. The dialog's focus scope is still tearing down, so focus on a + // frame of its own. + $effect(() => { + if (uiStore.manageModelsOpen) return; + + if (!uiStore.composerFocusRequested) return; + + uiStore.composerFocusRequested = false; + + setTimeout(() => inputRef?.focus(), 0); + }); + // Splice the mention link in place of the `@` token. Uses the // live cursor, not a stale snapshot - the token may have been edited. function handleMentionSelect(entry: FileMentionEntry) { @@ -628,7 +640,6 @@ isReasoning={chatStore.isReasoning} {isRecording} onFileUpload={handleFileUpload} - onMcpSettingsClick={() => (isMcpServersDialogOpen = true)} onMicClick={handleMicClick} {onStop} onSystemPromptClick={() => onSystemPromptClick?.({ files: uploadedFiles, message: value })} @@ -667,5 +678,3 @@ }} preSelectedUri={preSelectedResourceUri} /> - - diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddButton.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddButton.svelte index 60e7bd1d4aa2..443046615ba7 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddButton.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddButton.svelte @@ -15,11 +15,11 @@ + + + + {#if toolsPanel.totalToolCount > 0} + (toolsExpanded = open)} open={toolsExpanded}> + + {#if toolsExpanded} + + {:else} + + {/if} + + + + {ATTACHMENT_MENU_TEXT.TOOLS} + + + +
+ {#each toolsPanel.categoryGroups as group (group.key)} + {@render groupRow(group)} + {/each} + + {#each toolsPanel.mcpGroups as group (group.key)} + {@render groupRow(group)} + {/each} +
+ + + {/if} + + + + + +{#snippet groupRow(group: ToolGroup)} + {@const checkState = toolsPanel.getGroupCheckState(group)} + {@const enabledCount = toolsPanel.getEnabledToolCount(group)} + {@const favicon = toolsPanel.getFavicon(group)} + {@const groupDisabled = toolsPanel.isGroupDisabled(group)} + + +{/snippet} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte index 1ce731845fe7..2cae6e8caeba 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte @@ -1,23 +1,20 @@ - - - - {#if reasoning.isReasoningActive} - - {:else if reasoning.isOff} - - {:else} - - {/if} - - - Reasoning - - - {reasoning.currentEffort} - - - - - - {#each reasoning.levels as level (level.value)} - {@const tokenLabel = reasoning.tokenLabel(level)} - reasoning.select(level)} - > - {#if reasoning.isSelected(level)} - - {:else} -
- {/if} - - {level.label} - - {#if tokenLabel} - - {tokenLabel} - - {/if} - - {#if level.hasInfo} - - - - - - -

Maximum reasoning effort with extended context usage

- - - {/if} - - {/each} - - diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte deleted file mode 100644 index acd0f4d21161..000000000000 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte +++ /dev/null @@ -1,284 +0,0 @@ - - -
- - {@render trigger({ disabled: chatFormActions.disabled, onclick: () => (sheetOpen = true) })} - - - - Add to chat - - - Add files, system prompt or configure MCP servers - - - -
- {#if reasoning.modelSupportsThinking} - (reasoningExpanded = open)} - open={reasoningExpanded} - > - - {#if reasoningExpanded} - - {:else} - - {/if} - - {#if reasoning.thinkingEnabled} - - {:else if reasoning.isOff} - - {:else} - - {/if} - - Reasoning - - - {reasoning.currentEffort} - - - - -
- {#each reasoning.levels as level (level.value)} - {@const tokenLabel = reasoning.tokenLabel(level)} - - {/each} -
- - - {/if} - - (filesExpanded = open)} open={filesExpanded}> - - {#if filesExpanded} - - {:else} - - {/if} - - - - Add files - - - -
- {#each ATTACHMENT_FILE_ITEMS as item (item.id)} - {@const enabled = attachmentMenu.isItemEnabled(item.enabledWhen)} - {#if enabled} - - {:else if item.disabledTooltip} - - - - - - -

{item.disabledTooltip}

- - - {/if} - {/each} -
- - - - - - {#if toolsPanel.totalToolCount > 0} - (toolsExpanded = open)} open={toolsExpanded}> - - {#if toolsExpanded} - - {:else} - - {/if} - - - - Tools - - - {toolsPanel.totalToolCount} tool{toolsPanel.totalToolCount !== 1 ? 's' : ''} - - - - -
- {#each toolsPanel.categoryGroups as group (group.key)} - {@render sheetGroupRow(group)} - {/each} - - {#each toolsPanel.mcpGroups as group (group.key)} - {@render sheetGroupRow(group)} - {/each} -
- - - {/if} - - -
- - -
- -{#snippet sheetGroupRow(group: ToolGroup)} - {@const checkState = toolsPanel.getGroupCheckState(group)} - {@const enabledCount = toolsPanel.getEnabledToolCount(group)} - {@const favicon = toolsPanel.getFavicon(group)} - {@const groupDisabled = toolsPanel.isGroupDisabled(group)} - - -{/snippet} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte index f49544171480..fa40339be2be 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte @@ -4,7 +4,7 @@ import * as Collapsible from '$lib/components/ui/collapsible'; import * as DropdownMenu from '$lib/components/ui/dropdown-menu'; import * as Tooltip from '$lib/components/ui/tooltip'; - import { CLI_FLAGS, ICON_CLASS_DEFAULT } from '$lib/constants'; + import { ATTACHMENT_MENU_TEXT, CLI_FLAGS, ICON_CLASS_DEFAULT } from '$lib/constants'; import { useToolsPanel } from '$lib/hooks/use-tools-panel.svelte'; import { mcpStore, toolsStore } from '$lib/stores'; import type { ToolGroup } from '$lib/types'; @@ -17,7 +17,7 @@ - Tools + {ATTACHMENT_MENU_TEXT.TOOLS} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionsAdd.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionsAdd.svelte index b2581f11eed2..50eb9be1c017 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionsAdd.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionsAdd.svelte @@ -1,16 +1,16 @@ {#if deviceStore.isMobile} - + {#snippet trigger({ disabled, onclick })} {/snippet} - + {:else} {/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte deleted file mode 100644 index a4baa40bf994..000000000000 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte +++ /dev/null @@ -1,170 +0,0 @@ - - -{#if deviceStore.isMobile} - -{:else} - -{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionReasoning.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionReasoning.svelte new file mode 100644 index 000000000000..29de5ec26366 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionReasoning.svelte @@ -0,0 +1,147 @@ + + +{#snippet triggerLabel()} + + {#if reasoning.isOff} + + {:else} + + {/if} + + {#if !isDefault} + {reasoning.currentEffort} + {/if} + +{/snippet} + +{#snippet levelRow(level: ReasoningEffortLevel)} + {@const tokenLabel = reasoning.tokenLabel(level)} + + +{/snippet} + + +{#if isMobile} + + + + + + Reasoning effort + + +
+ {#each reasoning.levels as level (level.value)} + {@render levelRow(level)} + {/each} +
+ + +{:else} + + + {#snippet child({ props })} + + {/snippet} + + + + {#each reasoning.levels as level (level.value)} + {@const tokenLabel = reasoning.tokenLabel(level)} + + reasoning.select(level)}> + {#if reasoning.isSelected(level)} + + {:else} +
+ {/if} + + {level.label} + + {#if tokenLabel} + {tokenLabel} + {/if} + + {#if level.hasInfo} + + + + + + +

{MAX_EFFORT_HINT}

+ + + {/if} + + {/each} + + +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionRecord.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionRecord.svelte index d1dd3fe46c8b..7c305acbee9a 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionRecord.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionRecord.svelte @@ -27,7 +27,7 @@ + {/snippet} + + {:else} + + {/if} +{/snippet} + +{#snippet region()} + + {@render children()} +{/snippet} + +{#if triggerPosition === 'top'} + {@render triggerButton()} + + {@render region()} +{:else} + {@render region()} + + {@render triggerButton()} +{/if} diff --git a/tools/ui/src/lib/components/app/misc/GroupedList.svelte b/tools/ui/src/lib/components/app/misc/GroupedList.svelte new file mode 100644 index 000000000000..e249918c507c --- /dev/null +++ b/tools/ui/src/lib/components/app/misc/GroupedList.svelte @@ -0,0 +1,188 @@ + + + + +{#if groups} + {#each windowed as entry (entry.key)} + {@const expanded = !collapsed.has(entry.key)} + + {#if group} +
+ {@render group({ + depth: 0, + expanded, + group: entry.group, + key: entry.key, + toggle: () => toggleGroup(entry.key) + })} +
+ {/if} + + + {#each entry.rows as row (keyOf(row))} + {@render item({ depth: 1, entry: row })} + {/each} + + {#if entry.hidden > 0 && more} + {@render more({ + count: Math.min(groupWindow, entry.hidden), + onMore: () => growGroup(entry.key), + unit: GroupedListUnit.ENTRIES + })} + {/if} + + {/each} + + {#if windowed.length < (groups?.length ?? 0) && more} + {@render more({ + count: groupPage.count, + onMore: () => (sectionSteps += Math.max(1, groupPage.weight)), + unit: GroupedListUnit.GROUPS + })} + {/if} +{:else} + {#each flatRows as row (keyOf(row))} + {@render item({ depth: 0, entry: row })} + {/each} + + {#if flatRows.length < items.length && more} + {@render more({ + count: Math.min(sectionWindow, items.length - flatRows.length), + onMore: () => (sectionSteps += Math.min(sectionWindow, items.length - flatRows.length)), + unit: GroupedListUnit.ENTRIES + })} + {/if} +{/if} diff --git a/tools/ui/src/lib/components/app/misc/index.ts b/tools/ui/src/lib/components/app/misc/index.ts index a10410ef9289..a4ee2ebc35c7 100644 --- a/tools/ui/src/lib/components/app/misc/index.ts +++ b/tools/ui/src/lib/components/app/misc/index.ts @@ -59,3 +59,12 @@ export { default as CodeBlockActions } from './CodeBlockActions.svelte'; * `class` and `style` props and inherits color via `currentColor`. */ export { default as Logo } from './Logo.svelte'; + +/** Section with a trigger that expands in place. */ +export { default as CollapsibleSection } from './CollapsibleSection.svelte'; +export { + default as GroupedList, + type GroupedListGroup, + GroupedListUnit +} from './GroupedList.svelte'; +export { default as CollapsibleRegion } from './CollapsibleRegion.svelte'; diff --git a/tools/ui/src/lib/components/app/models/ModelAvatar.svelte b/tools/ui/src/lib/components/app/models/ModelAvatar.svelte new file mode 100644 index 000000000000..5e1234a27267 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelAvatar.svelte @@ -0,0 +1,106 @@ + + +{#if orgName && hubEnabled} + (isNearViewport = true)} + class={['inline-flex shrink-0', className]} + > + + +{:else} + {@render fallback?.()} +{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelBadge.svelte b/tools/ui/src/lib/components/app/models/ModelBadge.svelte index d5b723ff9039..53cdcc26a8f8 100644 --- a/tools/ui/src/lib/components/app/models/ModelBadge.svelte +++ b/tools/ui/src/lib/components/app/models/ModelBadge.svelte @@ -1,8 +1,8 @@ + + (isNearViewport = true)} class={className}> + + diff --git a/tools/ui/src/lib/components/app/models/ModelCapabilityIcons.svelte b/tools/ui/src/lib/components/app/models/ModelCapabilityIcons.svelte index c9505a95eb5e..3743a158fdca 100644 --- a/tools/ui/src/lib/components/app/models/ModelCapabilityIcons.svelte +++ b/tools/ui/src/lib/components/app/models/ModelCapabilityIcons.svelte @@ -1,6 +1,14 @@ - - - {#if supportsToolUse && !hideCapabilities} - - - - + let capabilities = $derived([ + ...(supportsToolUse ? [ModelCapability.TOOL_USE] : []), + ...(supportsThinking ? [ModelCapability.REASONING] : []) + ]); - -

Tool use

- - - {/if} + let shownModalities = $derived( + MODALITY_ORDER.filter((modality) => modalities?.[MODALITY_FLAG_KEYS[modality]]) + ); - {#if supportsThinking && !hideCapabilities} - - - - + // an icon-less box still takes the gap where its icons would sit, so a caller with + // nothing to show renders no box at all + let hasIcons = $derived( + (!hideCapabilities && capabilities.length > 0) || + (shownModalities.length > 0 && !hideModalities) + ); + - -

Reasoning

- - - {/if} +{#if hasIcons} + + {#if !hideCapabilities} + {#each capabilities as capability (capability)} + {@const Icon = CAPABILITY_ICONS[capability]} - {#if hasModalityIcons && !hideModalities} - - {#if modalities?.vision} - + -

Vision

+

{CAPABILITY_LABELS[capability]}

 - {/if} - - {#if modalities?.video} - - - + {/each} + {/if} - -

Video

- - - {/if} + {#if shownModalities.length > 0 && !hideModalities} + + {#each shownModalities as modality (modality)} + {@const Icon = MODALITY_ICONS[modality]} - {#if modalities?.audio} - - - - + + + + - -

Audio

- - - {/if} - - {/if} - + +

{MODALITY_LABELS[modality]}

+ + + {/each} + + {/if} + +{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelContext.svelte b/tools/ui/src/lib/components/app/models/ModelContext.svelte new file mode 100644 index 000000000000..6d5395fc4497 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelContext.svelte @@ -0,0 +1,58 @@ + + + (isNearViewport = true)} + class={['text-sm text-muted-foreground', className]} +> + {formatContextLength(values)} + diff --git a/tools/ui/src/lib/components/app/models/discover/DownloadProgressBar.svelte b/tools/ui/src/lib/components/app/models/ModelDownloadProgressBar.svelte similarity index 80% rename from tools/ui/src/lib/components/app/models/discover/DownloadProgressBar.svelte rename to tools/ui/src/lib/components/app/models/ModelDownloadProgressBar.svelte index 6f4973c12ed1..c3a89c0ba45c 100644 --- a/tools/ui/src/lib/components/app/models/discover/DownloadProgressBar.svelte +++ b/tools/ui/src/lib/components/app/models/ModelDownloadProgressBar.svelte @@ -18,14 +18,14 @@ {#if overlay}
{:else}
diff --git a/tools/ui/src/lib/components/app/models/ModelDraftSidecars.svelte b/tools/ui/src/lib/components/app/models/ModelDraftSidecars.svelte new file mode 100644 index 000000000000..9b655f8432aa --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelDraftSidecars.svelte @@ -0,0 +1,42 @@ + + +{#each badges as badge (badge.kind + (badge.quant ?? ''))} + + + + {badge.kind} + +{/each} diff --git a/tools/ui/src/lib/components/app/models/ModelId.svelte b/tools/ui/src/lib/components/app/models/ModelId.svelte index ba3a976aef8b..26d72f662ff2 100644 --- a/tools/ui/src/lib/components/app/models/ModelId.svelte +++ b/tools/ui/src/lib/components/app/models/ModelId.svelte @@ -1,13 +1,15 @@ + +
+ {#if isLoading} + + {:else} + + {#if showAction && revealOnHover} + {#if isFailed} + + {:else} + + {/if} + {:else} + {#if isFailed} + + {:else} + + {/if} + {/if} + + +
+ {#if isFailed} + modelsStore.status.load(option.model)} + stopPropagationOnClick + tooltip="Retry loading model" + tooltipAsTitle + /> + {:else if isLoaded || isSleeping} + modelsStore.status.unload(option.model)} + stopPropagationOnClick + tooltip="Unload model" + tooltipAsTitle + /> + {:else} + modelsStore.status.load(option.model)} + stopPropagationOnClick + tooltip="Load model" + tooltipAsTitle + /> + {/if} +
+ {/if} +
diff --git a/tools/ui/src/lib/components/app/models/ModelOrgAvatar.svelte b/tools/ui/src/lib/components/app/models/ModelOrgAvatar.svelte new file mode 100644 index 000000000000..5befb9306d07 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelOrgAvatar.svelte @@ -0,0 +1,130 @@ + + +{#if hubEnabled} + + {#if orgAvatarFailed} + + {:else} +
+ + { + failedAvatarOrgs.set(org, Date.now()); + }} + src={HuggingFaceService.getAvatarUrl(org)} + /> +
+ {/if} + + {#if quantOrg && quantOrg !== org} + + + {#if quantAvatarFailed} + + {:else} + { + failedAvatarOrgs.set(quantOrg ?? '', Date.now()); + }} + src={HuggingFaceService.getAvatarUrl(quantOrg)} + /> + {/if} + + {/if} + +{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelRowActions.svelte b/tools/ui/src/lib/components/app/models/ModelRowActions.svelte new file mode 100644 index 000000000000..8f173174aac3 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelRowActions.svelte @@ -0,0 +1,75 @@ + + +
event.stopPropagation()} + onkeydown={(event) => event.stopPropagation()} + role="presentation" +> + + // a phone has no manager: its information dialog takes the icon + deviceStore.isMobile + ? uiStore.openModelInformation(option) + : uiStore.openModelsManager(option.id)} + tooltip="Manage model" + tooltipAsTitle + /> + + {#if isFav} + + + modelsStore.toggleFavorite(option.model)} + tooltip="Remove from favorites" + tooltipAsTitle + /> + + + + + {:else} + modelsStore.toggleFavorite(option.model)} + tooltip="Add to favorites" + tooltipAsTitle + /> + {/if} +
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManager.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManager.svelte new file mode 100644 index 000000000000..e6020e917202 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManager.svelte @@ -0,0 +1,357 @@ + + +
+
+ (selectedId = option.id)} + {selectedId} + {toolbarEnd} + /> +
+ + +
+ +
+ {#if shownOption} + {#key shownId} + (selectedId = null)} + onToggleLoad={() => void toggleLoad(shownOption)} + onUseInChat={() => void useInChat(shownOption)} + onUseInNewChat={() => void useInNewChat(shownOption)} + option={shownOption} + /> + {/key} + {/if} +
+
+
+ + diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerDownloadControl.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerDownloadControl.svelte new file mode 100644 index 000000000000..0e81500f2c1f --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerDownloadControl.svelte @@ -0,0 +1,50 @@ + + + +
+ + {percent !== null ? `${percent}%` : isPaused ? 'Paused' : ''} + + + +
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerFilters.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerFilters.svelte new file mode 100644 index 000000000000..8b4c3bf37e62 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerFilters.svelte @@ -0,0 +1,135 @@ + + + + + (contextLimit = Number(value))} + type="single" + value={String(contextLimit)} + > + + Context: + + {contextLabel} + + + + {#each CONTEXT_STEPS as step (step.value)} + {step.label} + {/each} + + + + + {#each toggles as toggle (toggle.value)} + + + + {/each} + + diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfiguration.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfiguration.svelte new file mode 100644 index 000000000000..8a109a8db41e --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfiguration.svelte @@ -0,0 +1,110 @@ + + +
+ + +
+ + +
+ +
+
+ + + {#if deviceStore.isMobile} +
+ +
+ {/if} + + (deleteOpen = false)} + open={deleteOpen} + repoWithTag={option.model} + /> +
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationActions.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationActions.svelte new file mode 100644 index 000000000000..b7fc10997504 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationActions.svelte @@ -0,0 +1,85 @@ + + +
+ {#if isMobile} + + + {/if} + +
+ {#if !isMobile} + + {/if} + + {#if canToggleLoad} + + {/if} +
+
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationHeader.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationHeader.svelte new file mode 100644 index 000000000000..40f5d0d5481b --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationHeader.svelte @@ -0,0 +1,111 @@ + + + +
+
+
+ + + +
+ + + +

+ + + + {statusLabel} + +

+
+
+ + +
+ + + {#if !isMobile} + + {/if} +
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationInformation.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationInformation.svelte new file mode 100644 index 000000000000..7eed08516af6 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfigurationInformation.svelte @@ -0,0 +1,197 @@ + + +
+
+ Model + + + + {option.model} + + + + +
+ + {#each visibleRows as row (row.label)} +
+ {row.label} + + + {#if row.value === null} + - + {:else} + + {#if row.isBadge} + {row.value} + {:else} + {row.value} + {/if} + + + {#if row.isCopyable} + + {/if} + {/if} + +
+ {/each} + + {#if visibleRows.length === 0} +

+ The model reports its metadata once it is loaded. +

+ {/if} + + {#if modalities.length > 0} +
+ Modalities + + +
+ {/if} +
+ +{#if !serverProps} +

+ /props values - file path, slots, build info - appear once the model is loaded. Reading this + page never loads a model. +

+{/if} + +
+ + {#snippet trigger()} + Chat Template + {/snippet} + + 
{serverProps?.chat_template ??
+				gguf?.chat_template ??
+				'Shown once the model is loaded.'}
+ +
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelRow.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelRow.svelte new file mode 100644 index 000000000000..f4d86b7293aa --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelRow.svelte @@ -0,0 +1,104 @@ + + +
+ + + + + + + + {#if download} + + {/if} + +
+ +
+
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelsTable.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelsTable.svelte new file mode 100644 index 000000000000..2a3c2206b294 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelsTable.svelte @@ -0,0 +1,422 @@ + + +{#snippet sortHeader(key: ModelsTableSortKey, label: string)} + +{/snippet} + +{#snippet entryTree(entry: ModelQuantGroup, indent = 0)} + {#if entry.quants.length > 1} + toggleQuants(entry.key)} + /> + + + {#each entry.quants as quant (quant.id)} + + {/each} + + {:else} + + {/if} +{/snippet} + +{#snippet familyRow({ + expanded, + group: family, + toggle +}: { + expanded: boolean; + group: ModelFamilyGroup; + toggle: () => void; +})} + {@const countLabel = `${family.entries.length} model${family.entries.length === 1 ? '' : 's'}`} + +
handleFamilyKeydown(event, toggle)} + role="button" + tabindex="0" + > + + + + {family.label} + + {countLabel} + + + + + + + + + {#if expanded} + + {:else} + + {/if} + +
+{/snippet} + +{#snippet listItem({ depth, entry }: { depth: number; entry: ModelQuantGroup })} + {@render entryTree(entry, depth > 0 ? 16 : 0)} +{/snippet} + +{#snippet showMore({ + count, + onMore, + unit +}: { + count: number; + onMore: () => void; + unit: GroupedListUnit; +})} +
+ +
+{/snippet} + + (deleteOpen = false)} + open={deleteOpen} + repoWithTag={pendingDelete} +/> + +
+ + +
+ {@render sortHeader(ModelsTableSortKey.NAME, 'Model')} + + + {@render sortHeader(ModelsTableSortKey.CONTEXT, 'Context')} + + + + {@render sortHeader(ModelsTableSortKey.STATUS, 'Status')} + + + Actions +
+ +
+ {#each sections as group (group.key)} + {#if group.items.length > 0} + {#snippet groupIcon()} + {#if group.kind === ModelsTableGroupKind.DOWNLOADING} + + {:else if group.kind === ModelsTableGroupKind.FAVORITES} + + {:else if group.kind === ModelsTableGroupKind.LOADED} + + {:else if group.kind === ModelsTableGroupKind.HIDDEN} + + {:else if group.kind === ModelsTableGroupKind.LOCAL} + + {/if} + {/snippet} + + + entry.key} + more={showMore} + onCollapsedChange={(key, collapsed) => + modelsStore.setGroupCollapsed(group.key, key, collapsed)} + sectionWindow={MODEL_ROW_WINDOW} + stickyClass="bg-muted/30 backdrop-blur-lg" + stickyStyle={MODELS_TABLE_GROUP_STICKY_OFFSET} + weightOf={(entry) => entry.quants.length} + /> + + {/if} + {/each} + + {#if isEmpty} +

+ {hasFilters ? 'No models match these filters.' : 'No models found.'} +

+ {/if} +
+
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerQuantRow.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerQuantRow.svelte new file mode 100644 index 000000000000..6d76251cdedd --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerQuantRow.svelte @@ -0,0 +1,71 @@ + + +
+ + + + + + + +
+ +
+
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerRepoRow.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerRepoRow.svelte new file mode 100644 index 000000000000..8c26b201fa00 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerRepoRow.svelte @@ -0,0 +1,95 @@ + + +
+ + + + + + + + + + + {groupLabel} + + + + + + + + + + + {#if expanded} + + {:else} + + {/if} + +
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerStatusCell.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerStatusCell.svelte new file mode 100644 index 000000000000..a56f50518e07 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerStatusCell.svelte @@ -0,0 +1,43 @@ + + +{#if download} + +{:else if isRouter} + +{:else} + + - + +{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerTableToolbar.svelte b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerTableToolbar.svelte new file mode 100644 index 000000000000..fef282fcc18b --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerTableToolbar.svelte @@ -0,0 +1,88 @@ + + + +
+ + + + + {#if hasFilters} + + {/if} + +
+ {@render toolbarEnd?.()} +
+
diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/row-actions.ts b/tools/ui/src/lib/components/app/models/ModelsManager/row-actions.ts new file mode 100644 index 000000000000..9931dd3f475f --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/row-actions.ts @@ -0,0 +1,61 @@ +import { Eye, EyeOff, Heart, HeartOff, Trash2 } from '@lucide/svelte'; +import { MODEL_DOWNLOAD_ICONS } from '$lib/constants'; +import { ModelRowDownloadState } from '$lib/enums'; +import { modelsStore } from '$lib/stores'; +import type { ModelOption } from '$lib/types/models'; + +/** Row actions follow the app's dropdown pattern: icon, label, separators, variants. */ +export function modelRowActions( + option: ModelOption, + favorite: boolean, + isHidden: boolean, + onDelete: (option: ModelOption) => void, + /** Download state, when the row stands for a tracked download. */ + download?: ModelRowDownloadState | null +) { + // a tracked download is paused, resumed or dropped: it is not loaded, hidden or + // drafted. Deleting it stops the download and removes what is on disk, which is + // what the same wording offers on a cached model + if (download) { + const isPaused = download === ModelRowDownloadState.PAUSED; + + return [ + { + icon: isPaused ? MODEL_DOWNLOAD_ICONS.resume : MODEL_DOWNLOAD_ICONS.pause, + label: isPaused ? 'Resume downloading' : 'Pause downloading', + onclick: () => + void (isPaused + ? modelsStore.status.downloadModel(option.model) + : modelsStore.status.pauseDownload(option.model)) + }, + { + icon: Trash2, + label: 'Delete from disk', + onclick: () => onDelete(option), + separator: true, + variant: 'destructive' as const + } + ]; + } + + return [ + { + icon: favorite ? HeartOff : Heart, + label: favorite ? 'Remove from favorites' : 'Add to favorites', + onclick: () => modelsStore.toggleFavorite(option.model) + }, + { + icon: Trash2, + label: 'Delete from disk', + onclick: () => onDelete(option), + separator: true, + variant: 'destructive' as const + }, + { + icon: isHidden ? Eye : EyeOff, + label: isHidden ? 'Unhide model' : 'Hide model', + onclick: () => modelsStore.toggleHidden(option.id), + separator: true + } + ]; +} diff --git a/tools/ui/src/lib/components/app/models/ModelsManager/utils.ts b/tools/ui/src/lib/components/app/models/ModelsManager/utils.ts new file mode 100644 index 000000000000..3721ceca1091 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsManager/utils.ts @@ -0,0 +1,202 @@ +import { LOCAL_BACKEND_ID, type ModalityKey } from '$lib/constants'; +import { + ModelCapability, + ModelGroupKind, + ModelsTableGroupKind, + ServerModelStatus +} from '$lib/enums'; +import { HuggingFaceService, ModelsService } from '$lib/services'; +import { modelsStore } from '$lib/stores'; +import type { ModelDownloadEntry, ModelDownloadProgress, ModelOption } from '$lib/types/models'; +import { detectThinkingSupport, detectToolUseSupport, repoOf } from '$lib/utils'; + +/** One repo of the table, with the rows it ships as. */ +export interface ModelQuantGroup { + base: ModelOption; + key: string; + kind: ModelGroupKind; + quants: ModelOption[]; +} + +/** True when any of the manager's filters narrows the table. */ +export function hasActiveFilters( + contextLimit: number, + modalities: ModalityKey[], + capabilities: ModelCapability[] +): boolean { + return contextLimit > 0 || modalities.length > 0 || capabilities.length > 0; +} + +/** + * Order a status sorts behind: loaded first, then a model being worked on + * (loading, sleeping), then the rest. + */ +export function statusRank(option: ModelOption): number { + const status = modelsStore.getModelStatus(option.model); + + if (modelsStore.isModelRunning(option.model)) return 2; + + if (status === ServerModelStatus.LOADING || status === ServerModelStatus.SLEEPING) return 1; + + return 0; +} + +/** Byte counts of a tracked download: live while it runs, frozen while paused. */ +export function downloadProgressFor(repoWithTag: string): ModelDownloadProgress | null { + return ( + modelsStore.status.getDownloadProgress(repoWithTag) ?? + modelsStore.status.getPausedDownloadProgress(repoWithTag) + ); +} + +/** + * One table row per tracked download, so the table lists them the way it lists + * any other model. A paused download is already in the router's listing, so its + * own option is reused; a fresh one stands for its tag alone. + */ +export function downloadGroups( + entries: ModelDownloadEntry[], + models: ModelOption[] +): ModelQuantGroup[] { + return entries.map((entry) => { + const option = models.find((candidate) => candidate.model === entry.repoWithTag) ?? { + backendId: LOCAL_BACKEND_ID, + capabilities: [], + id: entry.repoWithTag, + model: entry.repoWithTag, + name: entry.repoWithTag + }; + + return { + base: option, + key: `download::${entry.repoWithTag}`, + kind: ModelGroupKind.QUANTS, + quants: [option] + }; + }); +} + +/** One collapsible block of the manager's table. */ +export interface ModelsTableGroup { + /** One entry per repo, its quants hanging off it. */ + items: ModelQuantGroup[]; + key: string; + kind: ModelsTableGroupKind; + label: string; +} + +/** Context the model runs with: what a loaded model reports. */ +export function configuredContext(option: ModelOption): number | null { + return modelsStore.isModelRunning(option.model) + ? modelsStore.props.getModelContextSize(option.model) + : null; +} + +/** + * Capability a model reports: true or false once its listing or its chat template + * answers, null while the Hub record that carries the template is not read yet. + */ +export function modelCapability(option: ModelOption, capability: ModelCapability): boolean | null { + if (option.capabilities.includes(capability)) return true; + + const details = HuggingFaceService.cachedDetails(option.model); + + // only the chat template answers this, and it arrives with the Hub record + if (details === undefined) return null; + + const template = details?.gguf?.chat_template ?? ''; + + return capability === ModelCapability.TOOL_USE + ? detectToolUseSupport(template) + : detectThinkingSupport(template); +} + +/** + * Context a model reports: the provider listing first, then the cached Hub record. + * Null while neither of them has answered. + */ +export function modelContextLength(option: ModelOption): number | null { + return ( + option.contextLength ?? + HuggingFaceService.cachedDetails(option.model)?.gguf?.context_length ?? + null + ); +} + +/** Fold the quants of one repo into a single entry, so the table shows one row per model. */ +export function groupModelQuants(models: ModelOption[]): ModelQuantGroup[] { + const groups = new Map(); + + for (const option of models) { + const key = repoOf(option.model); + const group = groups.get(key); + + if (group) { + group.quants.push(option); + + continue; + } + + groups.set(key, { base: option, key, kind: ModelGroupKind.QUANTS, quants: [option] }); + } + + return Array.from(groups.values()).flatMap((group) => { + const kind = groupKind(group.quants); + const modelIds = group.quants.map((option) => option.model); + + // the very same id twice is not a quant set; keep those rows apart + if (modelIds.length > 1 && modelIds.every((model) => model === modelIds[0])) { + return group.quants.map((option) => ({ + ...group, + base: option, + key: `${group.key}::${option.id}`, + kind: ModelGroupKind.VARIANTS, + quants: [option] + })); + } + + return [{ ...group, kind }]; + }); +} + +/** What a group folds: quants or variants of one repo. */ +function groupKind(quants: ModelOption[]): ModelGroupKind { + const isQuant = quants.every( + (option) => (option.parsedId ?? ModelsService.parseModelId(option.model)).quantization + ); + + return isQuant ? ModelGroupKind.QUANTS : ModelGroupKind.VARIANTS; +} + +/** + * Split the repos of a section into the quants that are hidden and the ones that + * are not, so a partly hidden repo lists in both blocks instead of dragging its + * visible quants into the hidden one. + */ +export function splitHiddenQuants( + entries: ModelQuantGroup[], + isHidden: (option: ModelOption) => boolean +): { hidden: ModelQuantGroup[]; local: ModelQuantGroup[] } { + const hidden: ModelQuantGroup[] = []; + const local: ModelQuantGroup[] = []; + + for (const entry of entries) { + const hiddenQuants = entry.quants.filter((quant) => isHidden(quant)); + const localQuants = entry.quants.filter((quant) => !isHidden(quant)); + + if (hiddenQuants.length > 0) { + hidden.push({ + ...entry, + base: hiddenQuants[0], + key: `${entry.key}::hidden`, + quants: hiddenQuants + }); + } + + if (localQuants.length > 0) { + local.push({ ...entry, base: localQuants[0], quants: localQuants }); + } + } + + return { hidden, local }; +} diff --git a/tools/ui/src/lib/components/app/models/ModelsSection.svelte b/tools/ui/src/lib/components/app/models/ModelsSection.svelte new file mode 100644 index 000000000000..749dc5b0705e --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSection.svelte @@ -0,0 +1,97 @@ + + + modelsStore.setGroupOpen(persistKey, open) : null} + {revealChevronOnHover} + {triggerClass} + {triggerStyle} +> + {#snippet trigger()} + {#if onBack} + + {/if} + + {#if icon} + {@render icon()} + {/if} + + {label} + + {#if loading} + + {:else if error} + + {/if} + + {#if count !== undefined} + {count} + {/if} + {/snippet} + + {@render children()} + diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelector.svelte b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelector.svelte new file mode 100644 index 000000000000..33b3b2f86e6c --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelector.svelte @@ -0,0 +1,405 @@ + + +{#snippet selectorTriggerInner()} + + + + {#if selectedOption} + + {:else} + Select model + {/if} + + + {#if ms.updating || ms.isLoadingModel || triggerLoading} + + {:else} + + {/if} + + {#if triggerLoading} + + {/if} +{/snippet} + +{#snippet modelOption(item: ModelItem, hideOrgName: boolean)} + {@const { option } = item} + {@const isSelected = ms.activeId === option.id} + {@const isHighlighted = option.id === highlightedId} + {@const isFav = ms.isFavorite(option.model)} + + { + if (event.key === KeyboardKey.ENTER || event.key === KeyboardKey.SPACE) { + event.preventDefault(); + void handleModelKeyAction(option.id, event.altKey); + } + }} + onMouseEnter={() => (highlightedId = option.id)} + onSelect={ms.handleSelect} + {option} + /> +{/snippet} + +{#snippet pickerList(listClass: string)} +
+ {#if !ms.isCurrentModelInCache && displayModel} + + + {/if} + + {#if ms.isEmpty} +

{ms.emptyMessage}

+ {/if} + + +
+{/snippet} + +{#if ms.loading && ms.options.length === 0 && ms.isMultiModel} +
+ + + Loading models... +
+{:else if ms.options.length === 0 && ms.isMultiModel} + {#if displayModel} + + + + {:else} + No models yet. + {/if} +{:else if ms.isMultiModel} + {#if isMobile} + + + ms.setSearchTerm(v)} + onSearchKeyDown={handleSearchKeyDown} + searchTerm={ms.searchTerm} + > + {@render pickerList('px-2.5')} + + {:else} + (highlightedId = id)} + onModelKeyAction={(id, unload) => void handleModelKeyAction(id, unload)} + onSearchKeyDown={handleSearchKeyDown} + {showOrgName} + /> + {/if} +{:else if isMobile} + + + + ms.setSearchTerm(v)} + onSearchKeyDown={handleSearchKeyDown} + searchTerm={ms.searchTerm} + > + {@render pickerList('px-2.5')} + +{:else} + + + + + {#snippet child({ props })} + + {/snippet} + + +{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDownloadItem.svelte b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDownloadItem.svelte new file mode 100644 index 000000000000..6ef1433b970f --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDownloadItem.svelte @@ -0,0 +1,130 @@ + + + +
+ + + + onRequestCancel?.(entry.repoWithTag)} + /> + + {#if entry.progress && entry.progress.totalBytes > 0} + + {/if} +
diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDrawer.svelte b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDrawer.svelte new file mode 100644 index 000000000000..1a2e8d3b82a7 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDrawer.svelte @@ -0,0 +1,53 @@ + + + + + +
+ + {@render children()} + +
+ + diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDropdown.svelte b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDropdown.svelte new file mode 100644 index 000000000000..a7a54e8602a1 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorDropdown.svelte @@ -0,0 +1,248 @@ + + +{#snippet selectorTriggerInner()} + + + + {#if selectedOption} + + {:else} + Select model + {/if} + + + {#if ms.updating || ms.isLoadingModel || triggerLoading} + + {:else} + + {/if} + + {#if triggerLoading} + + {/if} +{/snippet} + +{#snippet modelOption(item: ModelItem, hideOrgName: boolean)} + {@const { option } = item} + {@const isSelected = currentModel === option.model || ms.activeId === option.id} + {@const isHighlighted = option.id === highlightedId} + {@const isFav = ms.isFavorite(option.model)} + + { + if (event.key === KeyboardKey.ENTER || event.key === KeyboardKey.SPACE) { + event.preventDefault(); + onModelKeyAction?.(option.id, event.altKey); + } + }} + onMouseEnter={() => onHighlight?.(option.id)} + onSelect={ms.handleSelect} + {option} + /> +{/snippet} + + + + + + {#snippet child({ props })} + + {@render selectorTriggerInner()} + + {/snippet} + + + {#if selectedOption} + +

{triggerTooltipLabel(selectedOption)}

+ + {/if} + + + event.preventDefault()} + > + ms.setSearchTerm(v)} + {onSearchKeyDown} + placeholder="Search models..." + searchClass="bg-transparent" + searchValue={ms.searchTerm} + > + + {@render pickerContent('px-1.5')} + + {#snippet footer()} + + + + + Manage models + + + {/snippet} + + + + +{#snippet pickerContent(listClass: string)} +
+ {#if !ms.isCurrentModelInCache && currentModel} + + + {/if} + + {#if ms.isEmpty} +

{ms.emptyMessage}

+ {/if} + + +
+{/snippet} diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorList.svelte b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorList.svelte new file mode 100644 index 000000000000..6857e3018e64 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorList.svelte @@ -0,0 +1,235 @@ + + +{#snippet familyHeading({ group: family }: { group: ModelFamilyGroup })} +
+ + + + {family.label} + + + + {family.entries.length} model{family.entries.length === 1 ? '' : 's'} + +
+{/snippet} + +{#snippet listItem({ depth, entry }: { depth: number; entry: ModelItem })} +
{@render render(entry, !showOrgName)}
+{/snippet} + +{#snippet listRows(items: ModelItem[], prefix: string)} + {#if settingsStore.config[SETTINGS_KEYS.GROUP_MODELS_BY_FAMILY]} + row.option.model).map((family) => ({ + entries: family.entries, + group: family, + key: `${prefix}-${family.key}` + }))} + initialCollapsed={modelsStore.collapsedGroupsUnder(prefix)} + item={listItem} + keyOf={(row) => `${prefix}-${row.option.id}`} + onCollapsedChange={(key, collapsed) => modelsStore.setGroupCollapsed(prefix, key, collapsed)} + stickyStyle={SELECTOR_GROUP_STICKY_OFFSET} + /> + {:else} + `${prefix}-${row.option.id}`} /> + {/if} +{/snippet} + +{#snippet defaultOption(item: ModelItem, hideOrgName: boolean)} + {@const { option } = item} + {@const isSelected = currentModel === option.model || activeId === option.id} + {@const isFav = modelsStore.favoriteModelIds.has(option.model)} + + {}} + onMouseEnter={() => {}} + {onSelect} + {option} + /> +{/snippet} + +{#if loaded.length > 0} + + {#snippet icon()} + + {/snippet} + + {#each loaded as item (item.option.id)} + {@render render(item, !showOrgName)} + {/each} + +{/if} + +{#if favorites.length > 0} + + + {#snippet icon()} + + {/snippet} + + {#each favorites as item (`fav-${item.option.id}`)} + {@render render(item, !showOrgName)} + {/each} + +{/if} + +{#if downloadEntries.length > 0} +

Downloading

+ + {#each downloadEntries as entry (entry.repoWithTag)} + + {/each} +{/if} + +{#snippet localRows()} + {@render listRows( + localGroups.available.flatMap((group) => group.items), + 'local' + )} + + {#if hasMoreLocal} + + {/if} +{/snippet} + +{#if groups.available.length > 0} + + {#snippet icon()} + + {/snippet} + + {@render localRows()} + +{/if} + + (cancelOpen = false)} + open={cancelOpen} + repoWithTag={pendingCancel} +/> diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorOption.svelte b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorOption.svelte new file mode 100644 index 000000000000..2734573ca291 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorOption.svelte @@ -0,0 +1,105 @@ + + +
onSelect(option.id)} + onkeydown={onKeyDown} + onmouseenter={onMouseEnter} + role="option" + tabindex="0" + title={loadTitle} +> + + + + + {#if showCapabilities} + + {/if} + +
+ + + +
+ + {#if isLoading} + + {/if} +
diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorTriggerIcon.svelte b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorTriggerIcon.svelte new file mode 100644 index 000000000000..6c9621cddcb2 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/ModelsSelectorTriggerIcon.svelte @@ -0,0 +1,23 @@ + + +{#if option} + + {#snippet fallback()} + + {/snippet} + +{:else} + +{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelsSelector/index.ts b/tools/ui/src/lib/components/app/models/ModelsSelector/index.ts new file mode 100644 index 000000000000..e234fbac33e0 --- /dev/null +++ b/tools/ui/src/lib/components/app/models/ModelsSelector/index.ts @@ -0,0 +1,10 @@ +/** Model selection UI: a desktop serves an anchored dropdown menu and a phone a + * bottom drawer, over their shared parts. */ + +export { default as ModelsSelector } from './ModelsSelector.svelte'; +export { default as ModelsSelectorDropdown } from './ModelsSelectorDropdown.svelte'; +export { default as ModelsSelectorDrawer } from './ModelsSelectorDrawer.svelte'; +export { default as ModelsSelectorList } from './ModelsSelectorList.svelte'; +export { default as ModelsSelectorDownloadItem } from './ModelsSelectorDownloadItem.svelte'; +export { default as ModelsSelectorOption } from './ModelsSelectorOption.svelte'; +export { default as ModelsSelectorTriggerIcon } from './ModelsSelectorTriggerIcon.svelte'; diff --git a/tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte b/tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte deleted file mode 100644 index 30dedcbae90a..000000000000 --- a/tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte +++ /dev/null @@ -1,417 +0,0 @@ - - -
- {#if ms.loading && ms.options.length === 0 && ms.isRouter} -
- - - Loading models… -
- {:else if ms.options.length === 0 && ms.isRouter} - {#if currentModel} - - - - {:else} -

No models available.

- {/if} - {:else} - {@const selectedOption = ms.getDisplayOption()} - {@const triggerModel = selectedOption?.model} - {@const triggerStatus = triggerModel - ? modelsStore.routerModels.find((m) => m.id === triggerModel)?.status?.value - : undefined} - {@const triggerLoading = - !!triggerModel && - (triggerStatus === ServerModelStatus.LOADING || - modelsStore.status.isOperationInProgress(triggerModel))} - {@const triggerLoadPercent = triggerLoading - ? Math.round(modelLoadFraction(modelsStore.status.getLoadProgress(triggerModel)) * 100) - : 0} - - {#if ms.isRouter} - - - - - {#snippet child({ props })} - - - - - {#if selectedOption} - - {:else} - Select model - {/if} - - {#if reasoning.isReasoningActive} - - {/if} - - - {#if ms.updating || ms.isLoadingModel} - - {:else} - - {/if} - - {#if triggerLoading} - - {/if} - - {/snippet} - - - {#if selectedOption} - -

{selectedOption.model}

- - {/if} - - - event.preventDefault()} - > - - - - - {#if selectedOption} - - {:else} - No model - {/if} - - - - ms.setSearchTerm(v)} - onSearchKeyDown={handleSearchKeyDown} - placeholder="Search models..." - searchValue={ms.searchTerm} - > -
- {#if !ms.isCurrentModelInCache && currentModel} - - - {/if} - - {#if ms.filteredOptions.length === 0} -

No models found.

- {/if} - - {#snippet modelOption(item: ModelItem, hideOrgName: boolean)} - {@const { option } = item} - {@const isSelected = currentModel === option.model || ms.activeId === option.id} - {@const isHighlighted = option.id === highlightedId} - {@const isFav = ms.isFavorite(option.model)} - - { - if (event.key === KeyboardKey.ENTER || event.key === KeyboardKey.SPACE) { - event.preventDefault(); - void handleModelKeyAction(option.id, event.altKey); - } - }} - onMouseEnter={() => (highlightedId = option.id)} - onSelect={ms.handleSelect} - {option} - /> - {/snippet} - - -
- - - - - - - - {:else} - - - - {#snippet child({ props })} - - {/snippet} - - - {#if selectedOption} - -

{selectedOption.model}

- - {/if} - - {/if} - {/if} -
- -{#if ms.showModelDialog} - ms.setShowModelDialog(v)} - open={ms.showModelDialog} - /> -{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelsSelectorList.svelte b/tools/ui/src/lib/components/app/models/ModelsSelectorList.svelte deleted file mode 100644 index e40e33d0c087..000000000000 --- a/tools/ui/src/lib/components/app/models/ModelsSelectorList.svelte +++ /dev/null @@ -1,76 +0,0 @@ - - -{#snippet defaultOption(item: ModelItem, hideOrgName: boolean)} - {@const { option } = item} - {@const isSelected = currentModel === option.model || activeId === option.id} - {@const isFav = modelsStore.favoriteModelIds.has(option.model)} - - {}} - onMouseEnter={() => {}} - {onSelect} - {option} - /> -{/snippet} - -{#if groups.loaded.length > 0} -

Loaded models

- - {#each groups.loaded as item (`loaded-${item.option.id}`)} - {@render render(item, false)} - {/each} -{/if} - -{#if groups.favorites.length > 0} -

Favorite models

- - {#each groups.favorites as item (`fav-${item.option.id}`)} - {@render render(item, true)} - {/each} -{/if} - -{#if groups.available.length > 0} -

Available models

- - {#each groups.available as group (group.orgName)} - {#if group.orgName} -

{group.orgName}

- {/if} - - {#each group.items as item (item.option.id)} - {@render render(item, true)} - {/each} - {/each} -{/if} diff --git a/tools/ui/src/lib/components/app/models/ModelsSelectorOption.svelte b/tools/ui/src/lib/components/app/models/ModelsSelectorOption.svelte deleted file mode 100644 index cee7d0c3b57c..000000000000 --- a/tools/ui/src/lib/components/app/models/ModelsSelectorOption.svelte +++ /dev/null @@ -1,216 +0,0 @@ - - -
onSelect(option.id)} - onkeydown={onKeyDown} - onmouseenter={onMouseEnter} - role="option" - tabindex="0" - title={loadTitle} -> - - -
- - -
e.stopPropagation()} - > - {#if isFav} - modelsStore.toggleFavorite(option.model)} - tooltip="Remove from favorites" - /> - {:else} - modelsStore.toggleFavorite(option.model)} - tooltip="Add to favorites" - /> - {/if} - - - {#if isLoaded && onInfoClick} - onInfoClick(option.model)} - tooltip="Model information" - /> - {/if} -
- - {#if isLoading} -
- -
- {:else if isFailed} -
- - - -
- {:else if isSleeping} -
- - - -
- {:else if isLoaded} -
- - - -
- {:else} -
- - - -
- {/if} -
- - {#if isLoading} - - {/if} -
diff --git a/tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte b/tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte deleted file mode 100644 index c89ca186e3bb..000000000000 --- a/tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte +++ /dev/null @@ -1,210 +0,0 @@ - - -
- {#if ms.loading && ms.options.length === 0 && ms.isRouter} -
- - Loading models… -
- {:else if ms.options.length === 0 && ms.isRouter} -

No models available.

- {:else} - {@const selectedOption = ms.getDisplayOption()} - {@const triggerModel = selectedOption?.model} - {@const triggerStatus = triggerModel - ? modelsStore.routerModels.find((m) => m.id === triggerModel)?.status?.value - : undefined} - {@const triggerLoading = - !!triggerModel && - (triggerStatus === ServerModelStatus.LOADING || - modelsStore.status.isOperationInProgress(triggerModel))} - {@const triggerLoadPercent = triggerLoading - ? Math.round(modelLoadFraction(modelsStore.status.getLoadProgress(triggerModel)) * 100) - : 0} - - {#if ms.isRouter} - - - - - - Select Model - - - Choose a model to use for the conversation - - - -
-
- ms.setSearchTerm(v)} - placeholder="Search models..." - value={ms.searchTerm} - /> -
- -
- {#if !ms.isCurrentModelInCache && currentModel} - - -
- {/if} - - {#if ms.filteredOptions.length === 0} -

No models found.

- {/if} - - -
-
- - - {:else} - - {/if} - {/if} -
- -{#if ms.showModelDialog} - ms.setShowModelDialog(v)} - open={ms.showModelDialog} - /> -{/if} diff --git a/tools/ui/src/lib/components/app/models/discover/ModelsDiscoverAvatar.svelte b/tools/ui/src/lib/components/app/models/discover/ModelsDiscoverAvatar.svelte deleted file mode 100644 index f78e7458644b..000000000000 --- a/tools/ui/src/lib/components/app/models/discover/ModelsDiscoverAvatar.svelte +++ /dev/null @@ -1,109 +0,0 @@ - - - - {#if orgAvatarFailed} - - {:else} -
- { - failedAvatarOrgs.add(org); - }} - src={HuggingFaceService.getAvatarUrl(org)} - /> -
- {/if} - - {#if quantOrg && quantOrg !== org} - - - {#if quantAvatarFailed} - - {:else} - { - failedAvatarOrgs.add(quantOrg ?? ''); - }} - src={HuggingFaceService.getAvatarUrl(quantOrg)} - /> - {/if} - - {/if} - diff --git a/tools/ui/src/lib/components/app/models/index.ts b/tools/ui/src/lib/components/app/models/index.ts index f6a9f1f9ba6d..dbfb5c60e248 100644 --- a/tools/ui/src/lib/components/app/models/index.ts +++ b/tools/ui/src/lib/components/app/models/index.ts @@ -10,70 +10,6 @@ * */ -/** - * **ModelsSelectorDropdown** - Model selection dropdown (desktop) - * - * Dropdown for selecting AI models with status indicators, - * search, and model information display. Adapts UI based on server mode. - * - * **Architecture:** - * - Uses DropdownMenuSearchable for model list - * - Integrates with modelsStore for model options and selection - * - Detects router vs single mode from serverStore - * - Opens DialogModelInformation for model details - * - * **Features:** - * - Searchable model list with keyboard navigation - * - Model status indicators (loading/ready/error/updating) - * - Model capabilities badges (vision, tools, etc.) - * - Current/active model highlighting - * - Model information dialog on info button click - * - Router mode: shows all available models with status - * - Single mode: shows current model name only - * - Loading/updating skeleton states - * - Global selection support for form integration - * - * @example - * ```svelte - * updateModel(id)} - * useGlobalSelection - * /> - * ``` - */ -export { default as ModelsSelectorDropdown } from './ModelsSelectorDropdown.svelte'; - -/** - * **ModelsSelectorList** - Grouped model options list - * - * Renders grouped model options (loaded, favorites, available) with section - * headers and org subgroups. Shared between ModelsSelectorDropdown and ModelsSelectorSheet - * to avoid template duplication. - * - * Accepts an optional `renderOption` snippet to customize how each option is - * rendered (e.g., to add keyboard navigation or highlighting). - */ -export { default as ModelsSelectorList } from './ModelsSelectorList.svelte'; - -/** - * **ModelsSelectorOption** - Single model option row - * - * Renders a single model option with selection state, favorite toggle, - * load/unload actions, status indicators, and an info button. - * Used inside ModelsSelectorList or directly in custom render snippets. - */ -export { default as ModelsSelectorOption } from './ModelsSelectorOption.svelte'; - -/** - * **ModelsSelectorSheet** - Mobile model selection sheet - * - * Bottom sheet variant of ModelsSelectorDropdown optimized for touch interaction - * on mobile devices. Same functionality as ModelsSelectorDropdown but uses Sheet UI - * instead of DropdownMenu. - */ -export { default as ModelsSelectorSheet } from './ModelsSelectorSheet.svelte'; - /** * **ModelBadge** - Model name display badge * * Compact badge showing current model name with package icon. @@ -108,13 +44,19 @@ export { default as ModelBadge } from './ModelBadge.svelte'; * aliases, and tags. Supports raw mode to show the unprocessed model name. * Respects the user's `showRawModelNames` setting. */ +export { default as ModelCapabilities } from './ModelCapabilities.svelte'; +export { default as ModelContext } from './ModelContext.svelte'; export { default as ModelId } from './ModelId.svelte'; +export { default as ModelAvatar } from './ModelAvatar.svelte'; +export { default as ModelDownloadProgressBar } from './ModelDownloadProgressBar.svelte'; +export { default as ModelLoadControl } from './ModelLoadControl.svelte'; +export { default as ModelOrgAvatar } from './ModelOrgAvatar.svelte'; +export { default as ModelRowActions } from './ModelRowActions.svelte'; +export { default as ModelsSection } from './ModelsSection.svelte'; -/** - * **ModelCapabilityIcons** - Capability and modality icon row - * - * The shared tool-use / reasoning / vision / video / audio icon cluster with - * tooltips, used by ModelId and the discover details header so the order and - * styling stay consistent across every model-id surface. - */ +/** Capability and modality icon row, shared by every model-id surface. */ export { default as ModelCapabilityIcons } from './ModelCapabilityIcons.svelte'; + +/** Draft sidecar badges, rendered as `+ [KIND] [QUANT]` per sidecar. */ +export { default as ModelDraftSidecars } from './ModelDraftSidecars.svelte'; +export * from './ModelsSelector'; diff --git a/tools/ui/src/lib/components/app/navigation/DropdownMenuActions.svelte b/tools/ui/src/lib/components/app/navigation/DropdownMenuActions.svelte index 20df04eb423d..161462a04be1 100644 --- a/tools/ui/src/lib/components/app/navigation/DropdownMenuActions.svelte +++ b/tools/ui/src/lib/components/app/navigation/DropdownMenuActions.svelte @@ -1,7 +1,10 @@ - - - - - {#snippet child({ props })} - e.stopPropagation()} - > - {@render iconComponent(triggerIcon, 'h-3 w-3')} + // a phone shows the same actions in a drawer, where a thumb reaches them and every + // row is a touch target of its own + let isMobile = $derived(deviceStore.isMobile); - {#if triggerTooltip} - {triggerTooltip} - {/if} - - {/snippet} - + function runAction(action: ActionItem, event: Event) { + open = false; + + // let the drawer close before an action that moves the focus, such as a confirm dialog + setTimeout(() => action.onclick(event), 0); + } + + +{#if isMobile} + + + + + + {triggerTooltip ?? 'Actions'} + + +
+ {#each actions as action, index (action.label)} + {#if action.separator && index > 0} + + {/if} + + + {/each} +
+ + +{:else} + + + + + {#snippet child({ props })} + e.stopPropagation()} + > + {@render iconComponent(triggerIcon, 'h-3 w-3')} + + {#if triggerTooltip} + {triggerTooltip} + {/if} + + {/snippet} + - - {#each actions as action, index (action.label)} - {#if action.separator && index > 0} - + {#if triggerTooltip} + +

{triggerTooltip}

+ {/if} + - -
- {@render iconComponent( - action.icon, - `h-4 w-4 ${action.variant === 'destructive' ? 'text-destructive' : ''}` - )} - {action.label} -
- - {#if action.shortcut} - + + {#each actions as action, index (action.label)} + {#if action.separator && index > 0} + {/if} - - {/each} - - + + +
+ {@render iconComponent( + action.icon, + `h-4 w-4 ${action.variant === 'destructive' ? 'text-destructive' : ''}` + )} + {action.label} +
+ + {#if action.shortcut} + + {/if} + + {/each} + + +{/if} {#snippet iconComponent(IconComponent: Component, className: string)} diff --git a/tools/ui/src/lib/components/app/navigation/DropdownMenuSearchable.svelte b/tools/ui/src/lib/components/app/navigation/DropdownMenuSearchable.svelte index f14db5f23c0f..79336cf29d96 100644 --- a/tools/ui/src/lib/components/app/navigation/DropdownMenuSearchable.svelte +++ b/tools/ui/src/lib/components/app/navigation/DropdownMenuSearchable.svelte @@ -1,50 +1,79 @@ -
+
-
- {@render children()} +
+ +
+ {@render children()} - {#if isEmpty} -
{emptyMessage}
- {/if} + {#if isEmpty} +
{emptyMessage}
+ {/if} +
{#if footer} - +
+ - {@render footer()} + {@render footer()} +
{/if} diff --git a/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte b/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte index 8bf12279e3c7..2ce4824cf930 100644 --- a/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte +++ b/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte @@ -5,6 +5,9 @@ import { ActionIcon, DialogConversationRename, + DialogManageModels, + DialogMcpServers, + DialogModelInformation, DialogSettingsChat, Logo, SidebarNavigationActions, @@ -93,6 +96,7 @@ let renameDialogOpen = $state(false); let settingsDialogOpen = $state(false); + let mcpDialogOpen = $state(false); let renameTargetConversationId = $state(null); let renameDraft = $state(''); let renameOriginalTitle = $state(''); @@ -306,6 +310,15 @@ pendingCollapse = null; } } + + /** A dialog owns the screen on a phone, so the expanded list steps aside for it. */ + function openDialog(open: () => void): void { + if (deviceStore.isMobile) { + scheduleMobileCollapse(); + } + + open(); + } @@ -365,9 +378,9 @@ > 768 ? uiStore.isSidebarExpanded : true} + onManageModelsClick={() => openDialog(() => uiStore.openModelsManager())} + onMcpClick={() => openDialog(() => (mcpDialogOpen = true))} onNewChat={() => { if (deviceStore.isMobile) { scheduleMobileCollapse(); @@ -402,7 +417,7 @@ isSearchModeActive = false; searchQuery = ''; }} - onSettingsClick={() => (settingsDialogOpen = true)} + onSettingsClick={() => openDialog(() => (settingsDialogOpen = true))} /> {#if uiStore.isSidebarExpanded || isOnMobile} @@ -452,6 +467,12 @@ + + + + + + + + +
+

llama-server

+ + + connecting + + +
+
+ + + + diff --git a/tools/server/dashboard/embed.cmake b/tools/server/dashboard/embed.cmake new file mode 100644 index 000000000000..c49c29be9fc8 --- /dev/null +++ b/tools/server/dashboard/embed.cmake @@ -0,0 +1,21 @@ +# Embed a file as a C byte array header. +# usage: cmake -DINPUT= -DOUTPUT=
-DNAME= -P embed.cmake + +file(READ "${INPUT}" hex HEX) +string(LENGTH "${hex}" hex_len) +math(EXPR n_bytes "${hex_len} / 2") + +# 0xNN, per byte, a line break every 32 bytes +string(REGEX REPLACE "([0-9a-f][0-9a-f])" "0x\\1," bytes "${hex}") +string(REGEX REPLACE "((0x[0-9a-f][0-9a-f],){32})" "\\1\n" bytes "${bytes}") + +file(WRITE "${OUTPUT}.tmp" +"// Generated by tools/server/dashboard/embed.cmake from ${NAME} - do not edit.\n" +"#pragma once\n" +"#include \n" +"static const unsigned char ${NAME}[] = {\n${bytes}\n};\n" +"static const size_t ${NAME}_len = ${n_bytes};\n") + +# only touch the header when the contents change, to avoid needless rebuilds +execute_process(COMMAND ${CMAKE_COMMAND} -E copy_if_different "${OUTPUT}.tmp" "${OUTPUT}") +file(REMOVE "${OUTPUT}.tmp") diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index fac85eb19959..063ccf5b4a41 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -21,6 +21,8 @@ #include "mtmd-helper.h" #include +#include +#include #include #include #include @@ -1124,6 +1126,13 @@ struct server_context_impl { metrics.reset_bucket(); } + // live dashboard stats, see dash_to_json() + // note: not thread-safe, call from the main loop (or right before sleep) + // with_queue = false when the caller holds the task queue lock (sleeping state callback) + json get_dashboard_stats(bool with_queue = true) { + return dash_to_json(with_queue); + } + private: // note: accessing these fields outside of this class is not thread-safe // use server_context methods instead @@ -1180,6 +1189,38 @@ struct server_context_impl { int64_t t_prompt_start = 0; // start of the oldest queued prompt decode uint64_t n_prompt_queued = 0; + // dashboard stats, see GET /dashboard/stats + // note: only touched from the main loop (and from update_cached_responses right before sleep) + struct dash_request { + int64_t t_end_ms; // unix time, ms + int32_t id_slot; + uint64_t n_prompt; // processed prompt tokens (excluding cached) + uint64_t n_cached; // prompt tokens reused from the cache + uint64_t n_gen; // generated tokens + double t_prompt_ms; + double t_gen_ms; + double prompt_tps; + double gen_tps; + uint64_t n_draft; + uint64_t n_draft_accepted; + uint64_t n_draft_verif_steps; + int stop; // stop_type + bool truncated; + }; + struct dash_sample { + int64_t t_us; // ggml_time_us() + int64_t t_unix_ms; + uint64_t n_gen; // cumulative generated tokens (finished + in flight) + uint64_t n_prompt; // cumulative processed prompt tokens + }; + static constexpr size_t DASH_MAX_REQUESTS = 64; + static constexpr size_t DASH_MAX_SAMPLES = 300; // 10 minutes at 2 s + static constexpr int64_t DASH_SAMPLE_PERIOD_US = 2000000; + + std::deque dash_requests; + std::deque dash_samples; + uint64_t dash_n_requests = 0; + json json_ui_settings = json::object(); // Necessary similarity of prompt for slot selection @@ -1762,6 +1803,7 @@ struct server_context_impl { if (slot.stats.n_gen > 0) { metrics_on_prediction(slot); } + dash_on_request_done(slot); }; slot.reset(); @@ -2205,6 +2247,9 @@ struct server_context_impl { } bool launch_slot_with_task(server_slot & slot, server_task && task) { + // anchor the throughput series before the counters start moving after an idle period + dash_sample(); + // process per-request lora adapters if (!task.params.lora.empty()) { auto task_loras = construct_lora_list(task.params.lora); @@ -3085,7 +3130,8 @@ struct server_context_impl { // returns false to decline the task, it is offered again after the decode is done bool process_single_task(server_task && task, bool is_yielding) { // while yielding, an encode / decode is running and only reading the server state is safe - if (is_yielding && task.type != SERVER_TASK_TYPE_METRICS && task.type != SERVER_TASK_TYPE_SLOT_GET) { + if (is_yielding && task.type != SERVER_TASK_TYPE_METRICS && task.type != SERVER_TASK_TYPE_SLOT_GET && + task.type != SERVER_TASK_TYPE_DASHBOARD) { SRV_DBG("decoding, decline task, id_task = %d\n", task.id); return false; } @@ -3232,6 +3278,13 @@ struct server_context_impl { } queue_results.send(std::move(res)); } break; + case SERVER_TASK_TYPE_DASHBOARD: + { + auto res = std::make_unique(); + res->id = task.id; + res->data = dash_to_json(true); + queue_results.send(std::move(res)); + } break; case SERVER_TASK_TYPE_SLOT_GET: { json slots_data = json::array(); @@ -5023,6 +5076,8 @@ struct server_context_impl { slot.stats.set_prompt_last(t_now); } } + + dash_sample(); } // flush any queued prompt metrics if all slots are now idle @@ -5057,6 +5112,236 @@ struct server_context_impl { dst[i] += src[i]; } } + + // + // dashboard + // + + static int64_t dash_unix_ms() { + return std::chrono::duration_cast( + std::chrono::system_clock::now().time_since_epoch()).count(); + } + + static const char * dash_slot_state_str(slot_state state) { + switch (state) { + case SLOT_STATE_IDLE: return "idle"; + case SLOT_STATE_WAIT_OTHER: return "waiting"; + case SLOT_STATE_STARTED: return "started"; + case SLOT_STATE_PROCESSING_PROMPT: return "prompt"; + case SLOT_STATE_DONE_PROMPT: return "prompt_done"; + case SLOT_STATE_GENERATING: return "generating"; + } + return "unknown"; + } + + static const char * dash_stop_str(int stop) { + switch (stop) { + case STOP_TYPE_EOS: return "eos"; + case STOP_TYPE_WORD: return "word"; + case STOP_TYPE_LIMIT: return "limit"; + default: return "none"; + } + } + + // cumulative token counters, including the requests that are still running + uint64_t dash_gen_total() const { + uint64_t n = metrics.predict.count; + for (const auto & slot : slots) { + n += slot.stats.n_gen; + } + return n; + } + + uint64_t dash_prompt_total() const { + return metrics.prompt.count + n_prompt_queued; + } + + // take a throughput sample, at most once per DASH_SAMPLE_PERIOD_US + void dash_sample() { + const int64_t t_now = ggml_time_us(); + if (!dash_samples.empty() && t_now - dash_samples.back().t_us < DASH_SAMPLE_PERIOD_US) { + return; + } + dash_samples.push_back({ t_now, dash_unix_ms(), dash_gen_total(), dash_prompt_total() }); + while (dash_samples.size() > DASH_MAX_SAMPLES) { + dash_samples.pop_front(); + } + } + + void dash_on_request_done(const server_slot & slot) { + const auto & st = slot.stats; + if (!st.is_set() || (st.n_gen == 0 && st.n_prompt_processed == 0 && st.n_prompt_cached == 0)) { + return; + } + dash_n_requests++; + dash_requests.push_back({ + /* t_end_ms */ dash_unix_ms(), + /* id_slot */ slot.id, + /* n_prompt */ st.n_prompt_processed, + /* n_cached */ st.n_prompt_cached, + /* n_gen */ st.n_gen, + /* t_prompt_ms */ st.t_prompt_ms(), + /* t_gen_ms */ st.t_gen_ms(), + /* prompt_tps */ st.n_prompt_tps(), + /* gen_tps */ st.n_gen_tps(), + /* n_draft */ st.n_draft_tokens, + /* n_draft_accepted */ st.n_draft_accepted, + /* n_draft_verif_steps */ st.n_draft_verif_steps, + /* stop */ (int) slot.stop, + /* truncated */ slot.truncated, + }); + while (dash_requests.size() > DASH_MAX_REQUESTS) { + dash_requests.pop_front(); + } + } + + static json dash_draft_json(uint64_t n_draft, uint64_t n_accepted, uint64_t n_steps) { + return json { + {"n_draft", n_draft}, + {"n_accepted", n_accepted}, + {"n_steps", n_steps}, + {"acceptance", n_draft > 0 ? (double) n_accepted / (double) n_draft : 0.0}, + {"mean_accepted_len", n_steps > 0 ? 1.0 + (double) n_accepted / (double) n_steps : 0.0}, + }; + } + + // the live part of GET /dashboard/stats; no prompts or generated text in here + json dash_to_json(bool with_queue) { + dash_sample(); + + const int64_t t_now = ggml_time_us(); + + json jslots = json::array(); + int n_processing = 0; + uint64_t n_kv_used = 0; + for (const auto & slot : slots) { + const auto & st = slot.stats; + const bool busy = slot.is_processing(); + n_processing += busy ? 1 : 0; + n_kv_used += slot.prompt.n_tokens(); + + json js = { + {"id", slot.id}, + {"state", dash_slot_state_str(slot.state)}, + {"n_ctx", slot.n_ctx}, + {"n_tokens", slot.prompt.n_tokens()}, // tokens held in the KV cache by this slot + {"speculative", slot.can_speculate()}, + }; + if (busy && slot.task) { + js["id_task"] = slot.task->id; + js["n_prompt_tokens"] = slot.task->n_tokens(); + js["n_prompt_processed"] = st.n_prompt_processed; + js["n_prompt_cached"] = st.n_prompt_cached; + js["n_gen"] = st.n_gen; + js["n_remaining"] = slot.n_remaining(); + js["elapsed_ms"] = st.is_set() ? st.t_elapsed_us() / 1000.0 : 0.0; + if (slot.state == SLOT_STATE_PROCESSING_PROMPT && st.is_set()) { + // still processing: rate over the time spent so far + const double t_ms = (t_now - st.t_start) / 1000.0; + js["prompt_tps"] = t_ms > 0.0 ? 1e3 / t_ms * st.n_prompt_processed : 0.0; + } else { + js["prompt_tps"] = st.n_prompt_tps(); + } + js["gen_tps"] = st.n_gen_tps(); + js["draft"] = dash_draft_json(st.n_draft_tokens, st.n_draft_accepted, st.n_draft_verif_steps); + } + jslots.push_back(std::move(js)); + } + + // totals: finished requests (server_metrics) + the in-flight ones + uint64_t n_draft = metrics.n_draft_tokens, n_acc = metrics.n_draft_accepted, n_steps = metrics.n_draft_verif_steps; + for (const auto & slot : slots) { + n_draft += slot.stats.n_draft_tokens; + n_acc += slot.stats.n_draft_accepted; + n_steps += slot.stats.n_draft_verif_steps; + } + json jdraft = dash_draft_json(n_draft, n_acc, n_steps); + { + json per_pos = json::array(); + if (metrics.n_draft_verif_steps > 0) { + // drop the trailing never-accepted positions (the vector is sized for n_max) + size_t n = metrics.n_accepted_per_pos.size(); + while (n > 0 && metrics.n_accepted_per_pos[n - 1] == 0) { + n--; + } + for (size_t i = 0; i < n; ++i) { + per_pos.push_back((double) metrics.n_accepted_per_pos[i] / (double) metrics.n_draft_verif_steps); + } + } + jdraft["acceptance_per_pos"] = per_pos; + } + + json jrecent = json::array(); + for (auto it = dash_requests.rbegin(); it != dash_requests.rend(); ++it) { + const auto & r = *it; + jrecent.push_back({ + {"t_end_ms", r.t_end_ms}, + {"id_slot", r.id_slot}, + {"n_prompt", r.n_prompt}, + {"n_cached", r.n_cached}, + {"n_gen", r.n_gen}, + {"t_prompt_ms", r.t_prompt_ms}, + {"t_gen_ms", r.t_gen_ms}, + {"prompt_tps", r.prompt_tps}, + {"gen_tps", r.gen_tps}, + {"draft", dash_draft_json(r.n_draft, r.n_draft_accepted, r.n_draft_verif_steps)}, + {"stop", dash_stop_str(r.stop)}, + {"truncated", r.truncated}, + }); + } + + // throughput series: rate between consecutive samples + json s_t = json::array(), s_gen = json::array(), s_prompt = json::array(); + for (size_t i = 1; i < dash_samples.size(); ++i) { + const auto & a = dash_samples[i - 1]; + const auto & b = dash_samples[i]; + const double dt = (b.t_us - a.t_us) / 1e6; + if (dt <= 0.0) { + continue; + } + s_t .push_back(b.t_unix_ms); + s_gen .push_back(b.n_gen >= a.n_gen ? (b.n_gen - a.n_gen) / dt : 0.0); + s_prompt.push_back(b.n_prompt >= a.n_prompt ? (b.n_prompt - a.n_prompt) / dt : 0.0); + } + + const uint64_t t_start_unix_ms = dash_unix_ms() - (t_now - metrics.t_start) / 1000; + + return json { + {"t_now_ms", dash_unix_ms()}, + {"t_start_ms", t_start_unix_ms}, + {"slots", jslots}, + {"kv", { + {"n_ctx", n_ctx}, + {"n_ctx_slot", n_ctx_slot()}, + {"n_used", n_kv_used}, + {"unified", params_base.kv_unified}, + }}, + {"queue", { + {"n_slots", (int) slots.size()}, + {"n_processing", n_processing}, + {"n_deferred", with_queue ? queue_tasks.queue_tasks_deferred_size() : 0}, + }}, + {"totals", { + {"n_requests", dash_n_requests}, + {"n_prompt", metrics.prompt.count}, + {"n_prompt_cached", metrics.n_prompt_cached}, + {"t_prompt_s", metrics.prompt.time / 1e6}, + {"prompt_tps", metrics.prompt.n_per_second()}, + {"n_gen", dash_gen_total()}, + {"t_gen_s", metrics.predict.time / 1e6}, + {"gen_tps", metrics.predict.n_per_second()}, + {"n_decode", metrics.n_decode}, + {"draft", jdraft}, + }}, + {"recent", jrecent}, + {"series", { + {"period_s", DASH_SAMPLE_PERIOD_US / 1e6}, + {"t_ms", s_t}, + {"gen_tps", s_gen}, + {"prompt_tps", s_prompt}, + }}, + }; + } }; // @@ -5095,6 +5380,13 @@ server_context_meta server_context::get_meta() const { const char * ftype_name = llama_ftype_name(llama_model_ftype(impl->model_tgt)); + char model_arch[128] = {0}; + if (llama_model_meta_val_str(impl->model_tgt, "general.architecture", model_arch, sizeof(model_arch)) < 0) { + model_arch[0] = '\0'; + } + char model_desc[256] = {0}; + llama_model_desc(impl->model_tgt, model_desc, sizeof(model_desc)); + return server_context_meta { /* build_info */ std::string(llama_build_info()), /* model_name */ impl->model_name, @@ -5124,13 +5416,15 @@ server_context_meta server_context::get_meta() const { /* logit_bias_eog */ impl->params_base.sampling.logit_bias_eog, - /* model_vocab_type */ llama_vocab_type(impl->vocab), - /* model_vocab_n_tokens */ llama_vocab_n_tokens(impl->vocab), - /* model_n_ctx_train */ llama_model_n_ctx_train(impl->model_tgt), - /* model_n_embd_inp */ llama_model_n_embd(impl->model_tgt), - /* model_n_params */ llama_model_n_params(impl->model_tgt), - /* model_size */ llama_model_size(impl->model_tgt), - /* model_ftype */ ftype_name, + /* model_vocab_type */ llama_vocab_type(impl->vocab), + /* model_vocab_n_tokens */ llama_vocab_n_tokens(impl->vocab), + /* model_n_ctx_train */ llama_model_n_ctx_train(impl->model_tgt), + /* model_n_embd_inp */ llama_model_n_embd(impl->model_tgt), + /* model_n_params */ llama_model_n_params(impl->model_tgt), + /* model_size */ llama_model_size(impl->model_tgt), + /* model_ftype */ ftype_name, + /* model_arch */ std::string(model_arch), + /* model_desc */ std::string(model_desc), }; } @@ -5554,6 +5848,66 @@ static json get_res_props(const server_context_meta & meta, const common_params return props; } +// static model / config info for GET /dashboard/stats +// note: do NOT use ctx_server here, this must be accessible during sleep +static json get_res_dashboard_info(const server_context_meta & meta, const common_params & params) { + const auto & spec = params.speculative; + const bool has_spec = !(spec.types.empty() || + (spec.types.size() == 1 && spec.types[0] == COMMON_SPECULATIVE_TYPE_NONE)); + + json jspec = nullptr; + if (has_spec) { + std::string types; + for (const auto t : spec.types) { + if (t == COMMON_SPECULATIVE_TYPE_NONE) { + continue; + } + types += (types.empty() ? "" : ",") + common_speculative_type_to_str(t); + } + jspec = { + {"types", types}, + {"n_max", spec.draft.n_max}, + {"n_min", spec.draft.n_min}, + }; + if (spec.has_dft()) { + jspec["draft_model"] = std::filesystem::path(spec.draft.mparams.path).filename().string(); + } + if (spec.draft.mtp_vocab > 0) { + jspec["mtp_vocab"] = spec.draft.mtp_vocab; + } + } + + return json { + {"name", meta.model_name}, + {"aliases", meta.model_aliases}, + {"file", std::filesystem::path(meta.model_path).filename().string()}, + {"arch", meta.model_arch}, + {"desc", meta.model_desc}, + {"ftype", meta.model_ftype}, + {"size_bytes", meta.model_size}, + {"n_params", meta.model_n_params}, + {"n_ctx_train", meta.model_n_ctx_train}, + {"n_vocab", meta.model_vocab_n_tokens}, + {"modalities", { + {"vision", meta.has_inp_image}, + {"audio", meta.has_inp_audio}, + {"video", meta.has_inp_video}, + }}, + {"config", { + {"n_ctx", params.n_ctx}, + {"n_ctx_slot", meta.slot_n_ctx}, + {"n_parallel", params.n_parallel}, + {"kv_unified", params.kv_unified}, + {"cache_type_k", ggml_type_name(params.cache_type_k)}, + {"cache_type_v", ggml_type_name(params.cache_type_v)}, + {"flash_attn", llama_flash_attn_type_name(params.flash_attn_type)}, + {"n_gpu_layers", params.n_gpu_layers}, + {"speculative", jspec}, + {"ple_on_disk", params.ple_on_disk}, + }}, + }; +} + json server_routes::get_model_info() const { return get_res_model_info(*meta); } @@ -5637,6 +5991,59 @@ void server_routes::init_routes() { return res; }; + this->get_dashboard_stats = [this](const server_http_req & req) { + // must not wake the model up, see task_resets_idle_timer() + auto res = create_response(true); + + json stats = nullptr; + bool is_sleeping = queue_tasks.is_sleeping(); + + if (!is_sleeping) { + { + server_task task(SERVER_TASK_TYPE_DASHBOARD); + task.id = res->rd.get_new_id(); + res->rd.post_task(std::move(task), true); // high-priority task + } + + // a task posted right before sleeping is never processed, do not wait for it + auto result = res->rd.next([&]{ + return req.should_stop() || queue_tasks.is_sleeping(); + }); + if (!result) { + if (req.should_stop()) { + return res; + } + is_sleeping = true; + } else if (result->is_error()) { + res->error(result->to_json()); + return res; + } else { + stats = result->to_json(); + } + } + + if (is_sleeping) { + std::unique_lock lock(mutex_cache); + stats = cached_dashboard; + } + + json model = { + {"id", meta->model_name}, + {"status", is_sleeping ? "sleeping" : "loaded"}, + {"info", get_res_dashboard_info(*meta, params)}, + {"stats", stats}, + }; + + res->ok({ + {"mode", "single"}, + {"build_info", meta->build_info}, + {"t_now_ms", std::chrono::duration_cast( + std::chrono::system_clock::now().time_since_epoch()).count()}, + {"models", json::array({ model })}, + }); + return res; + }; + this->get_slots = [this](const server_http_req & req) { auto res = create_response(); if (!params.endpoint_slots) { @@ -6583,6 +6990,8 @@ void server_routes::update_cached_responses(bool is_sleeping) { cached_models = get_res_models(*meta); cached_props = get_res_props(*meta, params, true); cached_metrics = ctx_server.get_metrics(); + // note: called with the task queue locked, so the queue size must not be read + cached_dashboard = ctx_server.get_dashboard_stats(false); should_reset_buckets = false; diff --git a/tools/server/server-context.h b/tools/server/server-context.h index ed23f7d4e6ce..01899f09648b 100644 --- a/tools/server/server-context.h +++ b/tools/server/server-context.h @@ -53,6 +53,8 @@ struct server_context_meta { uint64_t model_n_params; uint64_t model_size; std::string model_ftype; + std::string model_arch; // general.architecture + std::string model_desc; // llama_model_desc() }; enum server_state { @@ -156,6 +158,7 @@ struct server_routes { server_http_context::handler_t post_systemone; server_http_context::handler_t get_lora_adapters; server_http_context::handler_t post_lora_adapters; + server_http_context::handler_t get_dashboard_stats; // to be used in router mode json get_model_info() const; @@ -189,6 +192,7 @@ struct server_routes { json cached_models = nullptr; json cached_props = nullptr; server_metrics cached_metrics; + json cached_dashboard = nullptr; // set when a scrape during sleep already reported the throughput buckets bool should_reset_buckets = false; // call right before sleep to update the cached responses diff --git a/tools/server/server-http.cpp b/tools/server/server-http.cpp index 33cbb156c559..9df71883de48 100644 --- a/tools/server/server-http.cpp +++ b/tools/server/server-http.cpp @@ -1,6 +1,7 @@ #include "common.h" #include "http.h" #include "server-http.h" +#include "dashboard.html.hpp" #include "server-common.h" #include "ui.h" @@ -240,7 +241,7 @@ bool server_http_context::init_listener(const common_params & params) { // Frontend paths - all embedded UI assets static const std::unordered_set frontend_paths = []() { - std::unordered_set paths { "/" }; + std::unordered_set paths { "/", "/dashboard" }; for (const llama_ui_asset & a : llama_ui_get_assets()) { paths.insert("/" + a.name); } @@ -368,6 +369,15 @@ bool server_http_context::init_listener(const common_params & params) { // Web UI setup // + // stats dashboard: a single self-contained page, independent of the chat UI + // the page is public (it holds no data), the JSON it polls (/dashboard/stats) needs the API key + srv->Get(params.api_prefix + "/dashboard", [](const httplib::Request &, httplib::Response & res) { + res.set_header("Cache-Control", "no-cache"); + res.set_header("X-Content-Type-Options", "nosniff"); + res.set_content(reinterpret_cast(dashboard_html), dashboard_html_len, "text/html; charset=utf-8"); + return false; + }); + // Use new `params.ui` field (backed by old `params.webui` for compat) if (!params.ui) { SRV_INF("%s", "The UI is disabled\n"); diff --git a/tools/server/server-models.cpp b/tools/server/server-models.cpp index d6ededbe6383..f2b24d01ac32 100644 --- a/tools/server/server-models.cpp +++ b/tools/server/server-models.cpp @@ -26,6 +26,7 @@ #include #include #include +#include #include #include #include @@ -2152,6 +2153,99 @@ void server_models_routes::init_routes() { return res; }; + this->get_router_dashboard_stats = [this](const server_http_req &) { + // children have no API key (see unset_reserved_args), they only listen on CHILD_ADDR + static constexpr int DASHBOARD_CHILD_TIMEOUT_MS = 1500; + + auto res = std::make_unique(); + auto all_models = models.get_all_meta(); + + // query the running children in parallel, a busy one must not stall the others + std::vector> child_stats; + child_stats.reserve(all_models.size()); + for (const auto & meta : all_models) { + if (meta.hidden || !meta.is_ready_or_sleep() || meta.port <= 0) { + child_stats.push_back({}); + continue; + } + child_stats.push_back(std::async(std::launch::async, [port = meta.port]() -> json { + httplib::Client cli(CHILD_ADDR, port); + cli.set_connection_timeout(0, DASHBOARD_CHILD_TIMEOUT_MS * 1000); + cli.set_read_timeout(0, DASHBOARD_CHILD_TIMEOUT_MS * 1000); + cli.set_write_timeout(0, DASHBOARD_CHILD_TIMEOUT_MS * 1000); + auto resp = cli.Get("/dashboard/stats"); + if (!resp || resp->status != 200) { + return nullptr; + } + try { + json body = json::parse(resp->body); + const auto & arr = body.at("models"); + if (arr.is_array() && !arr.empty()) { + return arr.at(0); + } + } catch (const std::exception &) { + } + return nullptr; + })); + } + + json models_json = json::array(); + for (size_t i = 0; i < all_models.size(); ++i) { + const auto & meta = all_models[i]; + if (meta.hidden) { + continue; + } + + json child = nullptr; + if (child_stats[i].valid()) { + child = child_stats[i].get(); + } + + json entry = child.is_object() ? child : json::object(); + entry["id"] = meta.name; + entry["aliases"] = meta.aliases; + entry["source"] = server_model_source_to_string(meta.source); + // the router's view of the status wins, it also knows unloaded / loading / failed + entry["status"] = server_model_status_to_string(meta.status); + if (meta.is_failed()) { + entry["failed"] = true; + entry["exit_code"] = meta.exit_code; + } + if (meta.status == SERVER_MODEL_STATUS_LOADING && !meta.progress.is_null()) { + entry["progress"] = meta.progress; + } + if (!entry.contains("stats")) { + entry["stats"] = nullptr; + } + if (!entry.contains("info") || !entry["info"].is_object()) { + // not running: what the preset says about the file + json info = json::object(); + std::string model_path; + if (meta.preset.get_option("LLAMA_ARG_MODEL", model_path) && !model_path.empty()) { + info["file"] = std::filesystem::path(model_path).filename().string(); + std::error_code ec; + auto size = std::filesystem::file_size(model_path, ec); + if (!ec) { + info["size_bytes"] = (uint64_t) size; + } + } else if (std::string hf_repo; meta.preset.get_option("LLAMA_ARG_HF_REPO", hf_repo)) { + info["file"] = hf_repo; + } + entry["info"] = info; + } + models_json.push_back(entry); + } + + res_ok(res, { + {"mode", "router"}, + {"build_info", std::string(llama_build_info())}, + {"t_now_ms", std::chrono::duration_cast( + std::chrono::system_clock::now().time_since_epoch()).count()}, + {"models", models_json}, + }); + return res; + }; + this->post_router_models_unload = [this](const server_http_req & req) { auto res = std::make_unique(); json body = json::parse(req.body); diff --git a/tools/server/server-models.h b/tools/server/server-models.h index f415b7c5fd61..89997086d9eb 100644 --- a/tools/server/server-models.h +++ b/tools/server/server-models.h @@ -369,6 +369,7 @@ struct server_models_routes { server_http_context::handler_t proxy_get; server_http_context::handler_t proxy_post; server_http_context::handler_t get_router_models; + server_http_context::handler_t get_router_dashboard_stats; // aggregates /dashboard/stats of the children server_http_context::handler_t post_router_models_load; server_http_context::handler_t post_router_models_unload; // management API diff --git a/tools/server/server-queue.cpp b/tools/server/server-queue.cpp index 3352c3cb5d17..2bed87c3d5f9 100644 --- a/tools/server/server-queue.cpp +++ b/tools/server/server-queue.cpp @@ -22,7 +22,8 @@ // static bool task_resets_idle_timer(server_task_type type) { - return type != SERVER_TASK_TYPE_METRICS; + // polling stats endpoints must not keep the model awake + return type != SERVER_TASK_TYPE_METRICS && type != SERVER_TASK_TYPE_DASHBOARD; } int server_queue::post(server_task && task, bool front) { diff --git a/tools/server/server-task.h b/tools/server/server-task.h index 8515e3a7266c..b3207742dcc8 100644 --- a/tools/server/server-task.h +++ b/tools/server/server-task.h @@ -23,6 +23,7 @@ enum server_task_type { SERVER_TASK_TYPE_NEXT_RESPONSE, SERVER_TASK_TYPE_METRICS, SERVER_TASK_TYPE_SLOT_GET, + SERVER_TASK_TYPE_DASHBOARD, // used by /dashboard/stats SERVER_TASK_TYPE_SLOT_SAVE, SERVER_TASK_TYPE_SLOT_RESTORE, SERVER_TASK_TYPE_SLOT_ERASE, @@ -563,6 +564,15 @@ struct server_task_result_metrics : server_task_result { std::string to_metrics(); }; +// used by /dashboard/stats +struct server_task_result_dashboard : server_task_result { + json data; + + virtual json to_json() override { + return data; + } +}; + // used by /slots API struct server_task_result_slots : server_task_result { int n_idle_slots = 0; diff --git a/tools/server/server.cpp b/tools/server/server.cpp index 871e827bef1d..ce8ab4042f0a 100644 --- a/tools/server/server.cpp +++ b/tools/server/server.cpp @@ -252,6 +252,7 @@ int llama_server(common_params & params, int argc, char ** argv, server_child & // custom routes for router routes.get_props = models_routes->get_router_props; routes.get_models = models_routes->get_router_models; + routes.get_dashboard_stats = models_routes->get_router_dashboard_stats; ctx_http.post("/models", ex_wrapper(models_routes->post_router_models)); ctx_http.post("/models/load", ex_wrapper(models_routes->post_router_models_load)); @@ -301,6 +302,8 @@ int llama_server(common_params & params, int argc, char ** argv, server_child & ctx_http.post("/lora-adapters", ex_wrapper(routes.post_lora_adapters)); // Save & load slots ctx_http.get ("/slots", ex_wrapper(routes.get_slots)); + // dashboard (the page itself is served by server_http_context, see dashboard.html) + ctx_http.get ("/dashboard/stats", ex_wrapper(routes.get_dashboard_stats)); ctx_http.post("/slots/:id_slot", ex_wrapper(routes.post_slots)); // resumable streaming: a child binds the local session factories, the router binds From faef76999fa6034e0c169a88b025ca881fb3cbe2 Mon Sep 17 00:00:00 2001 From: Laurent Zuijdwijk Date: Fri, 2 Oct 2026 15:12:25 +0100 Subject: [PATCH 330/400] server : dashboard reports lazy_mode (strix) instead of ple_on_disk Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_013YKZhadEzjDiTwDc4tzi2g --- tools/server/README.md | 2 +- tools/server/dashboard/dashboard.html | 2 +- tools/server/server-context.cpp | 5 ++++- 3 files changed, 6 insertions(+), 3 deletions(-) diff --git a/tools/server/README.md b/tools/server/README.md index fe9ea7ad9fe4..c8aef6dd5165 100644 --- a/tools/server/README.md +++ b/tools/server/README.md @@ -1187,7 +1187,7 @@ Returns the data behind `/dashboard`. Polling it does not reset the `--sleep-idl "file": "Qwen3.5-2B-UD-Q4_K_XL.gguf", "arch": "qwen35", "ftype": "Q4_K - Medium", "size_bytes": 1328790784, "n_params": 1881825088, "n_ctx_train": 262144, "config": { "n_ctx": 8192, "n_ctx_slot": 4096, "n_parallel": 2, "cache_type_k": "f16", - "speculative": { "types": "ngram-simple", "n_max": 3, "n_min": 0 }, "ple_on_disk": false, ... } + "speculative": { "types": "ngram-simple", "n_max": 3, "n_min": 0 }, "lazy_mode": "auto", ... } }, "stats": { // null when the model is not running "slots": [{ "id": 0, "state": "generating", "n_ctx": 4096, "n_tokens": 118, "n_gen": 45, diff --git a/tools/server/dashboard/dashboard.html b/tools/server/dashboard/dashboard.html index d44d280577cd..a350f663a770 100644 --- a/tools/server/dashboard/dashboard.html +++ b/tools/server/dashboard/dashboard.html @@ -409,7 +409,7 @@

llama-server

['GPU layers', c.n_gpu_layers], ['speculative', spec ? spec.types + ' · n_max ' + spec.n_max + (spec.n_min ? ' · n_min ' + spec.n_min : '') + (spec.draft_model ? ' · draft ' + spec.draft_model : '') + (spec.mtp_vocab ? ' · mtp vocab ' + spec.mtp_vocab : '') : (c.n_ctx !== undefined ? 'off' : null)], - ['n-gram table on disk (PLE)', c.ple_on_disk ? 'yes' : null], + ['lazy tensor reads (PLE)', c.lazy_mode || null], ['vocab', info.n_vocab ? int(info.n_vocab) : null], ['modalities', info.modalities ? ['text'].concat(Object.keys(info.modalities).filter(function (k) { return info.modalities[k]; })).join(', ') : null], ['aliases', m.aliases && m.aliases.length ? m.aliases.join(', ') : null], diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index 063ccf5b4a41..1392d8135781 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -5903,7 +5903,10 @@ static json get_res_dashboard_info(const server_context_meta & meta, const commo {"flash_attn", llama_flash_attn_type_name(params.flash_attn_type)}, {"n_gpu_layers", params.n_gpu_layers}, {"speculative", jspec}, - {"ple_on_disk", params.ple_on_disk}, + {"lazy_mode", params.lazy_mode == LLAMA_LAZY_MODE_OFF ? "off" + : params.lazy_mode == LLAMA_LAZY_MODE_AUTO ? "auto" + : params.lazy_mode == LLAMA_LAZY_MODE_ON ? "on" + : params.lazy_mode == LLAMA_LAZY_MODE_DIRECT ? "direct" : "unknown"}, }}, }; } From b35431512095821359a803a943e19e2147b38ea9 Mon Sep 17 00:00:00 2001 From: Laurent Zuijdwijk Date: Fri, 2 Oct 2026 15:51:30 +0100 Subject: [PATCH 331/400] server : link the chat UI and the /dashboard page to each other MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - dashboard: "← Chat" link at the start of the header (relative "./", so it follows --api-prefix), styled with the page's light/dark tokens and a focus-visible outline. - ui: "Server dashboard" entry in the sidebar action strip (Gauge icon, between Search and Settings), opening ./dashboard in a new tab. Sidebar items get an optional externalHref; ActionIcon gains a target prop. - ui: proxy /dashboard in the vite dev server. The UI change ships only when the UI is built from source (-DLLAMA_BUILD_UI=ON); the prebuilt dist downloaded from the ggml-org HF bucket does not contain it. Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_013YKZhadEzjDiTwDc4tzi2g --- tools/server/dashboard/dashboard.html | 8 ++++ tools/ui/README.md | 3 +- .../components/app/actions/ActionIcon.svelte | 4 ++ .../SidebarNavigationActions.svelte | 47 +++++++++++++++++-- .../ui/src/lib/constants/routes.constants.ts | 6 +++ tools/ui/src/lib/constants/ui.constants.ts | 10 ++-- tools/ui/src/lib/types/navigation.d.ts | 2 + tools/ui/vite.config.ts | 1 + 8 files changed, 73 insertions(+), 8 deletions(-) diff --git a/tools/server/dashboard/dashboard.html b/tools/server/dashboard/dashboard.html index a350f663a770..55fceb08b5cb 100644 --- a/tools/server/dashboard/dashboard.html +++ b/tools/server/dashboard/dashboard.html @@ -94,6 +94,13 @@ } button { cursor: pointer; } button:hover { border-color: var(--text-3); } +a.back { + display: inline-flex; align-items: center; gap: 4px; + color: var(--text-2); text-decoration: none; font-size: 13px; + padding: 3px 8px; margin-left: -8px; border-radius: 6px; border: 1px solid transparent; +} +a.back:hover { color: var(--text); background: var(--surface-2); border-color: var(--border); } +button:focus-visible, input:focus-visible, a.back:focus-visible { outline: 2px solid var(--decode); outline-offset: 1px; } main { max-width: 1280px; margin: 0 auto; padding: 16px; display: grid; gap: 16px; } .card { background: var(--surface); @@ -171,6 +178,7 @@
+ ← Chat

llama-server

diff --git a/tools/ui/README.md b/tools/ui/README.md index 2ecbc8f51255..3fbb971e0556 100644 --- a/tools/ui/README.md +++ b/tools/ui/README.md @@ -128,7 +128,8 @@ proxy: { '/models': SERVER_ORIGIN, '/tools': SERVER_ORIGIN, '/slots': SERVER_ORIGIN, - '/cors-proxy': SERVER_ORIGIN + '/cors-proxy': SERVER_ORIGIN, + '/dashboard': SERVER_ORIGIN }, ``` diff --git a/tools/ui/src/lib/components/app/actions/ActionIcon.svelte b/tools/ui/src/lib/components/app/actions/ActionIcon.svelte index 3116fc530d52..da45526a3cbc 100644 --- a/tools/ui/src/lib/components/app/actions/ActionIcon.svelte +++ b/tools/ui/src/lib/components/app/actions/ActionIcon.svelte @@ -14,6 +14,7 @@ onclick?: (e?: MouseEvent) => void; size?: ButtonSize; stopPropagationOnClick?: boolean; + target?: string; tooltip?: string; tooltipAsTitle?: boolean; tooltipSide?: TooltipSide; @@ -30,6 +31,7 @@ onclick, size = 'sm', stopPropagationOnClick = false, + target, tooltip, tooltipAsTitle = false, tooltipSide = TooltipSide.TOP, @@ -52,7 +54,9 @@ onclick?.(e); }} + rel={href && target === '_blank' ? 'noopener noreferrer' : undefined} {size} + {target} title={tooltipAsTitle ? tooltip : undefined} {variant} > diff --git a/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigationActions.svelte b/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigationActions.svelte index 95efa5372cfe..92d4e991732b 100644 --- a/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigationActions.svelte +++ b/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigationActions.svelte @@ -1,5 +1,5 @@