From f955e394bf94e01e5e36186d13c985727e5ef5b5 Mon Sep 17 00:00:00 2001 From: Michael Lamothe Date: Wed, 15 Jul 2026 18:46:56 +1000 Subject: [PATCH 01/22] ggml: add f16 out_prod support for CPU and out_prod op for Vulkan (#23997) --- ggml/src/ggml-cpu/ggml-cpu.c | 3 +- ggml/src/ggml-cpu/ggml-cpu.cpp | 5 +- ggml/src/ggml-cpu/ops.cpp | 71 +++++++++++++++++-- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 34 +++++++++ .../ggml-vulkan/vulkan-shaders/out_prod.comp | 59 +++++++++++++++ .../vulkan-shaders/vulkan-shaders-gen.cpp | 2 + 6 files changed, 167 insertions(+), 7 deletions(-) create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/out_prod.comp diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index d9347e3c2..0b8d9fac6 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -2859,7 +2859,8 @@ struct ggml_cplan ggml_graph_plan( } break; case GGML_OP_OUT_PROD: { - if (ggml_is_quantized(node->src[0]->type)) { + if (ggml_is_quantized(node->src[0]->type) || + node->src[0]->type == GGML_TYPE_F16) { cur = ggml_type_size(GGML_TYPE_F32) * node->src[0]->ne[0] * n_tasks; } } break; diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index 128883b41..49023ebf6 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -462,11 +462,12 @@ static bool ggml_backend_cpu_device_supports_op(ggml_backend_dev_t dev, const st return max_bias == 0.0f; } case GGML_OP_IM2COL_BACK: - return src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32; + return src0->type == GGML_TYPE_F32 && (src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16); case GGML_OP_GET_ROWS_BACK: return src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16; case GGML_OP_OUT_PROD: - return (src0->type == GGML_TYPE_F32 || (ggml_is_quantized(src0->type) && src0->ne[2] == src1->ne[2] && src0->ne[3] == src1->ne[3])) && + return (src0->type == GGML_TYPE_F32 || + ((src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && src0->ne[2] == src1->ne[2] && src0->ne[3] == src1->ne[3])) && src1->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; default: return true; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index c49719374..58636aa65 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -4449,6 +4449,70 @@ static void ggml_compute_forward_out_prod_q_f32( } } +static void ggml_compute_forward_out_prod_f16_f32( + const ggml_compute_params * params, + ggml_tensor * dst) { + + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * src1 = dst->src[1]; + + GGML_TENSOR_BINARY_OP_LOCALS; + + const int ith = params->ith; + const int nth = params->nth; + + GGML_ASSERT(src0->type == GGML_TYPE_F16); + GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + + GGML_ASSERT(ne02 == ne12); + GGML_ASSERT(ne03 == ne13); + GGML_ASSERT(ne2 == ne12); + GGML_ASSERT(ne3 == ne13); + + GGML_ASSERT(nb00 == sizeof(ggml_fp16_t)); + GGML_ASSERT(nb0 == sizeof(float)); + + GGML_ASSERT(ne0 == ne00); + GGML_ASSERT(ne1 == ne10); + GGML_ASSERT(ne2 == ne02); + GGML_ASSERT(ne3 == ne03); + + if (ith == 0) { + ggml_vec_set_f32(ne0*ne1*ne2*ne3, (float *)dst->data, 0); + } + ggml_barrier(params->threadpool); + + const int64_t nr = ne1*ne2*ne3; + const int64_t dr = (nr + nth - 1)/nth; + const int64_t ir0 = dr*ith; + const int64_t ir1 = MIN(ir0 + dr, nr); + + float * wdata = (float *) params->wdata + (ne0 + CACHE_LINE_SIZE_F32) * ith; + + for (int64_t ir = ir0; ir < ir1; ++ir) { + const int64_t i3 = ir/(ne2*ne1); + const int64_t i2 = (ir - i3*ne2*ne1)/ne1; + const int64_t i1 = (ir - i3*ne2*ne1 - i2*ne1); + + const int64_t i02 = i2; + const int64_t i03 = i3; + + const int64_t i12 = i2; + const int64_t i13 = i3; + + float * d = (float *) ((char *) dst->data + (i1*nb1 + i2*nb2 + i3*nb3)); + + for (int64_t i01 = 0; i01 < ne01; ++i01) { + const int64_t i11 = i01; + ggml_fp16_t * s0 = (ggml_fp16_t *) ((char *) src0->data + (i01*nb01 + i02*nb02 + i03*nb03)); + float * s1 = (float *) ((char *) src1->data + (i1*nb10 + i11*nb11 + i12*nb12 + i13*nb13)); + ggml_fp16_to_fp32_row(s0, wdata, ne0); + ggml_vec_mad_f32(ne0, d, wdata, *s1); + } + } +} + void ggml_compute_forward_out_prod( const ggml_compute_params * params, ggml_tensor * dst) { @@ -4486,9 +4550,8 @@ void ggml_compute_forward_out_prod( } break; case GGML_TYPE_F16: { - GGML_ABORT("fatal error"); // todo - // ggml_compute_forward_out_prod_f16_f32(params, dst); - } + ggml_compute_forward_out_prod_f16_f32(params, dst); + } break; case GGML_TYPE_F32: { ggml_compute_forward_out_prod_f32(params, dst); @@ -6469,7 +6532,7 @@ void ggml_compute_forward_im2col_back_f32( const ggml_tensor * src1 = dst->src[1]; // convolution kernel GGML_ASSERT(src0->type == GGML_TYPE_F32); - GGML_ASSERT(src1->type == GGML_TYPE_F32); + GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16); GGML_ASSERT( dst->type == GGML_TYPE_F32); GGML_TENSOR_BINARY_OP_LOCALS; diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 1704da07e..3c514930e 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -961,6 +961,7 @@ struct vk_device_struct { vk_pipeline pipeline_col2im_1d_f32; vk_pipeline pipeline_col2im_1d_f16; vk_pipeline pipeline_col2im_1d_bf16; + vk_pipeline pipeline_out_prod_f32; vk_pipeline pipeline_snake_f32; vk_pipeline pipeline_snake_f16; vk_pipeline pipeline_snake_bf16; @@ -5479,6 +5480,8 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_col2im_1d_f16, "col2im_1d_f16", col2im_1d_f16_len, col2im_1d_f16_data, "main", 2, sizeof(vk_op_col2im_1d_push_constants), {256, 1, 1}, {}, 1, true); ggml_vk_create_pipeline(device, device->pipeline_col2im_1d_bf16, "col2im_1d_bf16", col2im_1d_bf16_len, col2im_1d_bf16_data, "main", 2, sizeof(vk_op_col2im_1d_push_constants), {256, 1, 1}, {}, 1, true); + ggml_vk_create_pipeline(device, device->pipeline_out_prod_f32, "out_prod_f32", out_prod_f32_len, out_prod_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {256, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_snake_f32, "snake_f32", snake_f32_len, snake_f32_data, "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_snake_f16, "snake_f16", snake_f16_len, snake_f16_data, "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_snake_bf16, "snake_bf16", snake_bf16_len, snake_bf16_data, "main", 4, sizeof(vk_op_snake_push_constants), {256, 1, 1}, {}, 1); @@ -10745,6 +10748,11 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const return ctx->device->pipeline_add_id_f32; } return nullptr; + case GGML_OP_OUT_PROD: + if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { + return ctx->device->pipeline_out_prod_f32; + } + return nullptr; case GGML_OP_CONCAT: { if (src0->type != src1->type || src0->type != dst->type) { return nullptr; @@ -11701,6 +11709,7 @@ static void ggml_vk_op_f32(ggml_backend_vk_context * ctx, vk_context& subctx, co case GGML_OP_DIV: case GGML_OP_MUL: case GGML_OP_ADD1: + case GGML_OP_OUT_PROD: case GGML_OP_ARANGE: case GGML_OP_FILL: case GGML_OP_SCALE: @@ -12014,6 +12023,24 @@ static void ggml_vk_add(ggml_backend_vk_context * ctx, vk_context& subctx, const }); } +static void ggml_vk_out_prod(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { + const uint32_t src0_type_size = ggml_type_size(src0->type); + const uint32_t src1_type_size = ggml_type_size(src1->type); + const uint32_t dst_type_size = ggml_type_size(dst->type); + + ggml_vk_op_f32(ctx, subctx, src0, src1, nullptr, nullptr, dst, GGML_OP_OUT_PROD, { + (uint32_t)ggml_nelements(dst), + (uint32_t)src0->ne[0], (uint32_t)src0->ne[1], (uint32_t)src0->ne[2],(uint32_t)src0->ne[3], + (uint32_t)src0->nb[0] / src0_type_size, (uint32_t)src0->nb[1] / src0_type_size, (uint32_t)src0->nb[2] / src0_type_size, (uint32_t)src0->nb[3] / src0_type_size, + (uint32_t)src1->ne[0], (uint32_t)src1->ne[1], (uint32_t)src1->ne[2],(uint32_t)src1->ne[3], + (uint32_t)src1->nb[0] / src1_type_size, (uint32_t)src1->nb[1] / src1_type_size, (uint32_t)src1->nb[2] / src1_type_size, (uint32_t)src1->nb[3] / src1_type_size, + (uint32_t) dst->ne[0], (uint32_t) dst->ne[1], (uint32_t) dst->ne[2],(uint32_t) dst->ne[3], + (uint32_t) dst->nb[0] / dst_type_size, (uint32_t) dst->nb[1] / dst_type_size, (uint32_t) dst->nb[2] / dst_type_size, (uint32_t) dst->nb[3] / dst_type_size, + 0, + 0.0f, 0.0f, 0, + }); +} + static void ggml_vk_sub(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { const uint32_t src0_type_size = ggml_type_size(src0->type); const uint32_t src1_type_size = ggml_type_size(src1->type); @@ -14801,6 +14828,9 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr ggml_vk_add(ctx, compute_ctx, src0, src1, node); } break; + case GGML_OP_OUT_PROD: + ggml_vk_out_prod(ctx, compute_ctx, src0, src1, node); + break; case GGML_OP_SUB: ggml_vk_sub(ctx, compute_ctx, src0, src1, node); @@ -17655,6 +17685,10 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_OP_OPT_STEP_ADAMW: case GGML_OP_OPT_STEP_SGD: return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32; + case GGML_OP_OUT_PROD: + return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32 + && ggml_is_contiguous(op->src[1]) && op->src[1]->type == GGML_TYPE_F32 + && op->type == GGML_TYPE_F32; case GGML_OP_LOG: case GGML_OP_TRI: case GGML_OP_DIAG: diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/out_prod.comp b/ggml/src/ggml-vulkan/vulkan-shaders/out_prod.comp new file mode 100644 index 000000000..197316996 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/out_prod.comp @@ -0,0 +1,59 @@ +#version 450 + +#extension GL_EXT_shader_16bit_storage : require + +layout (push_constant) uniform parameter +{ + uint ne; + uint ne00; uint ne01; uint ne02; uint ne03; uint nb00; uint nb01; uint nb02; uint nb03; + uint ne10; uint ne11; uint ne12; uint ne13; uint nb10; uint nb11; uint nb12; uint nb13; + uint ne20; uint ne21; uint ne22; uint ne23; uint nb20; uint nb21; uint nb22; uint nb23; + uint misalign_offsets; + float param1; float param2; int param3; +} p; + +layout (binding = 0) readonly buffer A {float data_a[];}; +layout (binding = 1) readonly buffer B {float data_b[];}; +layout (binding = 2) writeonly buffer D {float data_d[];}; + +uint get_idx() { + return gl_GlobalInvocationID.z * 262144 + gl_GlobalInvocationID.y * 512 + gl_GlobalInvocationID.x; +} + +uint get_aoffset() { return p.misalign_offsets >> 16; } +uint get_boffset() { return (p.misalign_offsets >> 8) & 0xFF; } +uint get_doffset() { return p.misalign_offsets & 0xFF; } + +layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in; + +void main() { + uint idx = get_idx(); + if (idx >= p.ne) { + return; + } + + uint tmp = idx; + uint i0 = tmp % p.ne20; tmp /= p.ne20; + uint i1 = tmp % p.ne21; tmp /= p.ne21; + uint i2 = tmp % p.ne22; tmp /= p.ne22; + uint i3 = tmp; + + uint a_i0 = i0 % p.ne00; + uint a_i2 = i2 / (p.ne22 / p.ne02); + uint a_i3 = i3 / (p.ne23 / p.ne03); + + uint b_i0 = i1 % p.ne10; + uint b_i2 = i2; + uint b_i3 = i3; + + float sum = 0.0f; + uint K = p.ne01; + for (uint k = 0; k < K; k++) { + uint aoff = get_aoffset() + a_i3*p.nb03 + a_i2*p.nb02 + k*p.nb01 + a_i0*p.nb00; + uint boff = get_boffset() + b_i3*p.nb13 + b_i2*p.nb12 + k*p.nb11 + b_i0*p.nb10; + sum += data_a[aoff] * data_b[boff]; + } + + uint doff = get_doffset() + i3*p.nb23 + i2*p.nb22 + i1*p.nb21 + i0*p.nb20; + data_d[doff] = sum; +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index 240e1d1b3..bf3eeb22f 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -1036,6 +1036,8 @@ void process_shaders() { } } + string_to_spv("out_prod_f32", "out_prod.comp", {}); + string_to_spv("timestep_embedding_f32", "timestep_embedding.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}})); string_to_spv("conv_transpose_1d_f32", "conv_transpose_1d.comp", {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}}); From b3c9d1b846cc80a6360adb6aeaa4fcd8c4c8dcac Mon Sep 17 00:00:00 2001 From: Pascal Date: Wed, 15 Jul 2026 12:53:31 +0200 Subject: [PATCH 02/22] metal: fuse snake activation (mul, sin, sqr, mul, add) (#25459) * metal: fuse snake activation (mul, sin, sqr, mul, add) Mirror the CUDA, Vulkan and CPU snake fusion: same matcher on the naive 5-op chain, same F32 contract on a and inv_b, same F32/F16/BF16 kernel with F32 compute. Follows the Metal backend idioms: bf16 instantiation gated behind GGML_METAL_HAS_BF16 and concurrency ranges checked on the remaining chain nodes before encoding, as done by the bin fusion. Covered by the existing backend-agnostic SNAKE_FUSE tests. * metal: absorb snake fusion into ggml_metal_op_bin Extract the matcher to ggml_metal_op_can_fuse_snake, mirroring the Vulkan naming, and dispatch the fused path from ggml_metal_op_bin. The encode loop switch is back to a single call per case. Address review from ggerganov * metal: fix indentation in ggml_metal_op_can_fuse_snake --- ggml/src/ggml-metal/ggml-metal-device.cpp | 17 ++++ ggml/src/ggml-metal/ggml-metal-device.h | 1 + ggml/src/ggml-metal/ggml-metal-impl.h | 5 ++ ggml/src/ggml-metal/ggml-metal-ops.cpp | 100 ++++++++++++++++++++++ ggml/src/ggml-metal/ggml-metal-ops.h | 1 + ggml/src/ggml-metal/ggml-metal.metal | 29 +++++++ 6 files changed, 153 insertions(+) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 870f93df0..270c1411a 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -1834,6 +1834,23 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d(ggml_m return res; } +ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_snake(ggml_metal_library_t lib, enum ggml_type type) { + GGML_ASSERT(type == GGML_TYPE_F32 || type == GGML_TYPE_F16 || type == GGML_TYPE_BF16); + + char base[256]; + char name[256]; + + snprintf(base, 256, "kernel_snake_%s", ggml_type_name(type)); + snprintf(name, 256, "%s", base); + + ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); + if (!res.pipeline) { + res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); + } + + return res; +} + ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d(ggml_metal_library_t lib, const ggml_tensor * op) { assert(op->op == GGML_OP_CONV_TRANSPOSE_2D); diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h index 9d4aca121..b36fa8110 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.h +++ b/ggml/src/ggml-metal/ggml-metal-device.h @@ -151,6 +151,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_im2col struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_transpose_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_col2im_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); +struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_snake (ggml_metal_library_t lib, enum ggml_type type); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d (ggml_metal_library_t lib, const struct ggml_tensor * op); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_2d_dw (ggml_metal_library_t lib, const struct ggml_tensor * op, bool tiled); struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_conv_3d (ggml_metal_library_t lib, const struct ggml_tensor * op); diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index a45f3ac67..330278d00 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -616,6 +616,11 @@ typedef struct { int32_t p0; } ggml_metal_kargs_col2im_1d; +typedef struct { + int32_t T; + int32_t C; +} ggml_metal_kargs_snake; + typedef struct { int32_t IC; int32_t IH; diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 805bc4093..c716f118f 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -3077,7 +3077,58 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) { return 1; } +// Snake activation autofuse: mul -> sin -> sqr -> mul -> add +static bool ggml_metal_op_can_fuse_snake(ggml_metal_op_t ctx, int idx) { + static constexpr ggml_op snake_ops[5] = { GGML_OP_MUL, GGML_OP_SIN, GGML_OP_SQR, GGML_OP_MUL, GGML_OP_ADD }; + + if (ctx->node(idx)->op != GGML_OP_MUL || !ctx->can_fuse(idx, snake_ops, 5)) { + return false; + } + + const ggml_tensor * mul0 = ctx->node(idx + 0); + const ggml_tensor * sin_node = ctx->node(idx + 1); + const ggml_tensor * sqr = ctx->node(idx + 2); + const ggml_tensor * mul1 = ctx->node(idx + 3); + const ggml_tensor * add = ctx->node(idx + 4); + + // x carries the full activation shape, a is the broadcast operand + const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1]; + const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0]; + + // mul1 reads sqr and inv_b in either operand order + const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0]; + + // closure check: the trailing add reads the same x as the leading mul + const ggml_tensor * x_in_add = (add->src[0] == mul1) ? add->src[1] : add->src[0]; + + // x is in the supported whitelist and every chain intermediate shares x's type. + // a and inv_b bind as device const float * in the kernel, so they stay F32. + const bool types_ok = + (x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16) && + (a->type == GGML_TYPE_F32) && (inv_b->type == GGML_TYPE_F32) && + (mul0->type == x->type) && (sin_node->type == x->type) && + (sqr->type == x->type) && (mul1->type == x->type) && + (add->type == x->type); + // a / inv_b collapse to [1, C, 1, 1], x and add stay 2D + const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1]; + const bool dim_ok = + (x->ne[2] == 1) && (x->ne[3] == 1) && + (add->ne[2] == 1) && (add->ne[3] == 1) && + (a->ne[2] == 1) && (a->ne[3] == 1) && + (inv_b->ne[2] == 1) && (inv_b->ne[3] == 1); + // kernel reads x[idx] and a[c] / inv_b[c] linearly, so every operand is contiguous + const bool contig_ok = + ggml_is_contiguous(x) && ggml_is_contiguous(add) && + ggml_is_contiguous(a) && ggml_is_contiguous(inv_b); + + return types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x; +} + int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) { + if (ctx->use_fusion && ggml_metal_op_can_fuse_snake(ctx, idx)) { + return ggml_metal_op_snake_fused(ctx, idx); + } + ggml_tensor * op = ctx->node(idx); ggml_metal_library_t lib = ctx->lib; @@ -3984,6 +4035,55 @@ int ggml_metal_op_col2im_1d(ggml_metal_op_t ctx, int idx) { return 1; } +// Dispatch the fused snake kernel from the matched mul -> sin -> sqr -> mul -> add chain. +// idx points at the leading mul. The caller has validated the chain. +int ggml_metal_op_snake_fused(ggml_metal_op_t ctx, int idx) { + ggml_metal_library_t lib = ctx->lib; + ggml_metal_encoder_t enc = ctx->enc; + + const ggml_tensor * mul0 = ctx->node(idx + 0); + const ggml_tensor * sqr = ctx->node(idx + 2); + const ggml_tensor * mul1 = ctx->node(idx + 3); + ggml_tensor * add = ctx->node(idx + 4); + + const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1]; + const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0]; + const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0]; + + const int T = (int) x->ne[0]; + const int C = (int) x->ne[1]; + const int total = T * C; + + // the encode loop pre-checked the leading mul only, check the rest of the chain + for (int i = 1; i < 5; ++i) { + if (!ggml_metal_op_concurrency_check(ctx, ctx->node(idx + i))) { + ggml_metal_op_concurrency_reset(ctx); + + break; + } + } + + auto pipeline = ggml_metal_library_get_pipeline_snake(lib, x->type); + + ggml_metal_kargs_snake args = { + /*.T =*/ T, + /*.C =*/ C, + }; + + ggml_metal_encoder_set_pipeline(enc, pipeline); + ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(x), 1); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(a), 2); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(inv_b), 3); + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(add), 4); + + const int nth = 256; + const int ntg = (total + nth - 1) / nth; + ggml_metal_encoder_dispatch_threadgroups(enc, ntg, 1, 1, nth, 1, 1); + + return 5; +} + int ggml_metal_op_conv_transpose_2d(ggml_metal_op_t ctx, int idx) { ggml_tensor * op = ctx->node(idx); diff --git a/ggml/src/ggml-metal/ggml-metal-ops.h b/ggml/src/ggml-metal/ggml-metal-ops.h index 0bebd836a..89a6ad82f 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.h +++ b/ggml/src/ggml-metal/ggml-metal-ops.h @@ -80,6 +80,7 @@ int ggml_metal_op_conv_3d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_transpose_1d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_conv_transpose_2d (ggml_metal_op_t ctx, int idx); int ggml_metal_op_col2im_1d (ggml_metal_op_t ctx, int idx); +int ggml_metal_op_snake_fused (ggml_metal_op_t ctx, int idx); int ggml_metal_op_upscale (ggml_metal_op_t ctx, int idx); int ggml_metal_op_pad (ggml_metal_op_t ctx, int idx); int ggml_metal_op_pad_reflect_1d (ggml_metal_op_t ctx, int idx); diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 38919ebad..969fddfa5 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -5406,6 +5406,35 @@ template [[host_name("kernel_col2im_1d_bf16")]] kernel void kernel_col2im_1d +kernel void kernel_snake( + constant ggml_metal_kargs_snake & args, + device const T * x, + device const float * a, + device const float * inv_b, + device T * dst, + uint tgpig [[threadgroup_position_in_grid]], + uint tpitg [[thread_position_in_threadgroup]], + uint ntg [[threads_per_threadgroup]]) { + + const int idx = tgpig * ntg + tpitg; + if (idx >= args.T * args.C) { + return; + } + + const int c = idx / args.T; // x is [T, C], a / inv_b collapse to [1, C] + const float xi = float(x[idx]); + const float si = sin(a[c] * xi); + dst[idx] = T(xi + si * si * inv_b[c]); +} + +template [[host_name("kernel_snake_f32")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const float *, device const float *, device const float *, device float *, uint, uint, uint); +template [[host_name("kernel_snake_f16")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const half *, device const float *, device const float *, device half *, uint, uint, uint); +#if defined(GGML_METAL_HAS_BF16) +template [[host_name("kernel_snake_bf16")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const bfloat *, device const float *, device const float *, device bfloat *, uint, uint, uint); +#endif + + typedef void (conv_transpose_2d_t)( constant ggml_metal_kargs_conv_transpose_2d & args, device const float * src0, From c81029373dd72d8d0f8ebbc925d3c9554a274617 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Wed, 15 Jul 2026 14:34:53 +0300 Subject: [PATCH 03/22] ci : add HF_TOKEN to self-hosted workflows (#25706) * ci : add HF_TOKEN to self-hosted workflows Pass the HF_TOKEN_CI repo secret as HF_TOKEN env var in the self-hosted build and server workflows. Fix the stale build.yml path reference. Assisted-by: pi:llama.cpp/Qwen3.6-27B * cont : add comment --------- Co-authored-by: ggerganov --- .github/workflows/build-self-hosted.yml | 4 +++- .github/workflows/server-self-hosted.yml | 2 ++ 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/.github/workflows/build-self-hosted.yml b/.github/workflows/build-self-hosted.yml index 1a71ed827..441a897e5 100644 --- a/.github/workflows/build-self-hosted.yml +++ b/.github/workflows/build-self-hosted.yml @@ -6,7 +6,7 @@ on: branches: - master paths: [ - '.github/workflows/build.yml', + '.github/workflows/build-self-hosted.yml', '**/CMakeLists.txt', '**/.cmake', '**/*.h', @@ -48,6 +48,8 @@ concurrency: cancel-in-progress: true env: + # note: this is dud token to avoid rate limiting (https://github.com/ggml-org/llama.cpp/pull/25706#issuecomment-4979941302) + HF_TOKEN: ${{ secrets.HF_TOKEN_CI }} GGML_NLOOP: 3 GGML_N_THREADS: 1 LLAMA_ARG_LOG_COLORS: 1 diff --git a/.github/workflows/server-self-hosted.yml b/.github/workflows/server-self-hosted.yml index 2dcd6d742..d8266d7ee 100644 --- a/.github/workflows/server-self-hosted.yml +++ b/.github/workflows/server-self-hosted.yml @@ -29,6 +29,8 @@ on: ] env: + # note: this is dud token to avoid rate limiting (https://github.com/ggml-org/llama.cpp/pull/25706#issuecomment-4979941302) + HF_TOKEN: ${{ secrets.HF_TOKEN_CI }} LLAMA_ARG_LOG_COLORS: 1 LLAMA_ARG_LOG_PREFIX: 1 LLAMA_ARG_LOG_TIMESTAMPS: 1 From a3e5b96ac5e278c390df429df0b68efcee3ee1b5 Mon Sep 17 00:00:00 2001 From: fairydreaming <166155368+fairydreaming@users.noreply.github.com> Date: Wed, 15 Jul 2026 13:36:32 +0200 Subject: [PATCH 04/22] cuda : relax tensor contiguity requirements for quantized concat (#25678) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * cuda : relax tensor contiguity requirements for quantized concat * tests : add test cases for non-contiguous quantized concat * ggml : relax contiguity requirements for quantized concat --------- Co-authored-by: Stanisław Szymczyk --- ggml/src/ggml-cpu/ops.cpp | 4 ++-- ggml/src/ggml-cuda/concat.cu | 41 ++++++++++++++++++--------------- ggml/src/ggml-cuda/ggml-cuda.cu | 14 +++++++++-- tests/test-backend-ops.cpp | 22 +++++++++++++++--- 4 files changed, 55 insertions(+), 26 deletions(-) diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 58636aa65..8022ab49d 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -2081,8 +2081,8 @@ void ggml_compute_forward_concat( const ggml_tensor * src1 = dst->src[1]; if (ggml_is_quantized(src0->type)) { - GGML_ASSERT(ggml_is_contiguous(src0)); - GGML_ASSERT(ggml_is_contiguous(src1)); + GGML_ASSERT(ggml_is_contiguous_rows(src0)); + GGML_ASSERT(ggml_is_contiguous_rows(src1)); GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); } diff --git a/ggml/src/ggml-cuda/concat.cu b/ggml/src/ggml-cuda/concat.cu index 276ee64e8..6df89013c 100644 --- a/ggml/src/ggml-cuda/concat.cu +++ b/ggml/src/ggml-cuda/concat.cu @@ -141,27 +141,25 @@ static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) template static void concat_cuda(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, int dim, cudaStream_t stream) { - if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) { + if (dim != 3 && ggml_is_contiguous_to_3(src0) && ggml_is_contiguous_to_3(src1)) { const T * src0_d = (const T *) src0->data; const T * src1_d = (const T *) src1->data; T * dst_d = (T *) dst->data; - if (dim != 3) { - for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) { - concat_cont_cuda( - src0_d + i3*(src0->nb[3] / sizeof(T)), - src1_d + i3*(src1->nb[3] / sizeof(T)), - dst_d + i3*( dst->nb[3] / sizeof(T)), - ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2], - ggml_row_size(dst->type, dst->ne[0])/sizeof(T), dst->ne[1], dst->ne[2], dim, stream); - } - } else { - const size_t size0 = ggml_nbytes(src0); - const size_t size1 = ggml_nbytes(src1); - - CUDA_CHECK(cudaMemcpyAsync((char *) dst->data, src0->data, size0, cudaMemcpyDeviceToDevice, stream)); - CUDA_CHECK(cudaMemcpyAsync((char *) dst->data + size0, src1->data, size1, cudaMemcpyDeviceToDevice, stream)); + for (int64_t i3 = 0; i3 < dst->ne[3]; i3++) { + concat_cont_cuda( + src0_d + i3*(src0->nb[3] / sizeof(T)), + src1_d + i3*(src1->nb[3] / sizeof(T)), + dst_d + i3*( dst->nb[3] / sizeof(T)), + ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2], + ggml_row_size(dst->type, dst->ne[0])/sizeof(T), dst->ne[1], dst->ne[2], dim, stream); } + } else if (dim == 3 && ggml_is_contiguous(src0) && ggml_is_contiguous(src1)) { + const size_t size0 = ggml_nbytes(src0); + const size_t size1 = ggml_nbytes(src1); + + CUDA_CHECK(cudaMemcpyAsync((char *) dst->data, src0->data, size0, cudaMemcpyDeviceToDevice, stream)); + CUDA_CHECK(cudaMemcpyAsync((char *) dst->data + size0, src1->data, size1, cudaMemcpyDeviceToDevice, stream)); } else { GGML_ASSERT(!ggml_is_quantized(src0->type)); @@ -208,12 +206,17 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { GGML_ASSERT(dst->type == src0->type); if (ggml_is_quantized(src0->type)) { - GGML_ASSERT(ggml_is_contiguous(src0)); - GGML_ASSERT(ggml_is_contiguous(src1)); + if (dim == 3) { + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(src1)); + } else { + GGML_ASSERT(ggml_is_contiguous_to_3(src0)); + GGML_ASSERT(ggml_is_contiguous_to_3(src1)); + } GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); - // if tensors are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors + // if first 3 dimensions are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors concat_cuda(src0, src1, dst, dim, stream); } else { GGML_ASSERT(ggml_blck_size(src0->type) == 1); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 0878ab9c0..0e185e849 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4816,13 +4816,23 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g { ggml_type src0_type = op->src[0]->type; ggml_type src1_type = op->src[1]->type; + const int32_t dim = op->op_params[0]; return src0_type == src1_type && src0_type == op->type && ( ( ggml_is_quantized(src0_type) && - ggml_is_contiguous(op->src[0]) && - ggml_is_contiguous(op->src[1]) && + ( + ( + dim == 3 && + ggml_is_contiguous(op->src[0]) && + ggml_is_contiguous(op->src[1]) + ) || ( + dim != 3 && + ggml_is_contiguous_to_3(op->src[0]) && + ggml_is_contiguous_to_3(op->src[1]) + ) + ) && op->src[0]->ne[0] % ggml_blck_size(src0_type) == 0 && op->src[1]->ne[0] % ggml_blck_size(src0_type) == 0 ) || ( diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index ae49d2d0d..fae0e8eb9 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -5555,7 +5555,7 @@ struct test_concat : public test_case { const std::array ne_a; const int64_t ne_b_d; const int dim; - const int v; // view (1 << 0: non-cont a, 1 << 1: non-cont b) + const int v; // view (1 << 0: non-cont a (first 3 dim), 1 << 1: non-cont b (first 3 dim), 1 << 2: non-cont a (last 2 dim), 1 << 3: non-cont b (last 2 dim)) std::string vars() override { return VARS_TO_STR5(type, ne_a, ne_b_d, dim, v); @@ -5576,6 +5576,13 @@ struct test_concat : public test_case { a = ggml_new_tensor(ctx, type, 4, ne.data()); ggml_set_name(a, "a"); + a = ggml_view_4d(ctx, a, ne_a[0], ne_a[1], ne_a[2], ne_a[3], a->nb[1], a->nb[2], a->nb[3], 0); + ggml_set_name(a, "view_of_a"); + } else if (v & 4) { + auto ne = ne_a; ne[2] *= 2; ne[3] *= 4; + a = ggml_new_tensor(ctx, type, 4, ne.data()); + ggml_set_name(a, "a"); + a = ggml_view_4d(ctx, a, ne_a[0], ne_a[1], ne_a[2], ne_a[3], a->nb[1], a->nb[2], a->nb[3], 0); ggml_set_name(a, "view_of_a"); } else { @@ -5588,6 +5595,13 @@ struct test_concat : public test_case { b = ggml_new_tensor(ctx, type, 4, ne.data()); ggml_set_name(b, "b"); + b = ggml_view_4d(ctx, b, ne_b[0], ne_b[1], ne_b[2], ne_b[3], b->nb[1], b->nb[2], b->nb[3], 0); + ggml_set_name(b, "view_of_b"); + } else if (v & 8) { + auto ne = ne_b; ne[2] *= 3; ne[3] *= 2; + b = ggml_new_tensor(ctx, type, 4, ne.data()); + ggml_set_name(b, "b"); + b = ggml_view_4d(ctx, b, ne_b[0], ne_b[1], ne_b[2], ne_b[3], b->nb[1], b->nb[2], b->nb[3], 0); ggml_set_name(b, "view_of_b"); } else { @@ -9089,8 +9103,10 @@ static std::vector> make_test_cases_eval() { } for (ggml_type type_a : { GGML_TYPE_Q4_0, GGML_TYPE_Q4_1, GGML_TYPE_Q5_0, GGML_TYPE_Q5_1, GGML_TYPE_Q8_0 }) { - for (int dim : { 0, 1, 2, 3, }) { - test_cases.emplace_back(new test_concat(type_a, {128, 12, 13, 14}, dim == 0 ? 256 : 7, dim, 0)); + for (int v : { 0, 4, 8, 12 }) { + for (int dim : { 0, 1, 2, 3, }) { + test_cases.emplace_back(new test_concat(type_a, {128, 12, 13, 14}, dim == 0 ? 256 : 7, dim, v)); + } } } From a05df0a81a99d4c87023e98b294c48fdd9878833 Mon Sep 17 00:00:00 2001 From: Pascal Date: Wed, 15 Jul 2026 13:39:21 +0200 Subject: [PATCH 05/22] ui: fix thinking menu never appearing in single-model mode (#25637) In MODEL mode, modelPropsCache is never populated: fetchModelProps call sites are gated on router-only state (isRouterMode checks, routerModels always empty), so supportsThinking always reads an empty chat template once a model is auto-selected. Read serverStore.props.chat_template directly in non-router mode, since the global /props already describes the single loaded model. --- tools/ui/src/lib/stores/models.svelte.ts | 40 ++++++++++++++---------- 1 file changed, 23 insertions(+), 17 deletions(-) diff --git a/tools/ui/src/lib/stores/models.svelte.ts b/tools/ui/src/lib/stores/models.svelte.ts index 4c7c7cdfe..0b4d7b55d 100644 --- a/tools/ui/src/lib/stores/models.svelte.ts +++ b/tools/ui/src/lib/stores/models.svelte.ts @@ -245,21 +245,20 @@ class ModelsStore { * Whether the selected model's chat template supports thinking/reasoning. * Uses heuristic detection on the model's chat_template from /props. * - * - MODEL mode: uses serverStore.props.chat_template (single loaded model) - * - ROUTER mode: fetches /props?model= for the selected model (cached) - * - * Triggers an async fetch of model props if not yet cached in ROUTER mode. + * - MODEL mode: the global /props already describes the single loaded model, + * so its chat_template is used directly and no per-model cache is involved + * - ROUTER mode: fetches /props?model= for the selected model (cached), + * triggering an async fetch if not yet cached */ get supportsThinking(): boolean { - const modelId = this.selectedModelName; - if (!modelId) { - if (!isRouterMode()) { - return detectThinkingSupport(serverStore.props?.chat_template ?? ''); - } - return false; + if (!isRouterMode()) { + return detectThinkingSupport(serverStore.props?.chat_template ?? ''); } - if (isRouterMode() && !this.modelPropsCache.get(modelId)) { + const modelId = this.selectedModelName; + if (!modelId) return false; + + if (!this.modelPropsCache.get(modelId)) { this.fetchModelProps(modelId); } const props = this.getModelProps(modelId); @@ -268,12 +267,17 @@ class ModelsStore { /** * Check if a specific model supports thinking. - * Fetches model props if not cached (in router mode). + * In MODEL mode the global /props describes the single loaded model. + * In ROUTER mode, fetches model props if not cached. */ checkModelSupportsThinking(modelId: string): boolean { + if (!isRouterMode()) { + return detectThinkingSupport(serverStore.props?.chat_template ?? ''); + } + if (!modelId) return false; - if (isRouterMode() && !this.modelPropsCache.get(modelId)) { + if (!this.modelPropsCache.get(modelId)) { this.fetchModelProps(modelId); } @@ -285,14 +289,16 @@ class ModelsStore { * Detailed thinking support detection result with reason for debugging/UI. */ get thinkingSupportDetails(): { supported: boolean; reason: string } { + if (!isRouterMode()) { + return detectThinkingSupportWithReason(serverStore.props?.chat_template ?? ''); + } + const modelId = this.selectedModelName; if (!modelId) { - if (!isRouterMode()) { - return detectThinkingSupportWithReason(serverStore.props?.chat_template ?? ''); - } return { supported: false, reason: 'No model selected' }; } - if (isRouterMode() && !this.modelPropsCache.get(modelId)) { + + if (!this.modelPropsCache.get(modelId)) { this.fetchModelProps(modelId); } const props = this.getModelProps(modelId); From a5822222909b785f23ddc74ce3c8f85bd0e38562 Mon Sep 17 00:00:00 2001 From: Pascal Date: Wed, 15 Jul 2026 13:46:46 +0200 Subject: [PATCH 06/22] server: fix read_file append_loc space breaking edit_file match (#25705) read_file with append_loc emits "{n}\u2192 {line}". The space after the arrow is meant as a separator, but it is indistinguishable from real indentation. Models strip "{n}\u2192" yet keep the space, so the old_text passed to edit_file carries a phantom leading space and never matches (normalize_for_fuzzy_match trims trailing whitespace only, never leading). Drop the separator space so the arrow abuts content: stripping "{n}\u2192" now yields the exact line with its real indentation preserved, and the failure mode cannot occur by construction. Update the description example to match the new format. --- tools/server/server-tools.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tools/server/server-tools.cpp b/tools/server/server-tools.cpp index a8216d7db..9eb57abae 100644 --- a/tools/server/server-tools.cpp +++ b/tools/server/server-tools.cpp @@ -289,7 +289,7 @@ struct server_tool_read_file : server_tool { {"function", { {"name", name}, {"description", "Read the contents of a file. Optionally specify a 1-based line range. " - "If append_loc is true, each line is prefixed with its line number (e.g. \"1\u2192 ...\")."}, + "If append_loc is true, each line is prefixed with its line number (e.g. \"1\u2192...\")."}, {"parameters", { {"type", "object"}, {"properties", { @@ -339,7 +339,7 @@ struct server_tool_read_file : server_tool { std::string out_line; if (append_loc) { - out_line = std::to_string(lineno) + "\u2192 " + line + "\n"; + out_line = std::to_string(lineno) + "\u2192" + line + "\n"; } else { out_line = line + "\n"; } From 956973c76466b6c791d7bdbe6eed3aa3235b2dc1 Mon Sep 17 00:00:00 2001 From: Gaurav Garg Date: Wed, 15 Jul 2026 19:51:34 +0530 Subject: [PATCH 07/22] Fix crash with draft-simple (#25720) * Fix crash with draft-simple * Fix tests for spec decoding --- common/speculative.cpp | 5 ++++- tools/server/tests/unit/test_speculative.py | 11 +++++++---- tools/server/tests/utils.py | 3 +++ 3 files changed, 14 insertions(+), 5 deletions(-) diff --git a/common/speculative.cpp b/common/speculative.cpp index 580728a20..3cb08767b 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -260,7 +260,10 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { bool process(const llama_batch & batch) override { auto * ctx_dft = params.ctx_dft; - const int ret = llama_decode(ctx_dft, batch); + llama_batch batch_dft = batch; + batch_dft.logits = nullptr; + + const int ret = llama_decode(ctx_dft, batch_dft); if (ret != 0) { SPC_ERR("failed to decode draft batch, ret = %d\n", ret); diff --git a/tools/server/tests/unit/test_speculative.py b/tools/server/tests/unit/test_speculative.py index 84cd77e6f..c6568479c 100644 --- a/tools/server/tests/unit/test_speculative.py +++ b/tools/server/tests/unit/test_speculative.py @@ -12,8 +12,9 @@ def create_server(): server = ServerPreset.stories15m_moe() # set default values server.model_draft = download_file(MODEL_DRAFT_FILE_URL) - server.draft_min = 4 - server.draft_max = 8 + server.spec_type = "draft-simple" + server.spec_draft_n_min = 4 + server.spec_draft_n_max = 8 server.fa = "off" @@ -25,6 +26,7 @@ def fixture_create_server(): def test_with_and_without_draft(): global server server.model_draft = None # disable draft model + server.spec_type = None server.start() res = server.make_request("POST", "/completion", data={ "prompt": "I believe the meaning of life is", @@ -46,6 +48,7 @@ def test_with_and_without_draft(): "n_predict": 16, }) assert res.status_code == 200 + assert res.body["timings"]["draft_n"] > 0 content_draft = res.body["content"] assert content_no_draft == content_draft @@ -63,8 +66,8 @@ def test_different_draft_min_draft_max(): last_content = None for draft_min, draft_max in test_values: server.stop() - server.draft_min = draft_min - server.draft_max = draft_max + server.spec_draft_n_min = draft_min + server.spec_draft_n_max = draft_max server.start() res = server.make_request("POST", "/completion", data={ "prompt": "I believe the meaning of life is", diff --git a/tools/server/tests/utils.py b/tools/server/tests/utils.py index f4f0e61e6..5d5c873ac 100644 --- a/tools/server/tests/utils.py +++ b/tools/server/tests/utils.py @@ -95,6 +95,7 @@ class ServerProcess: no_models_autoload: bool | None = None lora_files: List[str] | None = None enable_ctx_shift: int | None = False + spec_type: str | None = None spec_draft_n_min: int | None = None spec_draft_n_max: int | None = None no_ui: bool | None = None @@ -226,6 +227,8 @@ class ServerProcess: server_args.extend(["--lora", lora_file]) if self.enable_ctx_shift: server_args.append("--context-shift") + if self.spec_type: + server_args.extend(["--spec-type", self.spec_type]) if self.api_key: server_args.extend(["--api-key", self.api_key]) if self.spec_draft_n_max: From f6f12e43fa869ef0e008b99ed97dc4006bbb8907 Mon Sep 17 00:00:00 2001 From: leonardHONG <2695316095@qq.com> Date: Wed, 15 Jul 2026 23:21:22 +0800 Subject: [PATCH 08/22] CUDA: tighter MMQ src1 buffer size for native fp4 (#25613) --- ggml/src/ggml-cuda/mmq.cu | 13 +++++++------ ggml/src/ggml-cuda/mmq.cuh | 13 ++++++++----- ggml/src/ggml-cuda/quantize.cu | 14 +++++++------- 3 files changed, 22 insertions(+), 18 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index bf9f5d526..34e5cbf8c 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -122,11 +122,12 @@ void ggml_cuda_mul_mat_q( const bool fallback = ne01 % 128 != 0; - // TODO: tighter pool buffer size vs q8 path const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4); + const size_t y_block_size = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq); + const size_t y_values_per_block = use_native_fp4 ? QK_FP4_MMQ : QK8_1_MMQ; if (!ids) { - const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1 + + const size_t nbytes_src1_q8_1 = ne13*ne12 * ne11*ne10_padded * y_block_size/y_values_per_block + ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq); ggml_cuda_pool_alloc src1_q8_1(ctx.pool(), nbytes_src1_q8_1); @@ -148,7 +149,7 @@ void ggml_cuda_mul_mat_q( // Stride depends on quantization format const int64_t s12 = use_native_fp4 ? - ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : // block_fp4_mmq holds 256 values + ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_FP4_MMQ * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; @@ -184,7 +185,7 @@ void ggml_cuda_mul_mat_q( CUDA_CHECK(cudaGetLastError()); } - const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * sizeof(block_q8_1)/QK8_1 + + const size_t nbytes_src1_q8_1 = ne12*n_expert_used*ne10_padded * y_block_size/y_values_per_block + ggml_cuda_mmq_get_J_max(src0->type, fallback, cc, ne11) * sizeof(block_q8_1_mmq); ggml_cuda_pool_alloc src1_q8_1(ctx.pool(), nbytes_src1_q8_1); @@ -207,8 +208,8 @@ void ggml_cuda_mul_mat_q( CUDA_CHECK(cudaGetLastError()); } - static_assert(QK_K == 8 * QK_MXFP4, "QK_K needs to be 8 * QK_MXFP4"); - const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_K * sizeof(int)) : + static_assert(QK_FP4_MMQ == 8 * QK_MXFP4, "QK_FP4_MMQ needs to be 8 * QK_MXFP4"); + const int64_t s12 = use_native_fp4 ? ne11 * ne10_padded * sizeof(block_fp4_mmq) / (QK_FP4_MMQ * sizeof(int)) : ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 607e433bf..ad3d6a5a6 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -21,6 +21,9 @@ enum mmq_q8_1_ds_layout { MMQ_Q8_1_DS_LAYOUT_D2S6, }; +static constexpr int QK8_1_MMQ = 4*QK8_1; +static constexpr int QK_FP4_MMQ = 2*QK8_1_MMQ; + struct block_q8_1_mmq { // The y float data is converted to a data layout that can simply be copied to shared memory as a contiguous block. // The y float data is first grouped as blocks of 128 values. @@ -39,7 +42,7 @@ struct block_q8_1_mmq { half d2s6[8]; // 1 16 bit scale per 64 values + 1 16 bit partial sum per 16 values for the first 96 values, // stored as d0,d1,s1,s2,s3,s4,s5 }; - int8_t qs[4*QK8_1]; // 128 values quantized to 8 bit each + int8_t qs[QK8_1_MMQ]; }; // this struct is used for fp4 data types (currently only used for Blackwell) @@ -47,10 +50,10 @@ struct block_q8_1_mmq { // nvfp4 has block size 16, each int32 of d4 contains 4 ue4m3 scales struct block_fp4_mmq { uint32_t d4[4]; - int8_t qs[4 * 32]; // 256 FP4 values packed as 4-bit pairs (2 per byte) + int8_t qs[QK_FP4_MMQ / 2]; }; -static_assert(sizeof(block_q8_1_mmq) == 4*QK8_1 + 4*sizeof(half2), "Unexpected block_q8_1_mmq size"); +static_assert(sizeof(block_q8_1_mmq) == QK8_1_MMQ + 4*sizeof(half2), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_q8_1_mmq) == 4*sizeof(block_q8_1), "Unexpected block_q8_1_mmq size"); static_assert(sizeof(block_fp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected block_fp4_mmq size"); @@ -833,9 +836,9 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( #if defined(BLACKWELL_MMA_AVAILABLE) // FP4 tile stores 8 blocks - constexpr int ne_block = (type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4) ? QK_K : 4 * QK8_1; + constexpr int ne_block = (type == GGML_TYPE_MXFP4 || type == GGML_TYPE_NVFP4) ? QK_FP4_MMQ : QK8_1_MMQ; #else - constexpr int ne_block = 4 * QK8_1; + constexpr int ne_block = QK8_1_MMQ; #endif // defined(BLACKWELL_MMA_AVAILABLE) constexpr int ITER_K = ggml_cuda_mmq_get_K_vram(type, J, fallback); diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 39a500a17..a7d450737 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -90,8 +90,8 @@ static __global__ void quantize_mmq_nvfp4( const int64_t i2 = blockIdx.z % ne2; const int64_t i3 = blockIdx.z / ne2; const int64_t i01 = ids ? ids[i1] : i1; - const int64_t k_block = i0_base / QK_K; - const int64_t blocks_per_col = (ne0 + QK_K - 1) / QK_K; + const int64_t k_block = i0_base / QK_FP4_MMQ; + const int64_t blocks_per_col = (ne0 + QK_FP4_MMQ - 1) / QK_FP4_MMQ; if (k_block >= blocks_per_col) { return; } @@ -100,7 +100,7 @@ static __global__ void quantize_mmq_nvfp4( block_fp4_mmq * y = (block_fp4_mmq *) vy; block_fp4_mmq * yb = y + ib; - const int sub = (i0_base % QK_K) / QK_NVFP4_SUB; + const int sub = (i0_base % QK_FP4_MMQ) / QK_NVFP4_SUB; float vals_raw[QK_NVFP4_SUB]; float amax_raw = 0.0f; @@ -207,7 +207,7 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, block_fp4_mmq * y = (block_fp4_mmq *) vy; - const int64_t block_fp4_mmq_size = 8 * QK_MXFP4; // 256 values + const int64_t block_fp4_mmq_size = QK_FP4_MMQ; const int64_t ib0 = blockIdx.z * ((int64_t) ne1 * (ne0 / block_fp4_mmq_size)); const int64_t ib = ib0 + (warp_start_offset / block_fp4_mmq_size) * ne1 + blockIdx.x; const int64_t quad_idx_in_block = (warp_start_offset % block_fp4_mmq_size) / vals_per_warp; @@ -303,8 +303,8 @@ static __global__ void quantize_mmq_q8_1( block_q8_1_mmq * y = (block_q8_1_mmq *) vy; const int64_t ib0 = blockIdx.z*((int64_t)gridDim.x*gridDim.y*blockDim.x/QK8_1); // first block of channel - const int64_t ib = ib0 + (i0 / (4*QK8_1))*ne1 + blockIdx.x; // block index in channel - const int64_t iqs = i0 % (4*QK8_1); // quant index in block + const int64_t ib = ib0 + (i0 / QK8_1_MMQ)*ne1 + blockIdx.x; // block index in channel + const int64_t iqs = i0 % QK8_1_MMQ; // quant index in block // Load 4 floats per thread and calculate max. abs. value between them: const float4 xi = i0 < ne00 ? x4[(i03*s03 + i02*s02 + i01*s01 + i00)/4] : make_float4(0.0f, 0.0f, 0.0f, 0.0f); @@ -394,7 +394,7 @@ void quantize_mmq_q8_1_cuda( const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { GGML_ASSERT(ne00 % 4 == 0); - GGML_ASSERT(ne0 % (4*QK8_1) == 0); + GGML_ASSERT(ne0 % QK8_1_MMQ == 0); // ne1 tends to assume the highest values, therefore use it as the "x" dimension of the CUDA grid: const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ); From c3d47e696b1187a27e896aa828d48ff9a33fc679 Mon Sep 17 00:00:00 2001 From: Hongqiang Wang Date: Wed, 15 Jul 2026 09:08:40 -0700 Subject: [PATCH 09/22] opencl: fix two issues on flash attention for Adreno a7x (#25697) * opencl: route `sub_group_shuffle_xor` to qcom ext when KHR ext is unavailable KHR `sub_group_shuffle_xor` is not defined by compiler when `cl_qcom_subgroup_shuffle` is present, causing certain FA kernels fail to build. Define the KHR shuffle_xor using the qcom extension. * opencl: skip FA kernels with mixed and quant types for A7x to avoid compiler crash --- ggml/src/ggml-opencl/ggml-opencl.cpp | 8 ++++++++ ggml/src/ggml-opencl/kernels/flash_attn_f32_f16.cl | 4 ++++ ggml/src/ggml-opencl/kernels/flash_attn_f32_q4_0.cl | 4 ++++ ggml/src/ggml-opencl/kernels/flash_attn_f32_q8_0.cl | 4 ++++ ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl | 4 ++++ 5 files changed, 24 insertions(+) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index b14ea8133..16e851bbd 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -7319,6 +7319,14 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te return false; } + // Some compilers for A7x (Adreno 740, compiler E031.41) crashes when + // building FA kernels with mixed or quant types (f32_f16, f32_q8_0, f32_q4_0) + // Here we skip all A7x for these kernels to avoid crash + if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::A7X && + (is_f32_f16 || is_f32_q8_0 || is_f32_q4_0)) { + return false; + } + if (dk == 512) { if (backend_ctx->gpu_family == INTEL) { return false; diff --git a/ggml/src/ggml-opencl/kernels/flash_attn_f32_f16.cl b/ggml/src/ggml-opencl/kernels/flash_attn_f32_f16.cl index 1cc0cc8c3..6e43ee81e 100644 --- a/ggml/src/ggml-opencl/kernels/flash_attn_f32_f16.cl +++ b/ggml/src/ggml-opencl/kernels/flash_attn_f32_f16.cl @@ -30,6 +30,10 @@ #elif defined(cl_qcom_subgroup_shuffle) #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable #define HAS_SUBGROUP_SHUFFLE 1 +// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR +// name, so calling it is an implicit declaration and the program fails to build. +// Route it to the qcom builtin. +#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f) #endif #define ACC_TYPE float diff --git a/ggml/src/ggml-opencl/kernels/flash_attn_f32_q4_0.cl b/ggml/src/ggml-opencl/kernels/flash_attn_f32_q4_0.cl index de09a1eaa..95d215971 100644 --- a/ggml/src/ggml-opencl/kernels/flash_attn_f32_q4_0.cl +++ b/ggml/src/ggml-opencl/kernels/flash_attn_f32_q4_0.cl @@ -10,6 +10,10 @@ #elif defined(cl_qcom_subgroup_shuffle) #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable #define HAS_SUBGROUP_SHUFFLE 1 +// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR +// name, so calling it is an implicit declaration and the program fails to build. +// Route it to the qcom builtin. +#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f) #endif // Flash attention: Q=f32, K=q4_0, V=q4_0. diff --git a/ggml/src/ggml-opencl/kernels/flash_attn_f32_q8_0.cl b/ggml/src/ggml-opencl/kernels/flash_attn_f32_q8_0.cl index 46bc4bc9d..7e89ed0bd 100644 --- a/ggml/src/ggml-opencl/kernels/flash_attn_f32_q8_0.cl +++ b/ggml/src/ggml-opencl/kernels/flash_attn_f32_q8_0.cl @@ -10,6 +10,10 @@ #elif defined(cl_qcom_subgroup_shuffle) #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable #define HAS_SUBGROUP_SHUFFLE 1 +// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR +// name, so calling it is an implicit declaration and the program fails to build. +// Route it to the qcom builtin. +#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f) #endif // Flash attention: Q=f32, K=q8_0, V=q8_0. diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl index da2e14ae9..97148d370 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl +++ b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl @@ -24,6 +24,10 @@ #elif defined(cl_qcom_subgroup_shuffle) #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable #define HAS_SUBGROUP_SHUFFLE 1 +// Adreno compilers that expose only cl_qcom_subgroup_shuffle do not declare the KHR +// name, so calling it is an implicit declaration and the program fails to build. +// Route it to the qcom builtin. +#define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f) #endif // Assumes row size (ne00) is a multiple of 4 From aff6eb6e7503538fec1532dec2f584bc7a4a4e4d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adrien=20Gallou=C3=ABt?= Date: Wed, 15 Jul 2026 18:41:51 +0200 Subject: [PATCH 10/22] tokenize : drop --stdin mutual-exclusion check (#25672) match cli and completion, which don't enforce it --- common/arg.cpp | 2 +- tools/tokenize/tokenize.cpp | 12 ++---------- 2 files changed, 3 insertions(+), 11 deletions(-) diff --git a/common/arg.cpp b/common/arg.cpp index b6fddae00..7bc770ab4 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -2932,7 +2932,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex ).set_examples({LLAMA_EXAMPLE_TOKENIZE})); add_opt(common_arg( {"--stdin"}, - string_format("read the prompt from stdin (mutually exclusive with -f/--file and -p/--prompt) (default: %s)", params.tokenize_stdin ? "true" : "false"), + string_format("read the prompt from stdin (takes precedence over -f/--file and -p/--prompt) (default: %s)", params.tokenize_stdin ? "true" : "false"), [](common_params & params) { params.tokenize_stdin = true; } diff --git a/tools/tokenize/tokenize.cpp b/tools/tokenize/tokenize.cpp index 23120ad2e..77b33c4a4 100644 --- a/tools/tokenize/tokenize.cpp +++ b/tools/tokenize/tokenize.cpp @@ -103,19 +103,11 @@ int main(int argc, char ** argv) { return 1; } - // which prompt source was requested? - // -p/--prompt and -f/--file both end up in params.prompt (common's -f also - // strips a single trailing newline), but -f additionally records the path - // in params.prompt_file, so we use that to tell them apart. + // -f and -p both land in params.prompt; -f also sets prompt_file. -f and -p + // resolve like the other tools (no mutual exclusion), --stdin takes precedence. const bool use_stdin = params.tokenize_stdin; const bool use_file = !params.prompt_file.empty(); - // sanity check: --stdin is mutually exclusive with -f/--file and -p/--prompt - if (use_stdin && (use_file || !params.prompt.empty())) { - LOG_ERR("error: --stdin is mutually exclusive with --file and --prompt\n"); - return 1; - } - // must have some prompt if (!use_stdin && !use_file && params.prompt.empty()) { LOG_ERR("error: must specify one of: --stdin, --file or --prompt\n"); From 3b53219361a61b53e7741c479b81b755ec6096b1 Mon Sep 17 00:00:00 2001 From: fairydreaming <166155368+fairydreaming@users.noreply.github.com> Date: Wed, 15 Jul 2026 19:57:52 +0200 Subject: [PATCH 11/22] cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) (#25545) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) * chore : remove indentation of #pragma unroll * cuda : remove unnecessary kernel template declarations * cuda : add WARPS_PER_BLOCK and K_VECS_PER_BLOCK template parameters in lightning indexer kernels to avoid duplication of constants. * cuda : relax MMA architecture requirements to Turing in lightning indexer implementation * chore : renamed variables * chore : rename ggml_cuda_op_lightning_indexer() to ggml_cuda_lightning_indexer() * chore : TODO for AMD rocWMMA * chore : whitespace formatting * chore : another variable rename to fix problems caused by shadowing * chore : yet another rename, this time uppercased all constants * cuda : added alignment checks for Q and K tensors in lightning indexer implementation --------- Co-authored-by: Stanisław Szymczyk --- ggml/src/ggml-cuda/ggml-cuda.cu | 6 + ggml/src/ggml-cuda/lightning-indexer.cu | 588 +++++++++++++++++++++++ ggml/src/ggml-cuda/lightning-indexer.cuh | 4 + 3 files changed, 598 insertions(+) create mode 100644 ggml/src/ggml-cuda/lightning-indexer.cu create mode 100644 ggml/src/ggml-cuda/lightning-indexer.cuh diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 0e185e849..e27ab1526 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -65,6 +65,7 @@ #include "ggml-cuda/tri.cuh" #include "ggml-cuda/cumsum.cuh" #include "ggml-cuda/fill.cuh" +#include "ggml-cuda/lightning-indexer.cuh" #include "ggml.h" #include @@ -2257,6 +2258,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_FILL: ggml_cuda_op_fill(ctx, dst); break; + case GGML_OP_LIGHTNING_INDEXER: + ggml_cuda_lightning_indexer(ctx, dst); + break; default: return false; } @@ -4987,6 +4991,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_DIAG: case GGML_OP_SOLVE_TRI: return true; + case GGML_OP_LIGHTNING_INDEXER: + return ggml_cuda_lightning_indexer_supported(dev_ctx->device, op); default: return false; diff --git a/ggml/src/ggml-cuda/lightning-indexer.cu b/ggml/src/ggml-cuda/lightning-indexer.cu new file mode 100644 index 000000000..5edc967e0 --- /dev/null +++ b/ggml/src/ggml-cuda/lightning-indexer.cu @@ -0,0 +1,588 @@ +#include "common.cuh" +#include "lightning-indexer.cuh" +#include "fattn-common.cuh" +#include "convert.cuh" + +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) +#if defined(TURING_MMA_AVAILABLE) + +typedef union { + int2 i2; + half2 h2[2]; +} half4; + +// TODO add support for AMD cards via rocWMMA +#include +namespace wmma = nvcuda::wmma; + +template +static __global__ void lightning_indexer_kernel_wmma( + const float * Q, const char * K, const float * W, const half * M, float * dst, + int64_t n_stream, int64_t n_batch, int64_t n_kv, + size_t nb1, size_t nb2, size_t nb3, + size_t nbq1, size_t nbq2, size_t nbq3, + size_t nbk1, size_t nbk2, size_t nbk3, + size_t nbw1, size_t nbw2, size_t nbw3, + size_t nbm1, size_t nbm2, size_t nbm3, + int64_t nem3 + ) { + + constexpr int THREADS_PER_BLOCK = WARPS_PER_BLOCK * WARP_SIZE; + constexpr int HEADS_PER_INNER_LOOP = 8; + constexpr int K_EMBD_PER_INNER_LOOP = 16; + constexpr int N_EMBD_PADDED = N_EMBD + 8; + + const int i_batch = blockIdx.y; + const int i_stream = blockIdx.z; + const int i_warp = threadIdx.y; + const int i_lane = threadIdx.x; + const int tid = i_warp * WARP_SIZE + i_lane; + + // each block processes K_VECS_PER_BLOCK K vectors + const int start_kv = blockIdx.x * K_VECS_PER_BLOCK; + + const char * q_base = (const char *) Q + i_batch*nbq2 + i_stream*nbq3; + const float * w_base = (const float *) ((const char *) W + i_batch*nbw1 + i_stream*nbw3); + + // phase 1 - load weights and first Q tile to shared memory + + __shared__ float w_shared[N_HEAD]; + __shared__ int2 q_shared_h[HEADS_PER_INNER_LOOP][N_EMBD_PADDED / 4]; + + if (tid < N_HEAD) { + w_shared[tid] = w_base[tid]; + } + + // total number of half4 elements in HEADS_PER_INNER_LOOP x N_EMBD Q tile + constexpr int N_Q_TILE = HEADS_PER_INNER_LOOP * (N_EMBD / 4); + // number of registers needed in each thread to store Q tile in thread block + constexpr int N_Q_NEXT = (N_Q_TILE + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; + +#pragma unroll + for (int i_q = tid; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) { + const int i_head = i_q / (N_EMBD / 4); + const int i_embd = i_q % (N_EMBD / 4); + const float4 q = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4)); + half4 q_packed; + q_packed.h2[0] = __float22half2_rn(make_float2(q.x, q.y)); + q_packed.h2[1] = __float22half2_rn(make_float2(q.z, q.w)); + q_shared_h[i_head][i_embd] = q_packed.i2; + } + + // phase 2 - load (and dequantize if needed) K to shared mem + + __shared__ half2 k_shared_h[K_VECS_PER_BLOCK][N_EMBD_PADDED / 4][2]; + + constexpr int n_k = K_VECS_PER_BLOCK * (N_EMBD / 4); + + if constexpr (TYPE_K == GGML_TYPE_F16) { +#pragma unroll + for (int i_k = tid; i_k < n_k; i_k += THREADS_PER_BLOCK) { + const int i_k_vec = i_k / (N_EMBD / 4); + const int i_embd = i_k % (N_EMBD / 4); + const int i_kv = start_kv + i_k_vec; + if (i_kv < n_kv) { + const int2 * k_base = (const int2 *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + *(int2*) &k_shared_h[i_k_vec][i_embd] = k_base[i_embd]; + } else { + *(int2*) &k_shared_h[i_k_vec][i_embd] = make_int2(0, 0); + } + } + } else { + constexpr dequantize_V_t dequantize_k = get_dequantize_V(); +#pragma unroll + for (int i_k = tid; i_k < n_k; i_k += THREADS_PER_BLOCK) { + const int i_k_vec = i_k / (N_EMBD / 4); + const int i_embd = i_k % (N_EMBD / 4); + const int i_kv = start_kv + i_k_vec; + if (i_kv < n_kv) { + const void * k_base = (const void *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + dequantize_k(k_base, &k_shared_h[i_k_vec][i_embd][0], i_embd * 4); + } else { + *(int2*) &k_shared_h[i_k_vec][i_embd] = make_int2(0, 0); + } + } + } + + __syncthreads(); + + // phase 3 - calculate lightning indexer scores + + __shared__ float qk_shared[WARPS_PER_BLOCK][HEADS_PER_INNER_LOOP][K_VECS_PER_BLOCK]; + + // load K fragment + wmma::fragment frag_k; + wmma::load_matrix_sync(frag_k, (half*) &k_shared_h[0][i_warp * K_EMBD_PER_INNER_LOOP / 4], N_EMBD_PADDED); + + float score_k = 0.0f; + + for (int i_head_0 = 0; i_head_0 < N_HEAD; i_head_0 += HEADS_PER_INNER_LOOP) { + const int i_head_next = i_head_0 + HEADS_PER_INNER_LOOP; + + // we don't use accumulator for anything, fill it with zeros + wmma::fragment frag_acc; + wmma::fill_fragment(frag_acc, 0.0f); + + // load Q fragment + wmma::fragment frag_q; + wmma::load_matrix_sync(frag_q, (half*) &q_shared_h[0][i_warp * K_EMBD_PER_INNER_LOOP / 4], N_EMBD_PADDED); + + // preload next Q tile to registers during matrix multiplication + float4 q_next[N_Q_NEXT]; + + if (i_head_next < N_HEAD) { +#pragma unroll + for (int i_q = tid, i_q_next = 0; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) { + const int i_head = i_head_next + i_q / (N_EMBD / 4); + const int i_embd = i_q % (N_EMBD / 4); + q_next[i_q_next++] = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4)); + } + } + + // perform matrix multiplication + wmma::mma_sync(frag_acc, frag_q, frag_k, frag_acc); + wmma::store_matrix_sync((float*) &qk_shared[i_warp][0][0], frag_acc, K_VECS_PER_BLOCK, wmma::mem_row_major); + + // make sure all threads finished using q_shared_h so we can store next tile + __syncthreads(); + + // write preloaded Q tile to shared memory + if (i_head_next < N_HEAD) { +#pragma unroll + for (int i_q = tid, i_q_next = 0; i_q < N_Q_TILE; i_q += THREADS_PER_BLOCK) { + const int i_head = i_q / (N_EMBD / 4); + const int i_embd = i_q % (N_EMBD / 4); + half4 q_packed; + q_packed.h2[0] = __float22half2_rn(make_float2(q_next[i_q_next].x, q_next[i_q_next].y)); + q_packed.h2[1] = __float22half2_rn(make_float2(q_next[i_q_next].z, q_next[i_q_next].w)); + q_shared_h[i_head][i_embd] = q_packed.i2; + ++i_q_next; + } + } + + // accumulate QK multiplication results from all block warps + // (there are 256 threads in block and 256 matmul outputs) + // TODO it will break if WARP_SIZE is not 32 + const int h = tid / K_VECS_PER_BLOCK; + const int k = tid % K_VECS_PER_BLOCK; + const float w_val = w_shared[i_head_0 + h]; + + float sum = 0.0f; +#pragma unroll + for (int w = 0; w < WARPS_PER_BLOCK; ++w) { + sum += qk_shared[w][h][k]; + } + + // ReLU, weight + sum = sum > 0.0f ? sum : 0.0f; + sum *= w_val; + + // wait until qk_shared[0] is no longer used + __syncthreads(); + + // reuse qk_shared[0] for storing partial results + qk_shared[0][h][k] = sum; + + // wait until all threads write their results + __syncthreads(); + + // accumulate result over heads + if (tid < K_VECS_PER_BLOCK) { +#pragma unroll + for (int i_head = 0; i_head < HEADS_PER_INNER_LOOP; ++i_head) { + score_k += qk_shared[0][i_head][tid]; + } + } + + // make sure all threads finished using qk_shared + __syncthreads(); + } + + // phase 4 - store output to VRAM + + if (tid < K_VECS_PER_BLOCK) { + const int i_kv = start_kv + tid; + if (i_kv < n_kv) { + const half * m_base = (const half *) ((const char *) M + i_batch*nbm1 + (i_stream%nem3)*nbm3); + float * dst_base = (float *) ((char *) dst + i_batch*nb1 + i_stream*nb3); + dst_base[i_kv] = score_k + __half2float(m_base[i_kv]); + } + } +} + +#else // defined(TURING_MMA_AVAILABLE) + +template +static __global__ void lightning_indexer_kernel_wmma( + const float * Q, const char * K, const float * W, const half * M, float * dst, + int64_t n_stream, int64_t n_batch, int64_t n_kv, + size_t nb1, size_t nb2, size_t nb3, + size_t nbq1, size_t nbq2, size_t nbq3, + size_t nbk1, size_t nbk2, size_t nbk3, + size_t nbw1, size_t nbw2, size_t nbw3, + size_t nbm1, size_t nbm2, size_t nbm3, + int64_t nem3 + ) { + GGML_UNUSED_VARS(Q, K, W, M, dst, + n_stream, n_batch, n_kv, + nb1, nb2, nb3, + nbq1, nbq2, nbq3, + nbk1, nbk2, nbk3, + nbw1, nbw2, nbw3, + nem3); + NO_DEVICE_CODE; +} + +#endif // defined(TURING_MMA_AVAILABLE) +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + +// TODO there is one ugly assumption used in this kernel - that WARP_SIZE is equal to 32 +// thanks to that one warp operating on float4 processes whole indexer K/Q vectors +// 32 * 4 = 128 (N_EMBD) + +template +static __global__ void lightning_indexer_kernel_vec( + const float * Q, const char * K, const float * W, const half * M, float * dst, + int64_t n_stream, int64_t n_batch, int64_t n_kv, + size_t nb1, size_t nb2, size_t nb3, + size_t nbq1, size_t nbq2, size_t nbq3, + size_t nbk1, size_t nbk2, size_t nbk3, + size_t nbw1, size_t nbw2, size_t nbw3, + size_t nbm1, size_t nbm2, size_t nbm3, + int64_t nem3 + ) { + + constexpr int K_VECS_PER_WARP = K_VECS_PER_BLOCK / WARPS_PER_BLOCK; + constexpr int THREADS_PER_BLOCK = WARPS_PER_BLOCK * WARP_SIZE; + + const int i_batch = blockIdx.y; + const int i_stream = blockIdx.z; + const int i_warp = threadIdx.y; + const int i_lane = threadIdx.x; + const int tid = i_warp * WARP_SIZE + i_lane; + + // each warp processes K_VECS_PER_WARP K vectors + const int start_kv_block = blockIdx.x * K_VECS_PER_BLOCK; + const int start_kv = start_kv_block + i_warp * K_VECS_PER_WARP; + + const char * q_base = (const char *) Q + i_batch*nbq2 + i_stream*nbq3; + const float * w_base = (const float *) ((const char *) W + i_batch*nbw1 + i_stream*nbw3); + + // phase 1 - load (and dequantize if needed) K to registers + + float4 k_reg_f[K_VECS_PER_WARP]; + + if constexpr (TYPE_K == GGML_TYPE_F32) { + // direct copy of float4 +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + int i_kv = start_kv + k; + if (i_kv < n_kv) { + const float4 * k_base = (const float4 *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + k_reg_f[k] = k_base[i_lane]; + } else { + k_reg_f[k] = make_float4(0, 0, 0, 0); + } + } + } else { + // dequantize remaining types to float + constexpr dequantize_V_t dequantize_k = get_dequantize_V(); +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + int i_kv = start_kv + k; + if (i_kv < n_kv) { + const void * k_base = (const void *) ((const char *) K + i_kv*nbk2 + i_stream*nbk3); + dequantize_k(k_base, &k_reg_f[k], i_lane * 4); + } else { + k_reg_f[k] = make_float4(0, 0, 0, 0); + } + } + } + + float score_k[K_VECS_PER_WARP] = { 0.0f }; + + // load weights and Q only for N_HEAD_INNER heads at once to reduce shared memory usage + constexpr int N_HEAD_INNER = N_HEAD / 4; + + for (int i_head_0 = 0; i_head_0 < N_HEAD; i_head_0 += N_HEAD_INNER) { + // phase 2 - load weights and Q to shared memory + + __shared__ float w_shared[N_HEAD_INNER]; + __shared__ float4 q_shared_f[N_HEAD_INNER][N_EMBD / 4]; + + if (tid < N_HEAD_INNER) { + w_shared[tid] = w_base[i_head_0 + tid]; + } + + constexpr int n_q = N_HEAD_INNER * (N_EMBD / 4); +#pragma unroll + for (int i_q = tid; i_q < n_q; i_q += THREADS_PER_BLOCK) { + const int i_head_inner = i_q / (N_EMBD / 4); + const int i_head = i_head_0 + i_head_inner; + const int i_embd = i_q % (N_EMBD / 4); + q_shared_f[i_head_inner][i_embd] = *(const float4 *) (q_base + i_head*nbq1 + i_embd*sizeof(float4)); + } + + __syncthreads(); + + // phase 3 - calculate lightning indexer scores + + for (int i_head_inner = 0; i_head_inner < N_HEAD_INNER; ++i_head_inner) { + const float w_val = w_shared[i_head_inner]; + float qk[K_VECS_PER_WARP] = { 0.0f }; + + // dot product of floats + const float4 q_vec = q_shared_f[i_head_inner][i_lane]; + +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + ggml_cuda_mad(qk[k], q_vec.x, k_reg_f[k].x); + ggml_cuda_mad(qk[k], q_vec.y, k_reg_f[k].y); + ggml_cuda_mad(qk[k], q_vec.z, k_reg_f[k].z); + ggml_cuda_mad(qk[k], q_vec.w, k_reg_f[k].w); + } + +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + float sum = warp_reduce_sum(qk[k]); + + // ReLU, weight + if (i_lane == 0) { + sum = (sum > 0.0f) ? sum : 0.0f; + score_k[k] += sum * w_val; + } + } + } + + __syncthreads(); + } + + // phase 4 - store outputs to shared memory + + __shared__ float dst_shared[K_VECS_PER_BLOCK]; + + if (i_lane == 0) { +#pragma unroll + for (int k = 0; k < K_VECS_PER_WARP; ++k) { + dst_shared[i_warp * K_VECS_PER_WARP + k] = score_k[k]; + } + } + + __syncthreads(); + + // phase 5 - write from shared memory to VRAM in coalesced manner + + if (tid < K_VECS_PER_BLOCK) { + int i_kv = start_kv_block + tid; + if (i_kv < n_kv) { + const half * m_base = (const half *) ((const char *) M + i_batch*nbm1 + (i_stream%nem3)*nbm3); + float * dst_base = (float *) ((char *) dst + i_batch*nb1 + i_stream*nb3); + dst_base[i_kv] = dst_shared[tid] + __half2float(m_base[i_kv]); + } + } +} + +#define LIGHTNING_INDEXER_CASE(lightning_indexer_kernel, n_embd, n_head, K, type_K) \ + if (K->type == (type_K)) { \ + lightning_indexer_kernel \ + <<>>( \ + q_d, k_d, w_d, m_d, dst_d, \ + n_stream, n_batch, n_kv, \ + nb1, nb2, nb3, \ + nbq1, nbq2, nbq3, \ + nbk1, nbk2, nbk3, \ + nbw1, nbw2, nbw3, \ + nbm1, nbm2, nbm3, \ + nem3 \ + ); \ + } else + +void ggml_cuda_lightning_indexer(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * q = dst->src[0]; + const ggml_tensor * k = dst->src[1]; + const ggml_tensor * w = dst->src[2]; // weights + const ggml_tensor * m = dst->src[3]; // mask + + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT( q->type == GGML_TYPE_F32); + GGML_ASSERT( w->type == GGML_TYPE_F32); + GGML_ASSERT( m->type == GGML_TYPE_F16); + + GGML_TENSOR_LOCALS(int64_t, neq, q, ne) + GGML_TENSOR_LOCALS(size_t, nbq, q, nb) + GGML_TENSOR_LOCALS(int64_t, nek, k, ne) + GGML_TENSOR_LOCALS(size_t, nbk, k, nb) + GGML_TENSOR_LOCALS(int64_t, new, w, ne) + GGML_TENSOR_LOCALS(size_t, nbw, w, nb) + GGML_TENSOR_LOCALS(int64_t, nem, m, ne) + GGML_TENSOR_LOCALS(size_t, nbm, m, nb) + GGML_TENSOR_LOCALS(int64_t, ne, dst, ne) + GGML_TENSOR_LOCALS(size_t, nb, dst, nb) + + // input tensor rows must be contiguous + GGML_ASSERT(nbq0 == ggml_type_size(q->type)); + GGML_ASSERT(nbk0 == ggml_type_size(k->type)); + GGML_ASSERT(nbw0 == ggml_type_size(w->type)); + GGML_ASSERT(nbm0 == ggml_type_size(m->type)); + + // dst cannot be transposed or permuted + GGML_ASSERT(nb0 == sizeof(float)); + GGML_ASSERT(nb0 <= nb1); + GGML_ASSERT(nb1 <= nb2); + GGML_ASSERT(nb2 <= nb3); + + const int n_embd = q->ne[0]; + const int n_head = q->ne[1]; + const int n_batch = q->ne[2]; + const int n_stream = q->ne[3]; + const int n_kv = k->ne[2]; + + const float * q_d = (const float *) q->data; + const char * k_d = (const char *) k->data; + const float * w_d = (const float *) w->data; + const half * m_d = (const half *) m->data; + float * dst_d = ( float *) dst->data; + + const int device = ggml_cuda_get_device(); + const int cc = ggml_cuda_info().devices[device].cc; + + if (n_embd == 128 && n_head == 64) { +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && k->type != GGML_TYPE_F32 && k->type != GGML_TYPE_BF16) { + // use wmma kernel + constexpr int K_VECS_PER_BLOCK = 32; + constexpr int WARPS_PER_BLOCK = 8; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 64, k, GGML_TYPE_Q8_0) + GGML_ABORT("fatal error"); + } else { +#else // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + { +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + // use vector kernel + constexpr int K_VECS_PER_WARP = 8; + constexpr int WARPS_PER_BLOCK = 8; + constexpr int K_VECS_PER_BLOCK = K_VECS_PER_WARP * WARPS_PER_BLOCK; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_Q8_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_BF16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 64, k, GGML_TYPE_F32) + GGML_ABORT("fatal error"); + } + } else if (n_embd == 128 && n_head == 32) { +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + if (GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) && k->type != GGML_TYPE_F32 && k->type != GGML_TYPE_BF16) { + // use wmma kernel + constexpr int K_VECS_PER_BLOCK = 32; + constexpr int WARPS_PER_BLOCK = 8; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_wmma, 128, 32, k, GGML_TYPE_Q8_0) + GGML_ABORT("fatal error"); + } else { +#else // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + { +#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + // use vector kernel + constexpr int K_VECS_PER_WARP = 8; + constexpr int WARPS_PER_BLOCK = 8; + constexpr int K_VECS_PER_BLOCK = K_VECS_PER_WARP * WARPS_PER_BLOCK; + + dim3 block(32, WARPS_PER_BLOCK); + int num_kv_blocks = (n_kv + (K_VECS_PER_BLOCK) - 1) / (K_VECS_PER_BLOCK); + dim3 grid(num_kv_blocks, n_batch, n_stream); + + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_F16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q4_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q4_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q5_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q5_1) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_Q8_0) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_BF16) + LIGHTNING_INDEXER_CASE(lightning_indexer_kernel_vec, 128, 32, k, GGML_TYPE_F32) + GGML_ABORT("fatal error"); + } + } else { + GGML_ABORT("fatal error"); + } +} + +bool ggml_cuda_lightning_indexer_supported(int device, const ggml_tensor * dst) { + GGML_UNUSED(device); + + const ggml_tensor * q = dst->src[0]; + const ggml_tensor * k = dst->src[1]; + const ggml_tensor * w = dst->src[2]; // weights + const ggml_tensor * m = dst->src[3]; // mask + + GGML_TENSOR_LOCALS(int64_t, neq, q, ne) + GGML_TENSOR_LOCALS(size_t, nbq, q, nb) + GGML_TENSOR_LOCALS(int64_t, nek, k, ne) + GGML_TENSOR_LOCALS(size_t, nbk, k, nb) + GGML_TENSOR_LOCALS(int64_t, new, w, ne) + GGML_TENSOR_LOCALS(size_t, nbw, w, nb) + GGML_TENSOR_LOCALS(int64_t, nem, m, ne) + GGML_TENSOR_LOCALS(size_t, nbm, m, nb) + GGML_TENSOR_LOCALS(int64_t, ne, dst, ne) + GGML_TENSOR_LOCALS(size_t, nb, dst, nb) + + if (neq0 != 128) { + return false; + } + + if (neq1 != 64 && neq1 != 32) { + return false; + } + + // alignment checks + for (const ggml_tensor * t : {q, k}) { + if (ggml_is_quantized(t->type)) { + continue; + } + for (size_t i = 1; i < GGML_MAX_DIMS; ++i) { + if (t->nb[i] % 16 != 0) { + return false; + } + } + } + + switch(k->type) { + case GGML_TYPE_F32: + case GGML_TYPE_BF16: + case GGML_TYPE_F16: + case GGML_TYPE_Q8_0: + case GGML_TYPE_Q5_1: + case GGML_TYPE_Q5_0: + case GGML_TYPE_Q4_1: + case GGML_TYPE_Q4_0: + return true; + default: + return false; + } +} diff --git a/ggml/src/ggml-cuda/lightning-indexer.cuh b/ggml/src/ggml-cuda/lightning-indexer.cuh new file mode 100644 index 000000000..f2fc95181 --- /dev/null +++ b/ggml/src/ggml-cuda/lightning-indexer.cuh @@ -0,0 +1,4 @@ +#include "common.cuh" + +void ggml_cuda_lightning_indexer(ggml_backend_cuda_context & ctx, ggml_tensor * dst); +bool ggml_cuda_lightning_indexer_supported(int device, const ggml_tensor * dst); From 32beb244f5c2ca91c583be15d4671643b54ba238 Mon Sep 17 00:00:00 2001 From: Aleksander Grygier Date: Wed, 15 Jul 2026 20:31:45 +0200 Subject: [PATCH 12/22] ui: Agentic Content UX improvements (#25450) * feat: Add shimmer text animation for processing state indicators * feat: Redesign CollapsibleContentBlock component with improved UX * feat: Add conditional setting display support with dependsOn field * feat: Add showAgenticTurnStats setting for per-turn statistics * feat: Update ChatMessageAgenticContent with improved UI and new features * feat: Enhance file read tool UI/UX * feat: Refine styling of collapsible content and code preview blocks * feat: add terminal variant to CollapsibleContentBlock * feat: add built-in tools UI registry * feat: extract ChatMessageReasoningBlock and ChatMessageToolCallBlock * refactor: simplify ChatMessageAgenticContent to use extracted blocks * fix: correct markdown content block margin spacing * fix: reorganize SettingsChatFields layout and reset button positioning * fix: use direct map access in agentic store session methods * refactor: remove reasoning preview/throttle system from CollapsibleContentBlock * feat: add auto-scroll to reasoning block and remove showThoughtInProgress * feat: add ChatMessageToolCallDateTime component and support for new tool types * feat: improve auto-scroll reliability in reasoning block with RAF coalescing and MutationObserver * feat: show MCP server favicon for tools without a built-in icon * feat: add search-results parsing utilities and tests * feat: add ChatMessageToolCallSearchResults component * feat: integrate search results rendering into ChatMessageAgenticContent * feat: display tool call input alongside output in ChatMessageToolCallBlock * style: use muted foreground color in reasoning block content * chore: Format * feat: Refine reasoning block layout and make pending thoughts display configurable * feat: Stream tool call code blocks with auto-scroll and handle partial JSON * feat: add streaming permission gate infrastructure * feat: wire permission gate into the agentic loop * fix: bail out on abort and skip already-approved tool calls * fix: clear partial tool calls on abort and savePartialResponse * test: cover partial tool call cleanup end-to-end * refactor: Remove streaming permission gate logic * fix: Correct autoscroll and streaming gates for tool calls and reasoning blocks * refactor: Chat Message Assistant componentization * fix: Show health metadata for disabled MCP servers and promote connections on enable * fix: Inherit global enabled state for missing MCP per-chat overrides * refactor: Cleanup * refactor: Split ChatMessageToolCallBlock into dedicated components * feat: Add live streaming and auto-scroll for tool execution output * feat: Add line numbers and change markers to file edit diffs * chore: Formatting * feat: Add type definitions and utilities for recommended MCP servers * feat: Add recommended MCP servers configuration and storage key * feat: Add McpServerCardCompact component for recommended servers * feat: Add recommended servers section to Add New Server dialog * feat: Update McpServerForm to support authorization requirements * feat: Add select-none classes for text selection prevention * feat: Add recommended MCP server icon assets * refactor: Store dismissed MCP recommendations as a boolean flag * feat: Render tool results as JSON or Markdown based on detected content type * feat: UI improvement * feat: Render search block early and update heading to show execution state * fix: Prevent non-web-search tools from triggering the search UI block * refactor: Cleanup * refactor: Extract hardcoded icon size classes into shared constants * refactor: Extract hardcoded tool result separator into a shared constant * refactor: Tool Calls UI/logic * refactor: Cleanup * refactor: Cleanup * refactor: Cleanup --- tools/ui/eslint.config.js | 7 + tools/ui/src/app.css | 27 ++ .../actions/ActionIconCopyToClipboard.svelte | 3 +- ...hatAttachmentsListItemThumbnailFile.svelte | 5 +- ...hatAttachmentsPreviewCurrentItemPdf.svelte | 9 +- ...hatAttachmentsPreviewThumbnailStrip.svelte | 7 +- .../ChatFormActionAddButton.svelte | 3 +- .../ChatFormActionAddDropdown.svelte | 15 +- .../ChatFormActionAddMcpServersSubmenu.svelte | 9 +- .../ChatFormActionAddReasoningSubmenu.svelte | 9 +- .../ChatFormActionAddSheet.svelte | 47 +- .../ChatFormActionAddToolsSubmenu.svelte | 15 +- .../ChatFormActionRecord.svelte | 5 +- .../ChatFormActions/ChatFormActions.svelte | 5 +- .../ChatFormReasoningEffortSubmenu.svelte | 8 +- .../ChatMessage/ChatMessage.svelte | 1 - .../ChatMessageAssistant.svelte | 201 +-------- .../ChatMessageAssistantModel.svelte | 46 ++ .../ChatMessageAssistantProcessingInfo.svelte | 25 ++ .../ChatMessageAssistantRawOutput.svelte | 33 ++ .../ChatMessageAssistantStatistics.svelte | 40 ++ .../ChatMessageToolCallBlock.svelte | 66 +++ .../ChatMessageToolCallBlockDefault.svelte | 124 ++++++ .../ChatMessageToolCallBlockEditFile.svelte | 166 +++++++ ...essageToolCallBlockExecShellCommand.svelte | 293 ++++++++++++ ...tMessageToolCallBlockFileGlobSearch.svelte | 61 +++ ...ChatMessageToolCallBlockGetDatetime.svelte | 57 +++ .../ChatMessageToolCallBlockGrepSearch.svelte | 67 +++ .../ChatMessageToolCallBlockReadFile.svelte | 44 ++ ...atMessageToolCallBlockRunJavascript.svelte | 69 +++ ...atMessageToolCallBlockSearchResults.svelte | 167 +++++++ .../ChatMessageToolCallBlockWriteFile.svelte | 55 +++ .../ChatMessageToolCall/ToolCallBlock.svelte | 129 ++++++ .../ChatMessageToolCall/parsers/_shared.ts | 49 +++ .../ChatMessageToolCall/parsers/edit-file.ts | 71 +++ .../parsers/exec-shell-command.ts | 23 + .../parsers/file-glob-search.ts | 58 +++ .../parsers/grep-search.ts | 108 +++++ .../ChatMessageToolCall/parsers/read-file.ts | 52 +++ .../parsers/run-javascript.ts | 56 +++ .../ChatMessageToolCall/parsers/write-file.ts | 54 +++ .../ChatMessageActionCard.svelte | 3 +- .../ChatMessageAgenticContent.svelte | 254 +++-------- .../ChatMessageReasoningBlock.svelte | 151 +++++++ .../ChatScreenActionScrollDown.svelte | 3 +- .../ChatScreen/ChatScreenServerError.svelte | 5 +- tools/ui/src/lib/components/app/chat/index.ts | 4 + .../content/CollapsibleContentBlock.svelte | 148 +++---- .../content/CollapsibleTerminalBlock.svelte | 97 ++++ .../MarkdownContent/markdown-content.css | 10 +- .../app/content/MermaidPreviewControls.svelte | 9 +- .../app/content/SyntaxHighlightedCode.svelte | 109 +++-- .../src/lib/components/app/content/index.ts | 17 +- .../dialogs/DialogMcpResourcesBrowser.svelte | 11 +- .../app/dialogs/DialogMcpServerAddNew.svelte | 148 ++++++- .../dialogs/DialogModelNotAvailable.svelte | 3 +- .../components/app/forms/KeyValuePairs.svelte | 5 +- .../components/app/forms/SearchInput.svelte | 5 +- .../app/mcp/McpActiveServersAvatars.svelte | 5 +- .../app/mcp/McpResourcePreview.svelte | 3 +- .../McpResourcesBrowserHeader.svelte | 5 +- .../McpResourcesBrowserServerItem.svelte | 5 +- .../mcp/McpServerCard/McpServerCard.svelte | 5 +- .../McpServerCard/McpServerCardCompact.svelte | 44 ++ .../components/app/mcp/McpServerForm.svelte | 46 +- tools/ui/src/lib/components/app/mcp/index.ts | 11 + .../app/misc/HorizontalScrollCarousel.svelte | 5 +- .../app/models/ModelsSelectorOption.svelte | 3 +- .../SidebarNavigationActions.svelte | 5 +- .../SidebarNavigationConversationItem.svelte | 3 +- .../app/server/ServerErrorSplash.svelte | 19 +- .../components/app/server/ServerStatus.svelte | 3 +- .../settings/SettingsChat/SettingsChat.svelte | 4 +- .../SettingsChat/SettingsChatFields.svelte | 387 ++++++++-------- .../SettingsChatImportExportSection.svelte | 3 +- .../SettingsChat/SettingsChatToolsTab.svelte | 7 +- .../SettingsChatDesktopSidebar.svelte | 7 +- .../settings/SettingsChatMobileHeader.svelte | 11 +- .../app/settings/SettingsMcpServers.svelte | 5 +- tools/ui/src/lib/constants/agentic.ts | 26 +- tools/ui/src/lib/constants/auto-scroll.ts | 14 + tools/ui/src/lib/constants/built-in-tools.ts | 59 +++ tools/ui/src/lib/constants/code.ts | 16 + tools/ui/src/lib/constants/css-classes.ts | 6 + tools/ui/src/lib/constants/formatters.ts | 27 -- tools/ui/src/lib/constants/index.ts | 2 + tools/ui/src/lib/constants/markdown.ts | 1 + tools/ui/src/lib/constants/mcp.ts | 6 + .../lib/constants/recommended-mcp-servers.ts | 38 ++ tools/ui/src/lib/constants/sandbox.ts | 4 +- tools/ui/src/lib/constants/settings-keys.ts | 1 + .../ui/src/lib/constants/settings-registry.ts | 17 +- tools/ui/src/lib/constants/storage.ts | 1 + tools/ui/src/lib/constants/ui.ts | 3 + tools/ui/src/lib/constants/url.ts | 3 + tools/ui/src/lib/enums/agentic.enums.ts | 18 + tools/ui/src/lib/enums/index.ts | 10 +- tools/ui/src/lib/enums/tools.enums.ts | 19 + tools/ui/src/lib/hooks/use-throttle.svelte.ts | 32 -- .../ui/src/lib/services/migration.service.ts | 2 +- tools/ui/src/lib/services/sandbox.service.ts | 6 +- tools/ui/src/lib/services/tools.service.ts | 90 ++++ tools/ui/src/lib/stores/agentic.svelte.ts | 110 ++++- tools/ui/src/lib/stores/chat.svelte.ts | 55 ++- .../ui/src/lib/stores/mcp-resources.svelte.ts | 4 +- tools/ui/src/lib/stores/mcp.svelte.ts | 140 +++--- tools/ui/src/lib/types/agentic.d.ts | 14 + tools/ui/src/lib/types/chat.d.ts | 5 + tools/ui/src/lib/types/index.ts | 1 + tools/ui/src/lib/types/mcp.d.ts | 20 + tools/ui/src/lib/types/settings.d.ts | 2 + tools/ui/src/lib/utils/agentic.ts | 185 +++++++- tools/ui/src/lib/utils/api-headers.ts | 9 +- tools/ui/src/lib/utils/api-key-validation.ts | 3 +- tools/ui/src/lib/utils/code.ts | 27 +- tools/ui/src/lib/utils/compute-line-diff.ts | 105 +++++ tools/ui/src/lib/utils/formatters.ts | 36 +- tools/ui/src/lib/utils/index.ts | 56 ++- tools/ui/src/lib/utils/model-names.ts | 4 +- .../src/lib/utils/parse-exec-shell-error.ts | 19 + .../src/lib/utils/parse-exec-shell-status.ts | 47 ++ .../src/lib/utils/parse-partial-json-args.ts | 83 ++++ tools/ui/src/lib/utils/search-results.ts | 229 ++++++++++ tools/ui/src/lib/utils/sse.ts | 71 +++ tools/ui/src/lib/utils/text.ts | 4 +- tools/ui/src/lib/utils/tool-call-meta.ts | 27 ++ tools/ui/src/lib/utils/url.ts | 42 +- tools/ui/src/routes/+layout.svelte | 19 +- tools/ui/static/recommended-mcp/context7.png | Bin 0 -> 1489 bytes tools/ui/static/recommended-mcp/exa.ico | Bin 0 -> 15154 bytes .../ui/static/recommended-mcp/github-dark.png | Bin 0 -> 584 bytes .../static/recommended-mcp/github-light.png | Bin 0 -> 958 bytes .../ui/static/recommended-mcp/huggingface.ico | Bin 0 -> 205556 bytes tools/ui/tests/unit/agentic-sections.test.ts | 23 + .../tests/unit/assistant-raw-output.test.ts | 77 ++++ .../tests/unit/classify-tool-result.test.ts | 130 ++++++ tools/ui/tests/unit/code.test.ts | 82 ++++ tools/ui/tests/unit/compute-line-diff.test.ts | 136 ++++++ .../tests/unit/mcp-override-fallback.test.ts | 143 ++++++ .../unit/parse-exec-shell-status.test.ts | 67 +++ .../unit/partial-tool-call-cleanup.test.ts | 95 ++++ .../tests/unit/search-results-fixture.test.ts | 44 ++ tools/ui/tests/unit/search-results.test.ts | 118 +++++ tools/ui/tests/unit/sse.test.ts | 77 ++++ tools/ui/tests/unit/tool-call-meta.test.ts | 30 ++ tools/ui/tests/unit/tool-calls.test.ts | 416 ++++++++++++++++++ 146 files changed, 5960 insertions(+), 1053 deletions(-) create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantModel.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockEditFile.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockFileGlobSearch.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGetDatetime.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGrepSearch.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadFile.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockRunJavascript.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockWriteFile.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts create mode 100644 tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageReasoningBlock.svelte create mode 100644 tools/ui/src/lib/components/app/content/CollapsibleTerminalBlock.svelte create mode 100644 tools/ui/src/lib/components/app/mcp/McpServerCard/McpServerCardCompact.svelte create mode 100644 tools/ui/src/lib/constants/built-in-tools.ts create mode 100644 tools/ui/src/lib/constants/recommended-mcp-servers.ts delete mode 100644 tools/ui/src/lib/hooks/use-throttle.svelte.ts create mode 100644 tools/ui/src/lib/utils/compute-line-diff.ts create mode 100644 tools/ui/src/lib/utils/parse-exec-shell-error.ts create mode 100644 tools/ui/src/lib/utils/parse-exec-shell-status.ts create mode 100644 tools/ui/src/lib/utils/parse-partial-json-args.ts create mode 100644 tools/ui/src/lib/utils/search-results.ts create mode 100644 tools/ui/src/lib/utils/sse.ts create mode 100644 tools/ui/src/lib/utils/tool-call-meta.ts create mode 100644 tools/ui/static/recommended-mcp/context7.png create mode 100644 tools/ui/static/recommended-mcp/exa.ico create mode 100644 tools/ui/static/recommended-mcp/github-dark.png create mode 100644 tools/ui/static/recommended-mcp/github-light.png create mode 100644 tools/ui/static/recommended-mcp/huggingface.ico create mode 100644 tools/ui/tests/unit/assistant-raw-output.test.ts create mode 100644 tools/ui/tests/unit/classify-tool-result.test.ts create mode 100644 tools/ui/tests/unit/code.test.ts create mode 100644 tools/ui/tests/unit/compute-line-diff.test.ts create mode 100644 tools/ui/tests/unit/mcp-override-fallback.test.ts create mode 100644 tools/ui/tests/unit/parse-exec-shell-status.test.ts create mode 100644 tools/ui/tests/unit/partial-tool-call-cleanup.test.ts create mode 100644 tools/ui/tests/unit/search-results-fixture.test.ts create mode 100644 tools/ui/tests/unit/search-results.test.ts create mode 100644 tools/ui/tests/unit/sse.test.ts create mode 100644 tools/ui/tests/unit/tool-call-meta.test.ts create mode 100644 tools/ui/tests/unit/tool-calls.test.ts diff --git a/tools/ui/eslint.config.js b/tools/ui/eslint.config.js index b90376b6c..54679a05b 100644 --- a/tools/ui/eslint.config.js +++ b/tools/ui/eslint.config.js @@ -29,6 +29,13 @@ export default ts.config( // This app uses hash-based routing (#/) where resolve() from $app/paths does not apply 'svelte/no-navigation-without-resolve': 'off', + // Snippet bodies often ignore one or more of the parent's params + // (e.g. `{#snippet children(_meta, ctx)}` when only ctx is read). + '@typescript-eslint/no-unused-vars': [ + 'error', + { argsIgnorePattern: '^_', varsIgnorePattern: '^_' } + ], + // Enforce empty line at end of file 'eol-last': 'error' } diff --git a/tools/ui/src/app.css b/tools/ui/src/app.css index 8c4056477..f9b544beb 100644 --- a/tools/ui/src/app.css +++ b/tools/ui/src/app.css @@ -193,6 +193,33 @@ -ms-overflow-style: none; scrollbar-width: none; } + + .shimmer-text { + background: linear-gradient( + 90deg, + var(--muted-foreground), + var(--foreground), + var(--muted-foreground) + ); + background-size: 200% 100%; + background-clip: text; + -webkit-background-clip: text; + -webkit-text-fill-color: transparent; + font-weight: 500; + animation: shimmer 1s linear infinite; + } + + @keyframes shimmer { + to { + background-position: -200% 0; + } + } + + @media (prefers-reduced-motion: reduce) { + .shimmer-text { + animation: none; + } + } } .mermaidTooltip { diff --git a/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte b/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte index 999f0cba9..9b7b370ad 100644 --- a/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte +++ b/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte @@ -1,4 +1,5 @@ + +{#if isRouter} + { + const status = modelsStore.getModelStatus(modelId); + + if (status !== ServerModelStatus.LOADED) { + pendingModel = modelId; + + try { + await modelsStore.loadModel(modelId); + } finally { + pendingModel = null; + } + } + + onRegenerate(modelName); + return true; + }} + /> +{:else} + +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte new file mode 100644 index 000000000..356512ecb --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantProcessingInfo.svelte @@ -0,0 +1,25 @@ + + +
+
+ + {modelLoadingText ?? + processingState.getPromptProgressText() ?? + processingState.getProcessingMessage() ?? + 'Processing...'} + +
+
diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte new file mode 100644 index 000000000..30ce16be9 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantRawOutput.svelte @@ -0,0 +1,33 @@ + + +
{rawOutputContent || ''}
+ + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte new file mode 100644 index 000000000..4cc4080c3 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistantStatistics.svelte @@ -0,0 +1,40 @@ + + +{#if showMessageStats && message.timings && message.timings.predicted_n && message.timings.predicted_ms} + {@const agentic = message.timings.agentic} + +{:else if isLoading && showMessageStats} + {@const liveStats = processingState.getLiveProcessingStats()} + {@const genStats = processingState.getLiveGenerationStats()} + + {#if genStats} + + {/if} +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte new file mode 100644 index 000000000..b1daedfc8 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlock.svelte @@ -0,0 +1,66 @@ + + +{#if isSearchCall} + +{:else if section.toolName === BuiltInTool.GET_DATETIME} + +{:else if section.toolName === BuiltInTool.READ_FILE} + +{:else if section.toolName === BuiltInTool.EDIT_FILE} + +{:else if section.toolName === BuiltInTool.WRITE_FILE} + +{:else if section.toolName === BuiltInTool.EXEC_SHELL_COMMAND} + +{:else if section.toolName === BuiltInTool.FILE_GLOB_SEARCH} + +{:else if section.toolName === BuiltInTool.GREP_SEARCH} + +{:else if section.toolName === BuiltInTool.RUN_JAVASCRIPT} + +{:else} + +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte new file mode 100644 index 000000000..acf2de12a --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockDefault.svelte @@ -0,0 +1,124 @@ + + + + {#snippet children(_meta, ctx)} + {#if ctx.isStreamingCall} +
+ Input + {#if ctx.isStreaming} + + {/if} +
+ {#if section.toolArgs} + + {:else if ctx.isStreaming} +
+ Receiving arguments... +
+ {:else} +
+ Response was truncated +
+ {/if} + {:else} + {@const showInput = Boolean(section.toolArgs)} + {#if showInput} +
+ Input +
+ + {/if} +
+ Output + {#if ctx.isPending} + + {/if} +
+ {#if ctx.isPending} +
+ Waiting for result... +
+ {:else if section.toolResult} + {#if outputKind === ToolResultKind.JSON} + + {:else if outputKind === ToolResultKind.MARKDOWN} + + {:else} +
+ {#each parsedLines as line, i (i)} +
+ {line.text} +
+ {#if line.image} + {line.image.name} + {/if} + {/each} +
+ {/if} + {:else} +
No output
+ {/if} + {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockEditFile.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockEditFile.svelte new file mode 100644 index 000000000..6f30060f5 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockEditFile.svelte @@ -0,0 +1,166 @@ + + + + {#snippet titleSnippet()} + Edit file + {editFileMeta?.filePath} + {#if editFileMeta?.errorMessage} + (failed) + {/if} + {/snippet} + + {#snippet children(meta, _ctx)} + {#if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta && meta.edits.length > 0} + {#each editDiffs as diffLines, ei (ei)} +
+
+ Edit {ei + 1} of {meta.edits.length} +
+
+
+ {#each diffLines as line, li (li)} +
+ {line.oldLine ?? ''} + {prefixFor(line.kind)} + {line.newLine ?? ''} + {line.text || ' '} +
+ {/each} +
+
+
+ {/each} +
+ {#if meta.resultMessage} + {meta.resultMessage}{meta.editsApplied != null ? RESULT_STAT_SEPARATOR : ''}{/if} + {#if meta.editsApplied != null} + {meta.editsApplied} + {meta.editsApplied === 1 ? 'edit' : 'edits'} applied + {/if} +
+ {:else} +
No edits
+ {/if} + {/snippet} + + + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte new file mode 100644 index 000000000..5de801d39 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte @@ -0,0 +1,293 @@ + + +{#snippet execShellTitle()} + {#if highlightedCommandHtml} + {@html highlightedCommandHtml} + {:else} + {execShellMeta?.command} + {/if} +{/snippet} + + + {#snippet titleSnippet()} + {@render execShellTitle()} + {/snippet} + + {#snippet children(_meta, ctx)} + {#if ctx.isPending} +
+ + Running... +
+ {:else if execShellError} +
+ + {execShellError} +
+ {:else if section.toolResult} +
+ {#each outputLines as line, i (i)} +
{line.text}
+ {#if line.image} + {line.image.name} + {/if} + {/each} + + {#if isExitCodeFinalLine && execShellExitStatus} +
+ {#if execShellExitStatus.timedOut} + + timed out + · + exit {execShellExitStatus.code} + {:else if execShellExitStatus.code === 0} + + exit 0 + {:else} + + exit {execShellExitStatus.code} + {/if} +
+ {/if} +
+ {/if} + {/snippet} + + + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockFileGlobSearch.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockFileGlobSearch.svelte new file mode 100644 index 000000000..ad082039f --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockFileGlobSearch.svelte @@ -0,0 +1,61 @@ + + + + {#snippet titleSnippet()} + {#if fileGlobMeta} + {fileGlobMeta.include === '**' ? 'List files' : 'Search files'}  + {#if fileGlobMeta.include !== '**'} + {fileGlobMeta.include} + {/if} +  in  + {fileGlobMeta.path} + {/if} + {/snippet} + + {#snippet children(meta, ctx)} + {#if ctx.isPending} +
+ Searching... +
+ {:else if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta && meta.matches.length > 0} +
+ {#each meta.matches as match, i (i)} +
{match}
+ {/each} +
+
+ Total matches: {meta.totalMatches ?? meta.matches.length} +
+ {:else} +
No matches
+
+ Total matches: {meta?.totalMatches ?? 0} +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGetDatetime.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGetDatetime.svelte new file mode 100644 index 000000000..e0c701dea --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGetDatetime.svelte @@ -0,0 +1,57 @@ + + +
+ + {#if showSpinner} + Current time + + {:else if dateMeta.errorMessage} + Current time  + - {dateMeta.errorMessage} + {:else if dateMeta.dateString} + Current time is  + {dateMeta.dateString} + {:else} + Current time + {/if} +
diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGrepSearch.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGrepSearch.svelte new file mode 100644 index 000000000..afb06fef7 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockGrepSearch.svelte @@ -0,0 +1,67 @@ + + + + {#snippet titleSnippet()} + {#if grepMeta} + Search for  + {grepMeta.pattern} +  in  + {grepMeta.path} + {/if} + {/snippet} + + {#snippet children(meta, ctx)} + {#if ctx.isPending} +
+ Searching... +
+ {:else if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta && meta.matches.length > 0} +
+ {#each meta.matches as match, mi (mi)} +
+ {match.file} + {#if meta.showLineNumbers && match.line != null} + :{match.line} + {/if} + : + {match.content} +
+ {/each} +
+
+ Total matches: {meta.totalMatches ?? meta.matches.length} + {#if meta.showLineNumbers} +  (with line numbers) + {/if} +
+ {:else} +
No matches
+
+ Total matches: {meta?.totalMatches ?? 0} +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadFile.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadFile.svelte new file mode 100644 index 000000000..a99ff9cee --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockReadFile.svelte @@ -0,0 +1,44 @@ + + + + {#snippet titleSnippet()} + Read file + {readFileMeta?.fileName} + {#if readFileMeta?.lineRange} +  (lines {readFileMeta.lineRange.start}-{readFileMeta.lineRange.end}) + {/if} + {/snippet} + + {#snippet children(_meta, _ctx)} + {#if section.toolResult} + + {:else} +
+ Waiting for file content... +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockRunJavascript.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockRunJavascript.svelte new file mode 100644 index 000000000..707d83d73 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockRunJavascript.svelte @@ -0,0 +1,69 @@ + + + + {#snippet children(meta, ctx)} + {#if ctx.isPending} +
Running...
+ {:else if meta?.errorMessage} +
+ + {meta.errorMessage} +
+
+ +
+ {:else if meta} + +
+ + Console + {#if meta.timeoutMs != null} + · timeout {meta.timeoutMs} ms + {/if} +
+ {#if section.toolResult} +
+ +
+ {:else} +
No output
+ {/if} + {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte new file mode 100644 index 000000000..e4b4adf15 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockSearchResults.svelte @@ -0,0 +1,167 @@ + + +{#snippet pill(result: SearchResult)} + {@const faviconUrl = faviconForUrl(result.url)} + {@const safeUrl = sanitizeExternalUrl(result.url)} + {@const showHoverCard = safeUrl !== null && hasDetails(result)} + {#if safeUrl} + + + {#if faviconUrl} + + {:else} + + {/if} + {result.title} + + {#if showHoverCard} + {@const publishDate = formatPublishDate(result.published)} + {@const host = hostFor(safeUrl)} + +
+ {result.title} + {#if publishDate || result.author} +
+ {#if publishDate} + {publishDate} + {/if} + {#if publishDate && result.author} + · + {/if} + {#if result.author} + {result.author} + {/if} +
+ {/if} + {#if result.highlights} +

+ {result.highlights} +

+ {/if} + {#if host} +
{host}
+ {/if} +
+ + {/if} + + {/if} +{/snippet} + + + {#if results.length > 0} +
+ {#each results as result (result.url)} + {@render pill(result)} + {/each} +
+ {:else if showSpinner} +
+ + Searching... +
+ {:else} +
No results
+ {/if} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockWriteFile.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockWriteFile.svelte new file mode 100644 index 000000000..eda067662 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockWriteFile.svelte @@ -0,0 +1,55 @@ + + + + {#snippet titleSnippet()} + Write file + {writeFileMeta?.filePath} + {#if writeFileMeta?.errorMessage} + (failed) + {/if} + {/snippet} + + {#snippet children(meta, ctx)} + {#if meta?.errorMessage} +
+ + {meta.errorMessage} +
+ {:else if meta} + +
+ {#if meta.resultMessage} + {meta.resultMessage}{meta.bytesWritten != null ? RESULT_STAT_SEPARATOR : ''}{/if} + {#if meta.bytesWritten != null} + {meta.bytesWritten} + bytes + {/if} +
+ {/if} + {/snippet} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte new file mode 100644 index 000000000..a17a74e16 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte @@ -0,0 +1,129 @@ + + + + {@render children(meta, { + isStreaming, + isPending, + isStreamingCall, + isCodeStreaming + })} + diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts new file mode 100644 index 000000000..6114f17b5 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/_shared.ts @@ -0,0 +1,49 @@ +// Helpers shared by the per-tool meta parsers under +// `src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/`. +// Each tool needs the same first three steps (tool-name check, +// args-present check, JSON parse) - keeping them here lets each parser +// stay focused on its own format quirks. + +import { BuiltInTool } from '$lib/enums'; +import { parsePartialJsonArgs } from '$lib/utils/parse-partial-json-args'; +import type { AgenticSection } from '$lib/utils/agentic'; + +/** + * Strict (final-state) JSON parser for a tool-args blob. Mirrors the + * behaviour the per-tool components used before extraction: an + * invalid JSON blob, a JSON array, or a JSON primitive all map to + * `null` so callers don't have to guard against surprise shapes. + */ +function parseFinalToolArgs(blob: string): Record | null { + try { + const parsed: unknown = JSON.parse(blob); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + return parsed as Record; + } + return null; + } catch { + return null; + } +} + +/** + * Parse a section's toolArgs against an expected tool name. Returns + * `null` when: + * - the section's toolName doesn't match (component isn't for this + * tool); + * - the section has no args yet (call hasn't started streaming); + * - or the args blob can't be parsed. + * + * Pass `{ partial: true }` for tools that need to render incrementally + * as each token lands (read_file, edit_file, write_file). + */ +export function parseToolArgs( + expected: BuiltInTool, + section: AgenticSection, + options: { partial?: boolean } = {} +): Record | null { + if (section.toolName !== expected || !section.toolArgs) return null; + return options.partial + ? parsePartialJsonArgs(section.toolArgs) + : parseFinalToolArgs(section.toolArgs); +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts new file mode 100644 index 000000000..4bff25bb5 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/edit-file.ts @@ -0,0 +1,71 @@ +// Meta parser for `edit_file` tool calls. Reads the file path and the +// array of edits from the streamed args (partial JSON for incremental +// rendering), plus the result blob for `result` / `edits_applied` / +// `error` fields. + +import { BuiltInTool } from '$lib/enums'; +import { FILE_PATH_SEPARATOR_REGEX } from '$lib/constants'; +import { tryParseToolResultObject, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type EditFileEdit = { + oldText: string; + newText: string; +}; + +export type EditFileMeta = { + fileName: string; + filePath: string; + edits: EditFileEdit[]; + resultMessage?: string; + editsApplied?: number; + errorMessage?: string; +}; + +export function parseEditFileMeta(section: AgenticSection): EditFileMeta | null { + const args = parseToolArgs(BuiltInTool.EDIT_FILE, section, { partial: true }); + if (!args) return null; + + const rawPath = args.path ?? args.file_path ?? args.filePath; + if (typeof rawPath !== 'string' || !rawPath) return null; + + const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath; + + // Filter the streamed edits array strictly: each entry must be an + // object with a non-empty `old_text`. Edits without an old_text + // would diff against empty and render as a full re-write. + const rawEdits = Array.isArray(args.edits) ? args.edits : []; + const edits: EditFileEdit[] = []; + for (const e of rawEdits) { + if (!e || typeof e !== 'object' || Array.isArray(e)) continue; + const obj = e as Record; + const oldText = typeof obj.old_text === 'string' ? obj.old_text : ''; + if (!oldText) continue; + const newText = typeof obj.new_text === 'string' ? obj.new_text : ''; + edits.push({ oldText, newText }); + } + + const resultObj = tryParseToolResultObject(section.toolResult); + let resultMessage: string | undefined; + let editsApplied: number | undefined; + let errorMessage: string | undefined; + if (typeof resultObj?.error === 'string') { + errorMessage = resultObj.error; + } else if (resultObj) { + if (typeof resultObj.result === 'string') { + resultMessage = resultObj.result; + } + if (Number.isFinite(Number(resultObj.edits_applied))) { + editsApplied = Number(resultObj.edits_applied); + } + } + + return { + fileName, + filePath: rawPath, + edits, + resultMessage, + editsApplied, + errorMessage + }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts new file mode 100644 index 000000000..e8adbd18b --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/exec-shell-command.ts @@ -0,0 +1,23 @@ +// Meta parser for `exec_shell_command` tool calls. Surfaces the +// command text from args `command` / `cmd` / `shell_command` aliases. +// The exit-status and error parsing live in their own utilities +// (`parse-exec-shell-status.ts` / `parse-exec-shell-error.ts`) - this +// file only deals with what's strictly about *calling* the tool, since +// the error / exit status elide from call-section to result-section. + +import { BuiltInTool } from '$lib/enums'; +import type { AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type ExecShellCommandMeta = { + command: string; +}; + +export function parseExecShellCommandMeta(section: AgenticSection): ExecShellCommandMeta | null { + const args = parseToolArgs(BuiltInTool.EXEC_SHELL_COMMAND, section); + if (!args) return null; + + const commandRaw = args.command ?? args.cmd ?? args.shell_command; + if (typeof commandRaw !== 'string' || !commandRaw) return null; + return { command: commandRaw }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts new file mode 100644 index 000000000..1ad92b74c --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/file-glob-search.ts @@ -0,0 +1,58 @@ +// Meta parser for `file_glob_search` tool calls. Reads the path, +// include pattern, and optional exclude from the args (strict parsing) +// and the matches from the result blob. Like grep_search, the result +// parser keeps the original raw-text fallback for MCP servers that +// emit unparseable output. + +import { BuiltInTool } from '$lib/enums'; +import { splitSearchSummaryList, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type FileGlobSearchMeta = { + path: string; + include: string; + exclude?: string; + matches: string[]; + totalMatches?: number; + errorMessage?: string; +}; + +export function parseFileGlobSearchMeta(section: AgenticSection): FileGlobSearchMeta | null { + const args = parseToolArgs(BuiltInTool.FILE_GLOB_SEARCH, section); + if (!args) return null; + + const path = typeof args.path === 'string' ? args.path : ''; + const include = typeof args.include === 'string' && args.include ? args.include : '**'; + const exclude = typeof args.exclude === 'string' && args.exclude ? args.exclude : undefined; + if (!path) return null; + + let matches: string[] = []; + let totalMatches: number | undefined; + let errorMessage: string | undefined; + + const toolResultString = section.toolResult; + if (toolResultString) { + try { + const parsed: unknown = JSON.parse(toolResultString); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + const obj = parsed as Record; + if (typeof obj.error === 'string') { + errorMessage = obj.error; + } else if (typeof obj.plain_text_response === 'string') { + const split = splitSearchSummaryList(obj.plain_text_response, (total) => { + totalMatches = total; + }); + matches = split.lines; + } + } + } catch { + // See grep-search.ts: same fallback used there. + const split = splitSearchSummaryList(toolResultString, (total) => { + totalMatches = total; + }); + matches = split.lines; + } + } + + return { path, include, exclude, matches, totalMatches, errorMessage }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts new file mode 100644 index 000000000..0e606e193 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/grep-search.ts @@ -0,0 +1,108 @@ +// Meta parser for `grep_search` tool calls. Reads the path/pattern +// triplet from args (strict parsing - we wait for the args to +// complete) and the matches from the result blob. The result parser +// keeps the original "scan result as raw text on JSON.parse failure" +// fallback so MCP servers that return unparseable output still get +// surfaced. + +import { BuiltInTool } from '$lib/enums'; +import { splitSearchSummaryList, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type GrepSearchMatch = { + file: string; + line?: number; + content: string; +}; + +export type GrepSearchMeta = { + path: string; + pattern: string; + include: string; + exclude?: string; + showLineNumbers: boolean; + matches: GrepSearchMatch[]; + totalMatches?: number; + errorMessage?: string; +}; + +export function parseGrepSearchMeta(section: AgenticSection): GrepSearchMeta | null { + const args = parseToolArgs(BuiltInTool.GREP_SEARCH, section); + if (!args) return null; + + const path = typeof args.path === 'string' ? args.path : ''; + const pattern = typeof args.pattern === 'string' ? args.pattern : ''; + if (!path || !pattern) return null; + + const include = typeof args.include === 'string' && args.include ? args.include : '**'; + const exclude = typeof args.exclude === 'string' && args.exclude ? args.exclude : undefined; + const showLineNumbers = args.return_line_numbers === true; + + let matches: GrepSearchMatch[] = []; + let totalMatches: number | undefined; + let errorMessage: string | undefined; + + const toolResultString = section.toolResult; + if (toolResultString) { + try { + const parsed: unknown = JSON.parse(toolResultString); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + const obj = parsed as Record; + if (typeof obj.error === 'string') { + errorMessage = obj.error; + } else if (typeof obj.plain_text_response === 'string') { + const split = splitSearchSummaryList(obj.plain_text_response, (total) => { + totalMatches = total; + }); + matches = split.lines.map((line) => parseGrepLine(line, showLineNumbers)); + } + } + } catch { + // Result wasn't JSON: keep behaviour for MCP servers that + // emit raw text and treat each line as a `:` + // (or `::`) match. + const split = splitSearchSummaryList(toolResultString, (total) => { + totalMatches = total; + }); + matches = split.lines.map((line) => parseGrepLine(line, showLineNumbers)); + } + } + + return { + path, + pattern, + include, + exclude, + showLineNumbers, + matches, + totalMatches, + errorMessage + }; +} + +function parseGrepLine(line: string, showLineNumbers: boolean): GrepSearchMatch { + // Server output: + // : when return_line_numbers=false + // :: when return_line_numbers=true + const firstColon = line.indexOf(':'); + if (firstColon === -1) { + return { file: line, content: '' }; + } + const file = line.slice(0, firstColon); + const tail = line.slice(firstColon + 1); + + if (!showLineNumbers) { + return { file, content: tail }; + } + + const secondColon = tail.indexOf(':'); + if (secondColon === -1) { + return { file, content: tail }; + } + const lineNum = parseInt(tail.slice(0, secondColon), 10); + return { + file, + line: Number.isFinite(lineNum) ? lineNum : undefined, + content: tail.slice(secondColon + 1) + }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts new file mode 100644 index 000000000..d37dcf501 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/read-file.ts @@ -0,0 +1,52 @@ +// Meta parser for `read_file` tool calls. Reads the file path and an +// optional line range (either `start_line`+`end_line` or +// `start_line`+`line_count`). Args are parsed partially so a header +// can render incrementally as the file path streams in. + +import { BuiltInTool } from '$lib/enums'; +import { + DEFAULT_LANGUAGE, + FILE_PATH_SEPARATOR_REGEX, + TEXT_LANGUAGE_PREFIX_REGEX +} from '$lib/constants'; +import { getFileTypeByExtension, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type ReadFileMeta = { + fileName: string; + lineRange: { start: number; end: number } | null; + language: string; +}; + +export function parseReadFileMeta(section: AgenticSection): ReadFileMeta | null { + const args = parseToolArgs(BuiltInTool.READ_FILE, section, { partial: true }); + if (!args) return null; + + const rawPath = args.path ?? args.file_path ?? args.filePath; + if (typeof rawPath !== 'string' || !rawPath) return null; + + const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath; + + // Models emit range arguments under several aliases. Accept all to + // stay forgiving across prompt variations. + const startRaw = args.start_line ?? args.line_start ?? args.startLine ?? args.from_line; + const endRaw = args.end_line ?? args.line_end ?? args.endLine ?? args.to_line; + const countRaw = args.line_count ?? args.count ?? args.num_lines; + + let lineRange: { start: number; end: number } | null = null; + const sNum = Number(startRaw); + const eNum = Number(endRaw); + if (startRaw != null && endRaw != null && Number.isFinite(sNum) && Number.isFinite(eNum)) { + lineRange = { start: sNum, end: eNum }; + } else if (startRaw != null && countRaw != null) { + const cNum = Number(countRaw); + if (Number.isFinite(sNum) && Number.isFinite(cNum)) { + lineRange = { start: sNum, end: sNum + cNum - 1 }; + } + } + + const fileType = getFileTypeByExtension(fileName); + const language = fileType ? fileType.replace(TEXT_LANGUAGE_PREFIX_REGEX, '') : DEFAULT_LANGUAGE; + + return { fileName, lineRange, language }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts new file mode 100644 index 000000000..9bcba8f03 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/run-javascript.ts @@ -0,0 +1,56 @@ +// Meta parser for `run_javascript` tool calls. Reads the JS code and +// optional timeout from args (strict parsing) and surfaces any error +// from the result blob. SandboxService.formatReply emits a JSON object +// containing an `error` field on failure, but a partial/non-JSON +// failure renders as a flat line beginning with `Error:`. Both shapes +// are handled. + +import { BuiltInTool } from '$lib/enums'; +import type { AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type RunJavascriptMeta = { + code: string; + timeoutMs?: number; + errorMessage?: string; +}; + +export function parseRunJavascriptMeta(section: AgenticSection): RunJavascriptMeta | null { + const args = parseToolArgs(BuiltInTool.RUN_JAVASCRIPT, section); + if (!args) return null; + + const code = typeof args.code === 'string' ? args.code : ''; + if (!code) return null; + + const timeoutRaw = Number(args.timeout_ms); + const timeoutMs = Number.isFinite(timeoutRaw) && timeoutRaw > 0 ? timeoutRaw : undefined; + + let errorMessage: string | undefined; + const toolResultString = section.toolResult; + if (toolResultString) { + // Branches matter here: a JSON object can carry `error`, but a + // JSON array always represents successful output (sandbox returns + // the array of values). Only when the result isn't a JSON object + // do we scan raw lines for the `Error:` prefix. + let parsedObject: Record | null = null; + try { + const parsed: unknown = JSON.parse(toolResultString); + if (parsed && typeof parsed === 'object' && !Array.isArray(parsed)) { + parsedObject = parsed as Record; + } + } catch { + parsedObject = null; + } + if (typeof parsedObject?.error === 'string') { + errorMessage = parsedObject.error; + } else if (!parsedObject) { + const errorLine = toolResultString + .split('\n') + .map((line) => line.trim()) + .find((line) => line.startsWith('Error:')); + if (errorLine) errorMessage = errorLine.slice('Error:'.length).trim(); + } + } + + return { code, timeoutMs, errorMessage }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts new file mode 100644 index 000000000..95edc3d95 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/parsers/write-file.ts @@ -0,0 +1,54 @@ +// Meta parser for `write_file` tool calls. Reads the path/content from +// the streamed args (partial JSON so we can render before the call +// finishes) and surfaces `bytes`, `result`, and `error` from the +// result blob. + +import { BuiltInTool } from '$lib/enums'; +import { + DEFAULT_LANGUAGE, + FILE_PATH_SEPARATOR_REGEX, + TEXT_LANGUAGE_PREFIX_REGEX +} from '$lib/constants'; +import { getFileTypeByExtension, tryParseToolResultObject, type AgenticSection } from '$lib/utils'; +import { parseToolArgs } from './_shared'; + +export type WriteFileMeta = { + fileName: string; + filePath: string; + language: string; + content: string; + bytesWritten?: number; + resultMessage?: string; + errorMessage?: string; +}; + +export function parseWriteFileMeta(section: AgenticSection): WriteFileMeta | null { + const args = parseToolArgs(BuiltInTool.WRITE_FILE, section, { partial: true }); + if (!args) return null; + + // Tool contracts drifted over time: some models emit `path`, + // others `file_path` / `filePath`. Accept all three. + const rawPath = args.path ?? args.file_path ?? args.filePath; + if (typeof rawPath !== 'string' || !rawPath) return null; + + const fileName = rawPath.split(FILE_PATH_SEPARATOR_REGEX).pop() || rawPath; + const content = typeof args.content === 'string' ? args.content : ''; + const language = + getFileTypeByExtension(rawPath)?.replace(TEXT_LANGUAGE_PREFIX_REGEX, '') ?? DEFAULT_LANGUAGE; + + const resultObj = tryParseToolResultObject(section.toolResult); + const bytesWritten = + resultObj && Number.isFinite(Number(resultObj.bytes)) ? Number(resultObj.bytes) : undefined; + const resultMessage = typeof resultObj?.result === 'string' ? resultObj.result : undefined; + const errorMessage = typeof resultObj?.error === 'string' ? resultObj.error : undefined; + + return { + fileName, + filePath: rawPath, + language, + content, + bytesWritten, + resultMessage, + errorMessage + }; +} diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte index 254031979..17d8e21d7 100644 --- a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageActions/ChatMessageActionCard/ChatMessageActionCard.svelte @@ -1,4 +1,5 @@ -{#snippet renderSection(section: (typeof sectionsParsed)[number], index: number)} +{#snippet renderSection(section: AgenticSection, index: number)} {#if section.type === AgenticSectionType.TEXT}
- {:else if section.type === AgenticSectionType.TOOL_CALL_STREAMING} - {@const streamingIcon = isStreaming ? Loader2 : Loader2} - {@const streamingIconClass = isStreaming ? 'h-4 w-4 animate-spin' : 'h-4 w-4'} - - toggleExpanded(index, section)} - > -
-
- Arguments: - - {#if isStreaming} - - {/if} -
- {#if section.toolArgs} - - {:else if isStreaming} -
- Receiving arguments... -
- {:else} -
- Response was truncated -
- {/if} -
- - {:else if section.type === AgenticSectionType.TOOL_CALL || section.type === AgenticSectionType.TOOL_CALL_PENDING} - {@const isPending = section.type === AgenticSectionType.TOOL_CALL_PENDING} - {@const toolIcon = isPending ? Loader2 : Wrench} - {@const toolIconClass = isPending ? 'h-4 w-4 animate-spin' : 'h-4 w-4'} - - + {:else if section.type === AgenticSectionType.TOOL_CALL || section.type === AgenticSectionType.TOOL_CALL_PENDING || section.type === AgenticSectionType.TOOL_CALL_STREAMING} + toggleExpanded(index, section)} - > - {#if section.toolArgs && section.toolArgs !== '{}'} -
-
Arguments:
- - -
- {/if} - -
-
- Result: - - {#if isPending} - - {/if} -
- {#if isPending} -
- Waiting for result... -
- {:else if section.toolResult} -
- {#each section.parsedLines as line, i (i)} -
- {line.text} -
- {#if line.image} - {line.image.name} - {/if} - {/each} -
- {:else} -
No output
- {/if} -
- - {:else if section.type === AgenticSectionType.REASONING} - {@const reasoningSubtitle = section.wasInterrupted - ? hasReasoningError - ? 'Error' - : 'Cancelled' - : isStreaming - ? '' - : undefined} - - toggleExpanded(index, section)} - > -
- {#if renderThinkingAsMarkdown} - - {:else} -
- {section.content} -
- {/if} -
- - {:else if section.type === AgenticSectionType.REASONING_PENDING} - {@const reasoningTitle = isStreaming ? 'Reasoning...' : 'Reasoning'} - {@const reasoningSubtitle = isStreaming ? '' : hasReasoningError ? 'Error' : 'Cancelled'} - - toggleExpanded(index, section)} - > -
- {#if renderThinkingAsMarkdown} - - {:else} -
- {section.content} -
- {/if} -
- + /> {/if} {/snippet} -
+
{#if turnGroups.length > 1} {#each turnGroups as turn, turnIndex (turnIndex)} {@const turnStats = message?.timings?.agentic?.perTurn?.[turnIndex]} -
+
{#each turn.sections as section, sIdx (turn.flatIndices[sIdx])} {@render renderSection(section, turn.flatIndices[sIdx])} {/each} - {#if turnStats && showMessageStats} -
+ {#if turnStats && showAgenticTurnStats} +
{/each} {:else} - {#each sectionsParsed as section, index (index)} + {#each sections as section, index (index)} {@render renderSection(section, index)} {/each} {/if} @@ -404,7 +267,6 @@ flex-direction: column; width: 100%; max-width: 48rem; - gap: 1rem; } .agentic-content > :global(*), diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageReasoningBlock.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageReasoningBlock.svelte new file mode 100644 index 000000000..833cae5db --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageReasoningBlock.svelte @@ -0,0 +1,151 @@ + + + +
+ {#if renderThinkingAsMarkdown} + + {:else} +
+ {section.content} +
+ {/if} +
+ + + diff --git a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte index c470ac729..dca24afd4 100644 --- a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenActionScrollDown.svelte @@ -1,4 +1,5 @@ @@ -76,59 +45,54 @@ open = value; onToggle?.(); }} - class="{className} my-0!" + class={cn('group/collapsible', 'my-0!', className)} > - - -
-
- {#if IconComponent} - - {/if} + +
+ {#if iconUrl} + + {:else if IconComponent} + + {/if} - {title} - - {#if subtitle} - {subtitle} - {/if} -
- - {#if displayedPreview && !showThoughtInProgress} -
-
- {displayedPreview} -
- {#if displayedOverflow > 0} - {displayedOverflow}+ chars - {/if} -
+ + {#if titleSnippet} + {@render titleSnippet()} + {:else} + {title} {/if} -
+ -
- + {#if subtitle} + {subtitle} + {/if} +
- Toggle content -
- + - -
+ Toggle content + + + +
+
{@render children()}
- - +
+ diff --git a/tools/ui/src/lib/components/app/content/CollapsibleTerminalBlock.svelte b/tools/ui/src/lib/components/app/content/CollapsibleTerminalBlock.svelte new file mode 100644 index 000000000..4370aea42 --- /dev/null +++ b/tools/ui/src/lib/components/app/content/CollapsibleTerminalBlock.svelte @@ -0,0 +1,97 @@ + + + { + open = value; + onToggle?.(); + }} + class={cn('group/collapsible', 'overflow-hidden rounded-md', className)} + style="background: var(--code-background); border: 1px solid color-mix(in oklch, var(--border) 30%, transparent);" +> + +
+ {#if iconUrl} + + {:else if IconComponent} + + {/if} + + + {#if titleSnippet} + {@render titleSnippet()} + {:else} + {title} + {/if} + + + {#if subtitle} + {subtitle} + {/if} +
+ + + + Toggle content + + + +
+ {@render children()} +
+ + diff --git a/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css b/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css index b0e04ca62..41813f4fd 100644 --- a/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css +++ b/tools/ui/src/lib/components/app/content/MarkdownContent/markdown-content.css @@ -19,8 +19,16 @@ line-height: 1.75; } +.markdown-content :global(.markdown-block:first-child p:first-child) { + margin-block-start: 0; +} + +.markdown-content :global(.markdown-block:last-child p:last-child) { + margin-block-end: 0; +} + .markdown-content :global(:is(h1, h2, h3, h4, h5, h6):first-child) { - margin-top: 0; + margin-top: 0.5rem; } /* Headers with consistent spacing */ diff --git a/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte b/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte index bb3185f40..39540e7a8 100644 --- a/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte +++ b/tools/ui/src/lib/components/app/content/MermaidPreviewControls.svelte @@ -1,4 +1,5 @@
- + 
{@html highlightedHtml}
diff --git a/tools/ui/src/lib/components/app/content/index.ts b/tools/ui/src/lib/components/app/content/index.ts index 5d2884bb2..5cfdd1b9c 100644 --- a/tools/ui/src/lib/components/app/content/index.ts +++ b/tools/ui/src/lib/components/app/content/index.ts @@ -68,7 +68,6 @@ export { default as SyntaxHighlightedCode } from './SyntaxHighlightedCode.svelte * ```svelte * @@ -78,6 +77,22 @@ export { default as SyntaxHighlightedCode } from './SyntaxHighlightedCode.svelte */ export { default as CollapsibleContentBlock } from './CollapsibleContentBlock.svelte'; +/** + * **CollapsibleTerminalBlock** - Expandable content card with a terminal-style frame + * + * Same shape as CollapsibleContentBlock, but with a `code-background` + * fill, subtle border, and tightened padding suited for shell command + * output and similar dense / monospace content. + * + * @example + * ```svelte + * + * 
{output}
+ * + * ``` + */ +export { default as CollapsibleTerminalBlock } from './CollapsibleTerminalBlock.svelte'; + /** * **MermaidPreview** - Interactive Mermaid diagram viewer * diff --git a/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte b/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte index eb162a557..f741b544b 100644 --- a/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte +++ b/tools/ui/src/lib/components/app/dialogs/DialogMcpResourcesBrowser.svelte @@ -1,4 +1,5 @@ - + - Add New Server + Add New MCP Server + {#if recommendationsToShow.length > 0} +
+
+

Recommended Servers

+ +
+ +
+ {#each recommendationsToShow as recommendation (recommendation.id)} + handleRecommendationClick(recommendation.id)} + selected={selectedRecommendationId === recommendation.id} + dimmed={hasSelection && selectedRecommendationId !== recommendation.id} + /> + {/each} +
+
+ {/if} +
(newServerUseProxy = v)} urlError={newServerUrl ? newServerUrlError : null} id="new-server" + bind:wantsAuthorization={newServerWantsAuthorization} + required={authRequired} />
diff --git a/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte b/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte index a6c20291f..89d23cd4b 100644 --- a/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte +++ b/tools/ui/src/lib/components/app/dialogs/DialogModelNotAvailable.svelte @@ -1,4 +1,5 @@ + + +
+ {#if activeIconUrl} + + {/if} + +

{server.name}

+
+ +

{server.description}

+ diff --git a/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte b/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte index 7f05d5fef..a7472add3 100644 --- a/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte +++ b/tools/ui/src/lib/components/app/mcp/McpServerForm.svelte @@ -5,9 +5,14 @@ import type { KeyValuePair } from '$lib/types'; import { parseHeadersToArray, serializeHeaders } from '$lib/utils'; import { UrlProtocol } from '$lib/enums'; - import { MCP_SERVER_URL_PLACEHOLDER } from '$lib/constants'; + import { + AUTHORIZATION_HEADER, + BEARER_PREFIX, + CLI_FLAGS, + MCP_SERVER_URL_PLACEHOLDER, + REDACTED_HEADERS + } from '$lib/constants'; import { mcpStore } from '$lib/stores/mcp.svelte'; - import { CLI_FLAGS } from '$lib/constants'; interface Props { url: string; @@ -18,6 +23,22 @@ onUseProxyChange?: (useProxy: boolean) => void; urlError?: string | null; id?: string; + /** + * "Wants Authorization" is the user's *intent* to add a Bearer token + * (separate from `hasAuthorization` which reflects what's already in + * the headers). Bindable so a parent - e.g. the recommendation cards + * on the "Add New Server" dialog - can flip the switch on when the + * picked server ships a `needsAuthorization: true` flag. + */ + wantsAuthorization?: boolean; + /** + * Marks the "Authorization" field as required. Locks the toggle so the + * user can't dismiss it, and visually marks the field with a red + * asterisk. The parent is expected to gate its submit affordance on + * the bearer token actually being filled. Used by the "Add New Server" + * dialog for recommendations whose `needsAuthorization` flag is true. + */ + required?: boolean; } let { @@ -28,7 +49,9 @@ onHeadersChange, onUseProxyChange, urlError = null, - id = 'server' + id = 'server', + wantsAuthorization = $bindable(false), + required = false }: Props = $props(); let isWebSocket = $derived( @@ -38,14 +61,11 @@ let headerPairs = $derived(parseHeadersToArray(headers)); - const AUTHORIZATION_HEADER = 'Authorization'; - const BEARER_PREFIX = 'Bearer '; - // Heuristic: this dedicated UI only owns Authorization headers that already // carry a Bearer scheme. Anything else (e.g. Basic, raw tokens) stays in the // KV section so the user can still edit those values verbatim. const matchesAuthorizationKey = (key: string): boolean => - key.trim().toLowerCase() === AUTHORIZATION_HEADER.toLowerCase(); + REDACTED_HEADERS.has(key.trim().toLowerCase()); const isBearerScheme = (value: string): boolean => value.trim().toLowerCase().startsWith(BEARER_PREFIX.toLowerCase()); @@ -55,8 +75,6 @@ let hasAuthorization = $derived(headerPairs.some(ownedByBearerUi)); - let wantsAuthorization = $state(false); - let showAuthorization = $derived(hasAuthorization || wantsAuthorization); let urlInput: HTMLInputElement | null = $state(null); @@ -119,7 +137,7 @@
-
-