From bef69f1306285b95ff65fbd726f9dc4958f205d0 Mon Sep 17 00:00:00 2001 From: Winston Ma Date: Mon, 1 Jun 2026 20:03:32 +0800 Subject: [PATCH 01/20] vulkan: reduce host memory lock contention (#23376) * vulkan: reduces lock contention * replace unique_lock with lock_guard --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 3cf191f20..c3d4c7a71 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -62,6 +62,7 @@ typedef struct VkPhysicalDeviceCooperativeMatrixDecodeVectorFeaturesNV { #include #include #include +#include #include #include #include @@ -618,6 +619,7 @@ static constexpr std::initializer_list> rms_norm_mul_rope_vie struct vk_device_struct { std::recursive_mutex mutex; + mutable std::shared_mutex pinned_memory_mutex; vk::PhysicalDevice physical_device; vk::PhysicalDeviceProperties properties; @@ -7010,7 +7012,7 @@ static void * ggml_vk_host_malloc(vk_device& device, size_t size) { return nullptr; } - std::lock_guard guard(device->mutex); + std::lock_guard guard(device->pinned_memory_mutex); device->pinned_memory.push_back(std::make_tuple(buf->ptr, size, buf)); return buf->ptr; @@ -7021,7 +7023,7 @@ static void ggml_vk_host_free(vk_device& device, void* ptr) { return; } VK_LOG_MEMORY("ggml_vk_host_free(" << ptr << ")"); - std::lock_guard guard(device->mutex); + std::lock_guard guard(device->pinned_memory_mutex); vk_buffer buf; size_t index; @@ -7045,7 +7047,7 @@ static void ggml_vk_host_free(vk_device& device, void* ptr) { } static void ggml_vk_host_get(const vk_device& device, const void * ptr, vk_buffer& buf, size_t& buf_offset) { - std::lock_guard guard(device->mutex); + std::shared_lock guard(device->pinned_memory_mutex); buf = nullptr; buf_offset = 0; for (size_t i = 0; i < device->pinned_memory.size(); i++) { From 55ac0909e5526efa950ce69d06ad0e8e7ebc7e0a Mon Sep 17 00:00:00 2001 From: Jeff Bolz Date: Mon, 1 Jun 2026 07:04:01 -0500 Subject: [PATCH 02/20] vulkan: don't hold the device mutex while compiling pipelines (#23641) * vulkan: don't hold the device mutex while compiling pipelines We need to hold a lock while we traverse all pipelines and lazily initialize them, but we don't need to hold it while the pipeline is being compiled. And it doesn't need to be the same lock as the device mutex. We call load_shaders each time a pipeline is needed, so we only need to compile that one pipeline (and, for example, don't want to end up compiling a pipeline that another thread should be compiling). * remove 'needed' --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 144 ++++++++++++++++++--------- 1 file changed, 99 insertions(+), 45 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index c3d4c7a71..e7d04634b 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -65,6 +65,7 @@ typedef struct VkPhysicalDeviceCooperativeMatrixDecodeVectorFeaturesNV { #include #include #include +#include #include #if defined(_MSC_VER) @@ -159,8 +160,9 @@ struct vk_pipeline_struct { uint32_t align; // true if fields have been set by ggml_vk_create_pipeline bool initialized {}; - // set to true to request the pipeline is compiled - std::atomic needed {}; + // true while a compile is in flight, used to dedupe concurrent claims. + // Protected by device->compile_mutex. + bool compile_pending {}; // set to true when the shader has been compiled std::atomic compiled {}; // number of registers used, extracted from pipeline executable properties @@ -621,6 +623,13 @@ struct vk_device_struct { std::recursive_mutex mutex; mutable std::shared_mutex pinned_memory_mutex; + // Guards compile_pending, all_pipelines, and the dynamic pipeline maps + // (flash_attn, fa_mask_opt, solve_tri, conv2d, etc). The actual compile + // runs with no lock held, so different pipelines can compile in parallel. + // Lock order is device->mutex -> compile_mutex, never the reverse. + std::mutex compile_mutex; + std::condition_variable compile_cv; + vk::PhysicalDevice physical_device; vk::PhysicalDeviceProperties properties; std::string name; @@ -1729,7 +1738,7 @@ struct ggml_vk_garbage_collector { }; static void ggml_vk_preallocate_buffers(ggml_backend_vk_context * ctx, vk_context subctx); -static void ggml_vk_load_shaders(vk_device& device); +static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested = nullptr); static void ggml_pipeline_allocate_descriptor_sets(ggml_backend_vk_context * ctx); static bool vk_memory_logger_enabled = false; @@ -2196,11 +2205,6 @@ static void ggml_vk_wait_for_fence(ggml_backend_vk_context * ctx) { ctx->device->device.resetFences({ ctx->fence }); } -// variables to track number of compiles in progress -static uint32_t compile_count = 0; -static std::mutex compile_count_mutex; -static std::condition_variable compile_count_cond; - static constexpr uint32_t kSpvOpCooperativeMatrixLoadTensorNV = 5367; static constexpr uint32_t kSpvCapabilityCooperativeMatrixDecodeVectorNV = 5447; static constexpr uint32_t kSpvTensorAddressingDecodeVectorFuncBit = 0x4; @@ -2495,7 +2499,6 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin std::cerr << "ggml_vulkan: " << e.what() << std::endl; throw e; } - pipeline->compiled = true; if (vk_instance.debug_utils_support) { vk::DebugUtilsObjectNameInfoEXT duoni; @@ -2544,14 +2547,13 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin } } - device->all_pipelines.push_back(pipeline); - { - std::lock_guard guard(compile_count_mutex); - assert(compile_count > 0); - compile_count--; + std::lock_guard guard(device->compile_mutex); + device->all_pipelines.push_back(pipeline); + pipeline->compiled = true; + pipeline->compile_pending = false; } - compile_count_cond.notify_all(); + device->compile_cv.notify_all(); } static void ggml_vk_destroy_pipeline(vk::Device& device, vk_pipeline& pipeline) { @@ -2567,8 +2569,7 @@ static void ggml_pipeline_request_descriptor_sets(ggml_backend_vk_context *ctx, VK_LOG_DEBUG("ggml_pipeline_request_descriptor_sets(" << pipeline->name << ", " << n << ")"); ctx->pipeline_descriptor_set_requirements += n; if (!pipeline->compiled) { - pipeline->needed = true; - ggml_vk_load_shaders(ctx->device); + ggml_vk_load_shaders(ctx->device, pipeline); } ggml_pipeline_allocate_descriptor_sets(ctx); } @@ -3567,10 +3568,26 @@ static bool ggml_vk_fa_scalar_uses_mmq(const vk_device& device, ggml_type k_type #endif } -static void ggml_vk_load_shaders(vk_device& device) { +// load_shaders walks the pipeline list under compile_mutex and either claims +// the requested pipeline for compilation or, if another thread is already +// compiling it, drops the lock and waits on compile_cv. Compiles themselves +// run unlocked. +struct CompileTask { + vk_pipeline pipeline; + size_t spv_size; + const void * spv_data; + std::string entrypoint; + uint32_t parameter_count; + std::array wg_denoms; + std::vector specialization_constants; + bool disable_robustness; + bool require_full_subgroups; + uint32_t required_subgroup_size; +}; + +static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { VK_LOG_DEBUG("ggml_vk_load_shaders(" << device->name << ")"); - std::lock_guard guard(device->mutex); // some shaders have a minimum subgroup size const uint32_t subgroup_size_8 = std::max(device->subgroup_size, 8u); const uint32_t subgroup_size_16 = std::max(device->subgroup_size, 16u); @@ -3600,6 +3617,15 @@ static void ggml_vk_load_shaders(vk_device& device) { l_mmqid_wg_denoms, m_mmqid_wg_denoms, s_mmqid_wg_denoms; uint32_t l_align, m_align, s_align; + + vk_pipeline wait_pipeline; + CompileTask claimed_task {}; + bool has_claimed_task = false; + + // The rest of the walk reads and writes shared device state, so hold the + // lock until we're done deciding what to compile. + std::unique_lock compile_lock(device->compile_mutex); + if (device->coopmat2) { // spec constants and tile sizes for non-quant matmul/matmul_id l_warptile = { 256, 128, 256, 64, 1 }; @@ -3785,7 +3811,6 @@ static void ggml_vk_load_shaders(vk_device& device) { device->pipeline_matmul_id_bf16 = std::make_shared(); } - std::vector> compiles; auto const &ggml_vk_create_pipeline = [&](vk_device& device, vk_pipeline& base_pipeline, const char *name, size_t spv_size, const void* spv_data, const char *entrypoint, uint32_t parameter_count, uint32_t push_constant_size, std::array wg_denoms, const std::vector& specialization_constants, uint32_t align, bool disable_robustness = false, bool require_full_subgroups = false, uint32_t required_subgroup_size = 0) { @@ -3819,23 +3844,33 @@ static void ggml_vk_load_shaders(vk_device& device) { #endif } - if (!pipeline->needed || pipeline->compiled) { + // We only care about the pipeline this call asked for; the rest + // (including the 64-bit indexing variant) are handled by their + // own request_descriptor_sets / load_shaders calls. + if (pipeline.get() != requested.get()) { continue; } - // TODO: We're no longer benefitting from the async compiles (shaders are - // compiled individually, as needed) and this complexity can be removed. - { - // wait until fewer than N compiles are in progress - uint32_t N = std::max(1u, std::thread::hardware_concurrency()); - std::unique_lock guard(compile_count_mutex); - while (compile_count >= N) { - compile_count_cond.wait(guard); - } - compile_count++; + + if (pipeline->compiled) { + continue; } - compiles.push_back(std::async(ggml_vk_create_pipeline_func, std::ref(device), std::ref(pipeline), spv_size, spv_data, entrypoint, - parameter_count, wg_denoms, specialization_constants, disable_robustness, require_full_subgroups, required_subgroup_size)); + wait_pipeline = pipeline; + + if (!pipeline->compile_pending) { + pipeline->compile_pending = true; + claimed_task.pipeline = pipeline; + claimed_task.spv_size = spv_size; + claimed_task.spv_data = spv_data; + claimed_task.entrypoint = entrypoint; + claimed_task.parameter_count = parameter_count; + claimed_task.wg_denoms = wg_denoms; + claimed_task.specialization_constants = specialization_constants; + claimed_task.disable_robustness = disable_robustness; + claimed_task.require_full_subgroups = require_full_subgroups; + claimed_task.required_subgroup_size = required_subgroup_size; + has_claimed_task = true; + } } }; @@ -5332,8 +5367,25 @@ static void ggml_vk_load_shaders(vk_device& device) { } } - for (auto &c : compiles) { - c.wait(); + // Drop compile_mutex so other threads can walk while we compile. + compile_lock.unlock(); + + // Compile what we claimed; create_pipeline_func reacquires compile_mutex + // at the end to flip compile_pending/compiled and notify waiters. + if (has_claimed_task) { + auto & task = claimed_task; + ggml_vk_create_pipeline_func(device, task.pipeline, task.spv_size, task.spv_data, + task.entrypoint, task.parameter_count, task.wg_denoms, + task.specialization_constants, task.disable_robustness, + task.require_full_subgroups, task.required_subgroup_size); + } + + // Another thread may be compiling the pipeline we need; block on it here. + if (wait_pipeline) { + std::unique_lock wait_lock(device->compile_mutex); + device->compile_cv.wait(wait_lock, [&] { + return wait_pipeline->compiled.load(); + }); } } @@ -9722,7 +9774,7 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx vk_pipeline pipeline = nullptr; { - std::lock_guard guard(ctx->device->mutex); + std::lock_guard guard(ctx->device->compile_mutex); auto &pipelines = ctx->device->pipeline_flash_attn_f32_f16; auto it = pipelines.find(fa_pipeline_state); if (it != pipelines.end()) { @@ -9786,13 +9838,15 @@ static void ggml_vk_flash_attn(ggml_backend_vk_context * ctx, vk_context& subctx vk_pipeline pipeline_fa_mask_opt = nullptr; if (use_mask_opt) { - std::lock_guard guard(ctx->device->mutex); - auto &pipelines = ctx->device->pipeline_fa_mask_opt; - auto it = pipelines.find({Br, Bc}); - if (it != pipelines.end()) { - pipeline_fa_mask_opt = it->second; - } else { - pipelines[{Br, Bc}] = pipeline_fa_mask_opt = std::make_shared(); + { + std::lock_guard guard(ctx->device->compile_mutex); + auto &pipelines = ctx->device->pipeline_fa_mask_opt; + auto it = pipelines.find({Br, Bc}); + if (it != pipelines.end()) { + pipeline_fa_mask_opt = it->second; + } else { + pipelines[{Br, Bc}] = pipeline_fa_mask_opt = std::make_shared(); + } } assert(pipeline_fa_mask_opt); ggml_pipeline_request_descriptor_sets(ctx, pipeline_fa_mask_opt, 1); @@ -10326,7 +10380,7 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const vk_pipeline pipeline = nullptr; { - std::lock_guard guard(ctx->device->mutex); + std::lock_guard guard(ctx->device->compile_mutex); auto it = ctx->device->pipeline_solve_tri_f32.find(solve_tri_pipeline_state); if (it != ctx->device->pipeline_solve_tri_f32.end()) { pipeline = it->second; @@ -10485,7 +10539,7 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const vk_pipeline pipeline = nullptr; { - std::lock_guard guard(ctx->device->mutex); + std::lock_guard guard(ctx->device->compile_mutex); auto it = pipelines->find(conv2d_pipeline_state); if (it != pipelines->end()) { pipeline = it->second; From 95b8b8ec1a9e77bcc7b3fb04da82c9f35cb12a79 Mon Sep 17 00:00:00 2001 From: Shrivas Shankar <86219405+shrivasshankar@users.noreply.github.com> Date: Mon, 1 Jun 2026 07:40:28 -0500 Subject: [PATCH 03/20] metal: template GLU kernels to support f16/f32 (#23882) Drops the hardcoded f32 GLU kernels in favor of a single template. We now load/store in the native tensor type (half or float) to save memory bandwidth, but keep the actual ALU compute in float to avoid exploding math in geglu/swiglu. Also opened up the dispatch gate to allow f16 inputs. --- ggml/src/ggml-metal/ggml-metal-device.m | 2 +- ggml/src/ggml-metal/ggml-metal.metal | 96 +++++++++++++++++-------- 2 files changed, 67 insertions(+), 31 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 885344ec6..196af1026 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1107,7 +1107,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_GLU_OP_SWIGLU_OAI: case GGML_GLU_OP_GEGLU_ERF: case GGML_GLU_OP_GEGLU_QUICK: - return ggml_is_contiguous_1(op->src[0]) && op->src[0]->type == GGML_TYPE_F32; + return ggml_is_contiguous_1(op->src[0]) && (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16); default: return false; } diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 4adf4614a..2bd310d94 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -1421,7 +1421,8 @@ template [[host_name("kernel_repeat_f16")]] kernel kernel_repeat_t kernel_repeat template [[host_name("kernel_repeat_i32")]] kernel kernel_repeat_t kernel_repeat; template [[host_name("kernel_repeat_i16")]] kernel kernel_repeat_t kernel_repeat; -kernel void kernel_reglu_f32( +template +kernel void kernel_reglu( constant ggml_metal_kargs_glu & args, device const char * src0, device const char * src1, @@ -1429,19 +1430,25 @@ kernel void kernel_reglu_f32( uint tgpig[[threadgroup_position_in_grid]], uint tpitg[[thread_position_in_threadgroup]], uint ntg[[threads_per_threadgroup]]) { - device const float * src0_row = (device const float *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; - device const float * src1_row = (device const float *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; - device float * dst_row = (device float *) ((device char *) dst + tgpig*args.nb1); + device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; + device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; + device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { const float x0 = src0_row[i0]; const float x1 = src1_row[i0]; - dst_row[i0] = x0*x1*(x0 > 0.0f); + dst_row[i0] = (T)(x0*x1*(x0 > 0.0f)); } } -kernel void kernel_geglu_f32( +typedef decltype(kernel_reglu) kernel_reglu_t; + +template [[host_name("kernel_reglu_f32")]] kernel kernel_reglu_t kernel_reglu; +template [[host_name("kernel_reglu_f16")]] kernel kernel_reglu_t kernel_reglu; + +template +kernel void kernel_geglu( constant ggml_metal_kargs_glu & args, device const char * src0, device const char * src1, @@ -1449,9 +1456,9 @@ kernel void kernel_geglu_f32( uint tgpig[[threadgroup_position_in_grid]], uint tpitg[[thread_position_in_threadgroup]], uint ntg[[threads_per_threadgroup]]) { - device const float * src0_row = (device const float *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; - device const float * src1_row = (device const float *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; - device float * dst_row = (device float *) ((device char *) dst + tgpig*args.nb1); + device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; + device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; + device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { const float x0 = src0_row[i0]; @@ -1459,11 +1466,17 @@ kernel void kernel_geglu_f32( const float gelu = 0.5f*x0*(1.0f + precise::tanh(SQRT_2_OVER_PI*x0*(1.0f + GELU_COEF_A*x0*x0))); - dst_row[i0] = gelu*x1; + dst_row[i0] = (T)(gelu*x1); } } -kernel void kernel_swiglu_f32( +typedef decltype(kernel_geglu) kernel_geglu_t; + +template [[host_name("kernel_geglu_f32")]] kernel kernel_geglu_t kernel_geglu; +template [[host_name("kernel_geglu_f16")]] kernel kernel_geglu_t kernel_geglu; + +template +kernel void kernel_swiglu( constant ggml_metal_kargs_glu & args, device const char * src0, device const char * src1, @@ -1471,9 +1484,9 @@ kernel void kernel_swiglu_f32( uint tgpig[[threadgroup_position_in_grid]], uint tpitg[[thread_position_in_threadgroup]], uint ntg[[threads_per_threadgroup]]) { - device const float * src0_row = (device const float *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; - device const float * src1_row = (device const float *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; - device float * dst_row = (device float *) ((device char *) dst + tgpig*args.nb1); + device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; + device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; + device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { const float x0 = src0_row[i0]; @@ -1481,11 +1494,17 @@ kernel void kernel_swiglu_f32( const float silu = x0 / (1.0f + exp(-x0)); - dst_row[i0] = silu*x1; + dst_row[i0] = (T)(silu*x1); } } -kernel void kernel_swiglu_oai_f32( +typedef decltype(kernel_swiglu) kernel_swiglu_t; + +template [[host_name("kernel_swiglu_f32")]] kernel kernel_swiglu_t kernel_swiglu; +template [[host_name("kernel_swiglu_f16")]] kernel kernel_swiglu_t kernel_swiglu; + +template +kernel void kernel_swiglu_oai( constant ggml_metal_kargs_glu & args, device const char * src0, device const char * src1, @@ -1493,9 +1512,9 @@ kernel void kernel_swiglu_oai_f32( uint tgpig[[threadgroup_position_in_grid]], uint tpitg[[thread_position_in_threadgroup]], uint ntg[[threads_per_threadgroup]]) { - device const float * src0_row = (device const float *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; - device const float * src1_row = (device const float *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; - device float * dst_row = (device float *) ((device char *) dst + tgpig*args.nb1); + device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; + device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; + device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { float x0 = src0_row[i0]; @@ -1507,11 +1526,17 @@ kernel void kernel_swiglu_oai_f32( float out_glu = x0 / (1.0f + exp(-x0 * args.alpha)); out_glu = out_glu * (1.0f + x1); - dst_row[i0] = out_glu; + dst_row[i0] = (T)out_glu; } } -kernel void kernel_geglu_erf_f32( +typedef decltype(kernel_swiglu_oai) kernel_swiglu_oai_t; + +template [[host_name("kernel_swiglu_oai_f32")]] kernel kernel_swiglu_oai_t kernel_swiglu_oai; +template [[host_name("kernel_swiglu_oai_f16")]] kernel kernel_swiglu_oai_t kernel_swiglu_oai; + +template +kernel void kernel_geglu_erf( constant ggml_metal_kargs_glu & args, device const char * src0, device const char * src1, @@ -1519,9 +1544,9 @@ kernel void kernel_geglu_erf_f32( uint tgpig[[threadgroup_position_in_grid]], uint tpitg[[thread_position_in_threadgroup]], uint ntg[[threads_per_threadgroup]]) { - device const float * src0_row = (device const float *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; - device const float * src1_row = (device const float *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; - device float * dst_row = (device float *) ((device char *) dst + tgpig*args.nb1); + device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; + device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; + device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { const float x0 = src0_row[i0]; @@ -1529,11 +1554,17 @@ kernel void kernel_geglu_erf_f32( const float gelu_erf = 0.5f*x0*(1.0f+erf_approx(x0*SQRT_2_INV)); - dst_row[i0] = gelu_erf*x1; + dst_row[i0] = (T)(gelu_erf*x1); } } -kernel void kernel_geglu_quick_f32( +typedef decltype(kernel_geglu_erf) kernel_geglu_erf_t; + +template [[host_name("kernel_geglu_erf_f32")]] kernel kernel_geglu_erf_t kernel_geglu_erf; +template [[host_name("kernel_geglu_erf_f16")]] kernel kernel_geglu_erf_t kernel_geglu_erf; + +template +kernel void kernel_geglu_quick( constant ggml_metal_kargs_glu & args, device const char * src0, device const char * src1, @@ -1541,9 +1572,9 @@ kernel void kernel_geglu_quick_f32( uint tgpig[[threadgroup_position_in_grid]], uint tpitg[[thread_position_in_threadgroup]], uint ntg[[threads_per_threadgroup]]) { - device const float * src0_row = (device const float *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; - device const float * src1_row = (device const float *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; - device float * dst_row = (device float *) ((device char *) dst + tgpig*args.nb1); + device const T * src0_row = (device const T *) ((device const char *) src0 + tgpig*args.nb01) + args.i00; + device const T * src1_row = (device const T *) ((device const char *) src1 + tgpig*args.nb11) + args.i10; + device T * dst_row = (device T *) ((device char *) dst + tgpig*args.nb1); for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) { const float x0 = src0_row[i0]; @@ -1551,10 +1582,15 @@ kernel void kernel_geglu_quick_f32( const float gelu_quick = x0*(1.0f/(1.0f+exp(GELU_QUICK_COEF*x0))); - dst_row[i0] = gelu_quick*x1; + dst_row[i0] = (T)(gelu_quick*x1); } } +typedef decltype(kernel_geglu_quick) kernel_geglu_quick_t; + +template [[host_name("kernel_geglu_quick_f32")]] kernel kernel_geglu_quick_t kernel_geglu_quick; +template [[host_name("kernel_geglu_quick_f16")]] kernel kernel_geglu_quick_t kernel_geglu_quick; + kernel void kernel_op_sum_f32( constant ggml_metal_kargs_sum & args, device const float * src0, From de6f727aaec7dc477629946d80c803a0bb7af0a1 Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Mon, 1 Jun 2026 23:01:38 +0800 Subject: [PATCH 04/20] llama: limit max outputs of `llama_context` (#23861) * llama: save more VRAM by reserving n_outputs == n_seqs when possible * add n_outputs_per_seq * move n_outputs_max to server-context * change ubatch to batch everywhere --- common/common.cpp | 1 + common/common.h | 1 + include/llama.h | 1 + src/llama-context.cpp | 21 ++++++++---- src/llama-cparams.h | 1 + tools/server/server-context.cpp | 57 ++++++++++++++++++++++++++++++--- 6 files changed, 71 insertions(+), 11 deletions(-) diff --git a/common/common.cpp b/common/common.cpp index 97daf2817..81b8b7500 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1563,6 +1563,7 @@ struct llama_context_params common_context_params_to_llama(const common_params & cparams.n_ctx = params.n_ctx; cparams.n_seq_max = params.n_parallel; cparams.n_rs_seq = params.speculative.need_n_rs_seq(); + cparams.n_outputs_max = std::max(params.n_outputs_max, 0); cparams.n_batch = params.n_batch; cparams.n_ubatch = params.n_ubatch; cparams.n_threads = params.cpuparams.n_threads; diff --git a/common/common.h b/common/common.h index 99898800d..92064a0e4 100644 --- a/common/common.h +++ b/common/common.h @@ -431,6 +431,7 @@ struct common_params { int32_t n_chunks = -1; // max number of chunks to process (-1 = unlimited) int32_t n_parallel = 1; // number of parallel sequences to decode int32_t n_sequences = 1; // number of sequences to decode + int32_t n_outputs_max = 0; // max outputs in a batch (0 = n_batch) int32_t grp_attn_n = 1; // group-attention factor int32_t grp_attn_w = 512; // group-attention width int32_t n_print = -1; // print token count every n tokens (-1 = disabled) diff --git a/include/llama.h b/include/llama.h index e8374c53b..a79a491c5 100644 --- a/include/llama.h +++ b/include/llama.h @@ -339,6 +339,7 @@ extern "C" { uint32_t n_ubatch; // physical maximum batch size uint32_t n_seq_max; // max number of sequences (i.e. distinct states for recurrent models) uint32_t n_rs_seq; // number of recurrent-state snapshots per seq for rollback (0 = no rollback) [EXPERIMENTAL] + uint32_t n_outputs_max; // max outputs in a ubatch (0 = n_batch) int32_t n_threads; // number of threads to use for generation int32_t n_threads_batch; // number of threads to use for batch processing diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 99bdf092b..a7dfb2492 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -182,6 +182,8 @@ llama_context::llama_context( cparams.n_ubatch = std::min(cparams.n_batch, params.n_ubatch == 0 ? params.n_batch : params.n_ubatch); + cparams.n_outputs_max = params.n_outputs_max == 0 ? cparams.n_batch : params.n_outputs_max; + cparams.op_offload = params.op_offload; cparams.kv_unified = params.kv_unified; @@ -531,7 +533,7 @@ void llama_context::sched_reserve() { // note: n_outputs must match n_tokens for embedding models with mean/rank pooling, // because build_pooling creates inp_mean with shape [n_tokens, n_seqs] and multiplies // it with t_embd which is reduced to [n_outputs, ...] via out_ids. if n_outputs != n_tokens, - // the ggml_mul_mat assertion fails. this matches the pp reservation below (line ~553). + // the ggml_mul_mat assertion fails. const uint32_t n_tokens_ch = 16*n_seqs; auto * gf = graph_reserve(n_tokens_ch, n_seqs, n_tokens_ch, mctx.get(), true); if (!gf) { @@ -577,16 +579,18 @@ void llama_context::sched_reserve() { int n_splits_tg = -1; int n_nodes_tg = -1; + const uint32_t n_outputs_pp = std::min(n_tokens, cparams.n_outputs_max); + // reserve pp (prompt processing) graph first so that buffers are only allocated once { - auto * gf = graph_reserve(n_tokens, n_seqs, n_tokens, mctx.get(), + auto * gf = graph_reserve(n_tokens, n_seqs, n_outputs_pp, mctx.get(), model.hparams.no_alloc, model.hparams.no_alloc ? backend_buf_exp_size.data() : nullptr); if (!gf) { if (cparams.pipeline_parallel) { LLAMA_LOG_WARN("%s: compute buffer allocation failed, retrying without pipeline parallelism\n", __func__); cparams.pipeline_parallel = false; sched.reset(ggml_backend_sched_new(backend_ptrs.data(), backend_buft.data(), backend_ptrs.size(), max_nodes, false, cparams.op_offload)); - gf = graph_reserve(n_tokens, n_seqs, n_tokens, mctx.get()); + gf = graph_reserve(n_tokens, n_seqs, n_outputs_pp, mctx.get()); } if (!gf) { throw std::runtime_error("failed to allocate compute pp buffers"); @@ -614,7 +618,7 @@ void llama_context::sched_reserve() { // // auto * gf = graph_reserve(n_tokens, 1, n_tokens, mctx.get()); // - auto * gf = graph_reserve(n_tokens, n_seqs, n_tokens, mctx.get(), model.hparams.no_alloc); + auto * gf = graph_reserve(n_tokens, n_seqs, n_outputs_pp, mctx.get(), model.hparams.no_alloc); if (!gf) { throw std::runtime_error("failed to allocate compute pp buffers"); } @@ -774,7 +778,9 @@ bool llama_context::memory_update(bool optimize) { const uint32_t n_seqs = cparams.n_seq_max; const uint32_t n_tokens = std::min(cparams.n_ctx, cparams.n_ubatch); - auto * gf = graph_reserve(n_tokens, n_seqs, n_tokens, mctx.get()); + const uint32_t n_outputs_max = std::min(n_tokens, cparams.n_outputs_max); + + auto * gf = graph_reserve(n_tokens, n_seqs, n_outputs_max, mctx.get()); if (!gf) { LLAMA_LOG_ERROR("%s: failed to reserve graph after the memory update\n", __func__); } @@ -2140,6 +2146,8 @@ uint32_t llama_context::output_reserve(int32_t n_outputs) { this->n_outputs = 0; + GGML_ASSERT(n_outputs_max <= cparams.n_outputs_max); + return n_outputs_max; } @@ -2226,8 +2234,6 @@ ggml_cgraph * llama_context::graph_reserve( if (n_tokens % n_seqs != 0) { n_tokens = ((n_tokens + (n_seqs - 1)) / n_seqs) * n_seqs; // round to next multiple of n_seqs - n_outputs = std::max(n_outputs, n_tokens); - LLAMA_LOG_DEBUG("%s: making n_tokens a multiple of n_seqs - n_tokens = %u, n_seqs = %u, n_outputs = %u\n", __func__, n_tokens, n_seqs, n_outputs); } @@ -3337,6 +3343,7 @@ llama_context_params llama_context_default_params() { /*.n_ubatch =*/ 512, /*.n_seq_max =*/ 1, /*.n_rs_seq =*/ 0, + /*.n_outputs_max =*/ 0, /*.n_threads =*/ GGML_DEFAULT_N_THREADS, // TODO: better default /*.n_threads_batch =*/ GGML_DEFAULT_N_THREADS, /*.ctx_type =*/ LLAMA_CONTEXT_TYPE_DEFAULT, diff --git a/src/llama-cparams.h b/src/llama-cparams.h index 20ec59fe3..ba4951a09 100644 --- a/src/llama-cparams.h +++ b/src/llama-cparams.h @@ -13,6 +13,7 @@ struct llama_cparams { uint32_t n_ubatch; uint32_t n_seq_max; uint32_t n_rs_seq; // number of recurrent-state snapshots per seq for rollback + uint32_t n_outputs_max; // max outputs supported by the context int32_t n_threads; // number of threads to use for generation int32_t n_threads_batch; // number of threads to use for batch processing diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index bfe3443c1..a5a688254 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -37,6 +37,49 @@ using json = nlohmann::ordered_json; constexpr int HTTP_POLLING_SECONDS = 1; +static uint32_t server_n_outputs_max(const common_params & params) { + const uint32_t n_batch = params.n_batch; + + if (params.embedding || + (params.pooling_type != LLAMA_POOLING_TYPE_UNSPECIFIED && params.pooling_type != LLAMA_POOLING_TYPE_NONE)) { + return n_batch; + } + + uint32_t n_outputs_per_seq = 1; + + for (const auto type : params.speculative.types) { + switch (type) { + case COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE: + case COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3: + case COMMON_SPECULATIVE_TYPE_DRAFT_MTP: + n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + std::max(0, params.speculative.draft.n_max)); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE: + n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + params.speculative.ngram_simple.size_m); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K: + n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + params.speculative.ngram_map_k.size_m); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V: + n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + params.speculative.ngram_map_k4v.size_m); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_MOD: + n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + std::max(0, params.speculative.ngram_mod.n_max)); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_CACHE: + n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + 8); + break; + case COMMON_SPECULATIVE_TYPE_NONE: + case COMMON_SPECULATIVE_TYPE_COUNT: + break; + } + } + + const uint64_t n_outputs = (uint64_t) params.n_parallel * n_outputs_per_seq; + + return std::max(1, std::min(n_batch, n_outputs)); +} + // state diagram: https://github.com/ggml-org/llama.cpp/pull/9283 enum slot_state { SLOT_STATE_IDLE, @@ -753,6 +796,7 @@ private: SRV_INF("loading model '%s'\n", params.model.path.c_str()); params_base = params; + params_base.n_outputs_max = server_n_outputs_max(params_base); std::string & mmproj_path = params_base.mmproj.path; bool has_mmproj = !mmproj_path.empty(); @@ -818,6 +862,10 @@ private: measure_model_bytes = false; } + if (!has_draft) { + params_dft.n_outputs_max = params_base.n_parallel; + } + auto mparams_dft = common_model_params_to_llama(params_dft); auto cparams_dft = common_context_params_to_llama(params_dft); if (spec_mtp) { @@ -941,10 +989,11 @@ private: params_base.model.path.c_str()); auto cparams_mtp = common_context_params_to_llama(params_base); - cparams_mtp.ctx_type = LLAMA_CONTEXT_TYPE_MTP; - cparams_mtp.type_k = params_base.speculative.draft.cache_type_k; - cparams_mtp.type_v = params_base.speculative.draft.cache_type_v; - cparams_mtp.n_rs_seq = 0; + cparams_mtp.ctx_type = LLAMA_CONTEXT_TYPE_MTP; + cparams_mtp.type_k = params_base.speculative.draft.cache_type_k; + cparams_mtp.type_v = params_base.speculative.draft.cache_type_v; + cparams_mtp.n_rs_seq = 0; + cparams_mtp.n_outputs_max = params_base.n_parallel; ctx_dft.reset(llama_init_from_model(model_tgt, cparams_mtp)); if (ctx_dft == nullptr) { From 335abed17decc19ba9e975c0e913ace56d4d722b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adrien=20Gallou=C3=ABt?= Date: Mon, 1 Jun 2026 18:40:10 +0200 Subject: [PATCH 05/20] vendor : update cpp-httplib to 0.46.1 (#23980) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Adrien Gallouët --- scripts/sync_vendor.py | 2 +- vendor/cpp-httplib/httplib.cpp | 36 ++++++++++++++++++++++++---------- vendor/cpp-httplib/httplib.h | 4 ++-- 3 files changed, 29 insertions(+), 13 deletions(-) diff --git a/scripts/sync_vendor.py b/scripts/sync_vendor.py index fec05d97d..8306cf93e 100755 --- a/scripts/sync_vendor.py +++ b/scripts/sync_vendor.py @@ -5,7 +5,7 @@ import os import sys import subprocess -HTTPLIB_VERSION = "refs/tags/v0.46.0" +HTTPLIB_VERSION = "refs/tags/v0.46.1" vendor = { "https://github.com/nlohmann/json/releases/latest/download/json.hpp": "vendor/nlohmann/json.hpp", diff --git a/vendor/cpp-httplib/httplib.cpp b/vendor/cpp-httplib/httplib.cpp index f3555f2d4..b9d05d929 100644 --- a/vendor/cpp-httplib/httplib.cpp +++ b/vendor/cpp-httplib/httplib.cpp @@ -1832,7 +1832,7 @@ int getaddrinfo_with_timeout(const char *node, const char *service, #ifdef _WIN32 // Windows-specific implementation using GetAddrInfoEx with overlapped I/O - OVERLAPPED overlapped = {0}; + OVERLAPPED overlapped = {}; HANDLE event = CreateEventW(nullptr, TRUE, FALSE, nullptr); if (!event) { return EAI_FAIL; } @@ -1841,7 +1841,7 @@ int getaddrinfo_with_timeout(const char *node, const char *service, PADDRINFOEXW result_addrinfo = nullptr; HANDLE cancel_handle = nullptr; - ADDRINFOEXW hints_ex = {0}; + ADDRINFOEXW hints_ex = {}; if (hints) { hints_ex.ai_flags = hints->ai_flags; hints_ex.ai_family = hints->ai_family; @@ -9912,13 +9912,28 @@ bool ClientImpl::process_request(Stream &strm, Request &req, } #endif - // Handle Expect: 100-continue with timeout - if (expect_100_continue && CPPHTTPLIB_EXPECT_100_TIMEOUT_MSECOND > 0) { - time_t sec = CPPHTTPLIB_EXPECT_100_TIMEOUT_MSECOND / 1000; - time_t usec = (CPPHTTPLIB_EXPECT_100_TIMEOUT_MSECOND % 1000) * 1000; - auto ret = detail::select_read(strm.socket(), sec, usec); - if (ret <= 0) { - // Timeout or error: send body anyway (server didn't respond in time) + // Handle Expect: 100-continue. + // + // Wait for an interim/early response by attempting to read the status line + // under a short timeout, instead of trusting raw socket readability. Over + // TLS, post-handshake records (e.g. session tickets) make the socket + // readable without any HTTP response being available; relying on + // `select_read` there caused the body to be withheld forever and the + // request to fail with `Read` (#2458). If no status line arrives within the + // timeout, send the body anyway (matching curl's behavior). + auto status_line_read = false; + if (expect_100_continue && write_request_success) { + if (CPPHTTPLIB_EXPECT_100_TIMEOUT_MSECOND > 0) { + time_t sec = CPPHTTPLIB_EXPECT_100_TIMEOUT_MSECOND / 1000; + time_t usec = (CPPHTTPLIB_EXPECT_100_TIMEOUT_MSECOND % 1000) * 1000; + strm.set_read_timeout(sec, usec); + status_line_read = read_response_line(strm, req, res, false); + strm.set_read_timeout(read_timeout_sec_, read_timeout_usec_); + } + + if (!status_line_read) { + // No interim response within the timeout: send the body and handle the + // response as usual. if (!write_request_body(strm, req, error)) { return false; } expect_100_continue = false; // Switch to normal response handling } @@ -9926,7 +9941,8 @@ bool ClientImpl::process_request(Stream &strm, Request &req, // Receive response and headers // When using Expect: 100-continue, don't auto-skip `100 Continue` response - if (!read_response_line(strm, req, res, !expect_100_continue) || + if ((!status_line_read && + !read_response_line(strm, req, res, !expect_100_continue)) || !detail::read_headers(strm, res.headers)) { if (write_request_success) { error = Error::Read; } output_error_log(error, &req); diff --git a/vendor/cpp-httplib/httplib.h b/vendor/cpp-httplib/httplib.h index af856dd63..cbb549e71 100644 --- a/vendor/cpp-httplib/httplib.h +++ b/vendor/cpp-httplib/httplib.h @@ -8,8 +8,8 @@ #ifndef CPPHTTPLIB_HTTPLIB_H #define CPPHTTPLIB_HTTPLIB_H -#define CPPHTTPLIB_VERSION "0.46.0" -#define CPPHTTPLIB_VERSION_NUM "0x002e00" +#define CPPHTTPLIB_VERSION "0.46.1" +#define CPPHTTPLIB_VERSION_NUM "0x002e01" #ifdef _WIN32 #if defined(_WIN32_WINNT) && _WIN32_WINNT < 0x0A00 From 27d9ed839713e31c7a0ba45e342109a04549834f Mon Sep 17 00:00:00 2001 From: shaofeiqi Date: Mon, 1 Jun 2026 10:06:50 -0700 Subject: [PATCH 06/20] opencl: add basic support for q5_0 and q5_1 (#23548) * opencl: add general q5_0 support * opencl: add general q5_1 support * opencl: support non-uniform workgrp size --------- Co-authored-by: Li He --- ggml/src/ggml-opencl/CMakeLists.txt | 6 + ggml/src/ggml-opencl/ggml-opencl.cpp | 422 +++++++++++++++++- ggml/src/ggml-opencl/kernels/cvt.cl | 100 +++++ .../kernels/mul_mm_q5_0_f32_l4_lm.cl | 173 +++++++ .../kernels/mul_mm_q5_1_f32_l4_lm.cl | 175 ++++++++ .../ggml-opencl/kernels/mul_mv_q5_0_f32.cl | 241 ++++++++++ .../kernels/mul_mv_q5_0_f32_flat.cl | 243 ++++++++++ .../ggml-opencl/kernels/mul_mv_q5_1_f32.cl | 243 ++++++++++ .../kernels/mul_mv_q5_1_f32_flat.cl | 247 ++++++++++ 9 files changed, 1845 insertions(+), 5 deletions(-) create mode 100644 ggml/src/ggml-opencl/kernels/mul_mm_q5_0_f32_l4_lm.cl create mode 100644 ggml/src/ggml-opencl/kernels/mul_mm_q5_1_f32_l4_lm.cl create mode 100644 ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32.cl create mode 100644 ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32_flat.cl create mode 100644 ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32.cl create mode 100644 ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32_flat.cl diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt index 446fb7279..cd15d5732 100644 --- a/ggml/src/ggml-opencl/CMakeLists.txt +++ b/ggml/src/ggml-opencl/CMakeLists.txt @@ -87,6 +87,10 @@ set(GGML_OPENCL_KERNELS mul_mv_q4_1_f32_flat mul_mv_q4_k_f32 mul_mv_q4_k_f32_flat + mul_mv_q5_0_f32 + mul_mv_q5_0_f32_flat + mul_mv_q5_1_f32 + mul_mv_q5_1_f32_flat mul_mv_q5_k_f32 mul_mv_q5_k_f32_flat mul_mv_q6_k_f32 @@ -126,6 +130,8 @@ set(GGML_OPENCL_KERNELS mul_mm_f16_f32_l4_lm mul_mm_q4_0_f32_l4_lm mul_mm_q4_1_f32_l4_lm + mul_mm_q5_0_f32_l4_lm + mul_mm_q5_1_f32_l4_lm mul_mm_q8_0_f32_l4_lm mul_mm_iq4_nl_f32_l4_lm mul_mm_q4_k_f32_l4_lm diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 3f3643a4c..7cafbe0cd 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -576,7 +576,9 @@ struct ggml_backend_opencl_context { cl_kernel kernel_convert_block_q4_0_trans4_ns, kernel_restore_block_q4_0_trans4_ns; cl_kernel kernel_convert_block_q4_1, kernel_restore_block_q4_1; cl_kernel kernel_convert_block_q4_1_trans4_ns, kernel_restore_block_q4_1_trans4_ns; + cl_kernel kernel_convert_block_q5_0, kernel_restore_block_q5_0; cl_kernel kernel_convert_block_q5_0_trans4_ns, kernel_restore_block_q5_0_trans4_ns; + cl_kernel kernel_convert_block_q5_1, kernel_restore_block_q5_1; cl_kernel kernel_convert_block_q5_1_trans4_ns, kernel_restore_block_q5_1_trans4_ns; cl_kernel kernel_convert_block_q4_k_trans4_ns, kernel_restore_block_q4_k_trans4_ns; cl_kernel kernel_convert_block_q5_k_trans4_ns, kernel_restore_block_q5_k_trans4_ns; @@ -604,6 +606,10 @@ struct ggml_backend_opencl_context { cl_kernel kernel_mul_mat_q4_0_f32_1d_8x_flat, kernel_mul_mat_q4_0_f32_1d_16x_flat; cl_kernel kernel_mul_mv_q4_1_f32; cl_kernel kernel_mul_mv_q4_1_f32_flat; + cl_kernel kernel_mul_mv_q5_0_f32; + cl_kernel kernel_mul_mv_q5_0_f32_flat; + cl_kernel kernel_mul_mv_q5_1_f32; + cl_kernel kernel_mul_mv_q5_1_f32_flat; cl_kernel kernel_mul_mv_q4_K_f32; cl_kernel kernel_mul_mv_q4_K_f32_flat; cl_kernel kernel_mul_mv_q5_K_f32; @@ -662,6 +668,8 @@ struct ggml_backend_opencl_context { cl_kernel kernel_mul_mm_f16_f32_l4_lm; cl_kernel kernel_mul_mm_q4_0_f32_l4_lm; cl_kernel kernel_mul_mm_q4_1_f32_l4_lm; + cl_kernel kernel_mul_mm_q5_0_f32_l4_lm; + cl_kernel kernel_mul_mm_q5_1_f32_l4_lm; cl_kernel kernel_mul_mm_q8_0_f32_l4_lm; cl_kernel kernel_mul_mm_q4_k_f32_l4_lm; cl_kernel kernel_mul_mm_q5_k_f32_l4_lm; @@ -1141,8 +1149,12 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { CL_CHECK((backend_ctx->kernel_restore_block_q4_1 = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q4_1", &err), err)); CL_CHECK((backend_ctx->kernel_convert_block_q4_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q4_1_trans4_ns", &err), err)); CL_CHECK((backend_ctx->kernel_restore_block_q4_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q4_1_trans4_ns", &err), err)); + CL_CHECK((backend_ctx->kernel_convert_block_q5_0 = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_0", &err), err)); + CL_CHECK((backend_ctx->kernel_restore_block_q5_0 = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_0", &err), err)); CL_CHECK((backend_ctx->kernel_convert_block_q5_0_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_0_trans4_ns", &err), err)); CL_CHECK((backend_ctx->kernel_restore_block_q5_0_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_0_trans4_ns", &err), err)); + CL_CHECK((backend_ctx->kernel_convert_block_q5_1 = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_1", &err), err)); + CL_CHECK((backend_ctx->kernel_restore_block_q5_1 = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_1", &err), err)); CL_CHECK((backend_ctx->kernel_convert_block_q5_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_1_trans4_ns", &err), err)); CL_CHECK((backend_ctx->kernel_restore_block_q5_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_1_trans4_ns", &err), err)); CL_CHECK((backend_ctx->kernel_convert_block_q4_k_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q4_k_trans4_ns", &err), err)); @@ -1485,6 +1497,74 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { GGML_LOG_CONT("."); } + // mul_mv_q5_0_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_q5_0_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_q5_0_f32.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_mul_mv_q5_0_f32 = clCreateKernel(prog, "kernel_mul_mv_q5_0_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // mul_mv_q5_0_f32_flat + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_q5_0_f32_flat.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_q5_0_f32_flat.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_mul_mv_q5_0_f32_flat = clCreateKernel(prog, "kernel_mul_mv_q5_0_f32_flat", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // mul_mv_q5_1_f32 + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_q5_1_f32.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_q5_1_f32.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_mul_mv_q5_1_f32 = clCreateKernel(prog, "kernel_mul_mv_q5_1_f32", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + + // mul_mv_q5_1_f32_flat + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mv_q5_1_f32_flat.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mv_q5_1_f32_flat.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_mul_mv_q5_1_f32_flat = clCreateKernel(prog, "kernel_mul_mv_q5_1_f32_flat", &err), err)); + CL_CHECK(clReleaseProgram(prog)); + GGML_LOG_CONT("."); + } + // mul_mv_q5_k_f32 { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -1835,6 +1915,38 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { GGML_LOG_CONT("."); } + // mul_mm_q5_0_f32_l4_lm + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mm_q5_0_f32_l4_lm.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mm_q5_0_f32_l4_lm.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_mul_mm_q5_0_f32_l4_lm = clCreateKernel(prog, "kernel_mul_mm_q5_0_f32_l4_lm", &err), err)); + GGML_LOG_CONT("."); + } + + // mul_mm_q5_1_f32_l4_lm + { +#ifdef GGML_OPENCL_EMBED_KERNELS + const std::string kernel_src { + #include "mul_mm_q5_1_f32_l4_lm.cl.h" + }; +#else + const std::string kernel_src = read_file("mul_mm_q5_1_f32_l4_lm.cl"); +#endif + cl_program prog = + build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts); + + CL_CHECK((backend_ctx->kernel_mul_mm_q5_1_f32_l4_lm = clCreateKernel(prog, "kernel_mul_mm_q5_1_f32_l4_lm", &err), err)); + GGML_LOG_CONT("."); + } + // mul_mm_q8_0_f32_l4_lm { #ifdef GGML_OPENCL_EMBED_KERNELS @@ -5027,6 +5139,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te } else if (op->src[0]->type == GGML_TYPE_F32) { return op->src[1]->type == GGML_TYPE_F32; } else if (op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 || + op->src[0]->type == GGML_TYPE_Q5_0 || op->src[0]->type == GGML_TYPE_Q5_1 || op->src[0]->type == GGML_TYPE_MXFP4 || op->src[0]->type == GGML_TYPE_IQ4_NL || op->src[0]->type == GGML_TYPE_Q4_K || @@ -5977,7 +6090,24 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, return; } #endif // GGML_OPENCL_USE_ADRENO_KERNELS - return; + cl_kernel kernel = backend_ctx->kernel_convert_block_q5_0; + cl_ulong n_blk = ggml_nelements(tensor)/ggml_blck_size(tensor->type); + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qs)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->qh)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->d)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &n_blk)); + + size_t global_work_size[] = {(size_t)CEIL_DIV(n_blk, 64) * 64, 1, 1}; + size_t local_work_size[] = {64, 1, 1}; + + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseMemObject(data_device)); + + tensor->extra = extra; + return; } if (tensor->type == GGML_TYPE_Q5_1) { ggml_tensor_extra_cl * extra_orig = (ggml_tensor_extra_cl *)tensor->extra; @@ -6078,6 +6208,24 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, return; } #endif // GGML_OPENCL_USE_ADRENO_KERNELS + cl_kernel kernel = backend_ctx->kernel_convert_block_q5_1; + cl_ulong n_blk = ggml_nelements(tensor)/ggml_blck_size(tensor->type); + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qs)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->qh)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->d)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra->m)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &n_blk)); + + size_t global_work_size[] = {(size_t)CEIL_DIV(n_blk, 64) * 64, 1, 1}; + size_t local_work_size[] = {64, 1, 1}; + + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clReleaseMemObject(data_device)); + + tensor->extra = extra; return; } if (tensor->type == GGML_TYPE_MXFP4) { @@ -7135,8 +7283,29 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, return; } #endif // GGML_OPENCL_USE_ADRENO_KERNELS - // TODO: normal q5_0 - (void) extra; + + cl_int err; + cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE, + ggml_nbytes(tensor), NULL, &err); + CL_CHECK(err); + + cl_kernel kernel = backend_ctx->kernel_restore_block_q5_0; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &data_device)); + + size_t global_work_size[] = {(size_t)ggml_nelements(tensor)/ggml_blck_size(tensor->type), 1, 1}; + size_t local_work_size[] = {1, 1, 1}; + + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, + global_work_size, local_work_size, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clEnqueueReadBuffer( + queue, data_device, CL_TRUE, offset, + size, data, 0, NULL, NULL)); + CL_CHECK(clReleaseMemObject(data_device)); return; } if (tensor->type == GGML_TYPE_Q5_1) { @@ -7177,8 +7346,29 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, return; } #endif // GGML_OPENCL_USE_ADRENO_KERNELS - // TODO: normal q5_1 - (void) extra; + cl_int err; + cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE, + ggml_nbytes(tensor), NULL, &err); + CL_CHECK(err); + + cl_kernel kernel = backend_ctx->kernel_restore_block_q5_1; + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->m)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &data_device)); + + size_t global_work_size[] = {(size_t)ggml_nelements(tensor)/ggml_blck_size(tensor->type), 1, 1}; + size_t local_work_size[] = {1, 1, 1}; + + cl_event evt; + CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, + global_work_size, local_work_size, 0, NULL, &evt)); + CL_CHECK(clWaitForEvents(1, &evt)); + CL_CHECK(clEnqueueReadBuffer( + queue, data_device, CL_TRUE, offset, + size, data, 0, NULL, NULL)); + CL_CHECK(clReleaseMemObject(data_device)); return; } if (tensor->type == GGML_TYPE_MXFP4) { @@ -12936,6 +13126,8 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co #ifdef GGML_OPENCL_SOA_Q ggml_tensor_extra_cl_q4_0 * extra0_q4_0 = (ggml_tensor_extra_cl_q4_0 *)src0->extra; ggml_tensor_extra_cl_q4_1 * extra0_q4_1 = (ggml_tensor_extra_cl_q4_1 *)src0->extra; + ggml_tensor_extra_cl_q5_0 * extra0_q5_0 = (ggml_tensor_extra_cl_q5_0 *)src0->extra; + ggml_tensor_extra_cl_q5_1 * extra0_q5_1 = (ggml_tensor_extra_cl_q5_1 *)src0->extra; ggml_tensor_extra_cl_mxfp4 * extra0_mxfp4 = (ggml_tensor_extra_cl_mxfp4 *)src0->extra; ggml_tensor_extra_cl_q8_0 * extra0_q8_0 = (ggml_tensor_extra_cl_q8_0 *)src0->extra; ggml_tensor_extra_cl_iq4_nl * extra0_iq4_nl = (ggml_tensor_extra_cl_iq4_nl *)src0->extra; @@ -13271,6 +13463,93 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); return; } + case GGML_TYPE_Q5_0: { + if (ne11 < 32) { + break; + } + if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) { + break; + } + + kernel = backend_ctx->kernel_mul_mm_q5_0_f32_l4_lm; + nth0 = 128; // calculated as (BM*BN)/(TM*TN) + + int batch_stride_a = ne00*ne01; + int batch_stride_b = ne10*ne11; + int batch_stride_d = ne0*ne1; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_0->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_0->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_0->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne11)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne10)); // stride_a + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne10)); // stride_b + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne01)); // stride_d + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &batch_stride_a)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &batch_stride_b)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &batch_stride_d)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &r3)); + + // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed. + size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)(CEIL_DIV(ne11, 64)), (size_t)ne12*ne13}; + size_t local_work_size[] = {(size_t)nth0, 1, 1}; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + return; + } + case GGML_TYPE_Q5_1: { + if (ne11 < 32) { + break; + } + if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) { + break; + } + + kernel = backend_ctx->kernel_mul_mm_q5_1_f32_l4_lm; + nth0 = 128; // calculated as (BM*BN)/(TM*TN) + + int batch_stride_a = ne00*ne01; + int batch_stride_b = ne10*ne11; + int batch_stride_d = ne0*ne1; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_1->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_1->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_1->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_1->m)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne11)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne10)); // stride_a + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne10)); // stride_b + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne01)); // stride_d + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &batch_stride_a)); + CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &batch_stride_b)); + CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &batch_stride_d)); + CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &r3)); + + // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed. + size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)(CEIL_DIV(ne11, 64)), (size_t)ne12*ne13}; + size_t local_work_size[] = {(size_t)nth0, 1, 1}; + + backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst); + return; + } case GGML_TYPE_Q8_0: { if (ne11 < 32) { break; @@ -13807,6 +14086,137 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co #endif // GGML_OPENCL_SOA_Q break; } + case GGML_TYPE_Q5_0: { +#ifdef GGML_OPENCL_SOA_Q + if (backend_ctx->gpu_family == INTEL) { + nth0 = 16; + nth1 = 1; + ndst = 4; + } else if (backend_ctx->gpu_family == ADRENO) { + nth0 = 64; + nth1 = 1; + ndst = 4; + } else { + GGML_ASSERT(false && "TODO: Unknown GPU"); + } + + kernel = backend_ctx->kernel_mul_mv_q5_0_f32_flat; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_0->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_0->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_0->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne10)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &r3)); +#else + if (backend_ctx->gpu_family == INTEL) { + nth0 = 16; + nth1 = 1; + ndst = 4; + } else if (backend_ctx->gpu_family == ADRENO) { + nth0 = 64; + nth1 = 1; + ndst = 4; + } else { + GGML_ASSERT(false && "TODO: Unknown GPU"); + } + + kernel = backend_ctx->kernel_mul_mv_q5_0_f32; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne10)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &r3)); +#endif // GGML_OPENCL_SOA_Q + break; + } + case GGML_TYPE_Q5_1: { +#ifdef GGML_OPENCL_SOA_Q + if (backend_ctx->gpu_family == INTEL) { + nth0 = 16; + nth1 = 1; + ndst = 4; + } else if (backend_ctx->gpu_family == ADRENO) { + nth0 = 64; + nth1 = 1; + ndst = 4; + } else { + GGML_ASSERT(false && "TODO: Unknown GPU"); + } + + kernel = backend_ctx->kernel_mul_mv_q5_1_f32_flat; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_1->qs)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_1->qh)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_1->d)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_1->m)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne10)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &r3)); +#else + if (backend_ctx->gpu_family == INTEL) { + nth0 = 16; + nth1 = 1; + ndst = 4; + } else if (backend_ctx->gpu_family == ADRENO) { + nth0 = 64; + nth1 = 1; + ndst = 4; + } else { + GGML_ASSERT(false && "TODO: Unknown GPU"); + } + + kernel = backend_ctx->kernel_mul_mv_q5_1_f32; + + CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device)); + CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0)); + CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra1->data_device)); + CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset1)); + CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extrad->data_device)); + CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offsetd)); + CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne00)); + CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne01)); + CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne02)); + CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne10)); + CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne12)); + CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne0)); + CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne1)); + CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &r2)); + CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &r3)); +#endif // GGML_OPENCL_SOA_Q + break; + } case GGML_TYPE_Q8_0: { #ifdef GGML_OPENCL_SOA_Q kernel = backend_ctx->kernel_mul_mv_q8_0_f32_flat; @@ -14247,6 +14657,8 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co if (src0t == GGML_TYPE_Q4_0 || src0t == GGML_TYPE_MXFP4 || src0t == GGML_TYPE_Q4_1 || + src0t == GGML_TYPE_Q5_0 || + src0t == GGML_TYPE_Q5_1 || src0t == GGML_TYPE_Q8_0 || src0t == GGML_TYPE_IQ4_NL || src0t == GGML_TYPE_Q2_K) { diff --git a/ggml/src/ggml-opencl/kernels/cvt.cl b/ggml/src/ggml-opencl/kernels/cvt.cl index 4f01887ef..d07f0a1a0 100644 --- a/ggml/src/ggml-opencl/kernels/cvt.cl +++ b/ggml/src/ggml-opencl/kernels/cvt.cl @@ -537,6 +537,53 @@ kernel void kernel_restore_block_q4_1_trans4_ns( ((__global ushort8 *)(&(b->qs[0])))[0] = pre_block; } +//------------------------------------------------------------------------------ +// kernel_convert_block_q5_0 +// Convert the block_q5_0 format to 3 separate arrays (AOS -> SOA). +// This kernel does not deshuffle the bits. +//------------------------------------------------------------------------------ +kernel void kernel_convert_block_q5_0( + global struct block_q5_0 * src0, + global uchar * dst_qs, + global uint * dst_qh, + global half * dst_d, + ulong n_blk +) { + if (get_global_id(0) >= n_blk) { + return; + } + + global struct block_q5_0 * b = (global struct block_q5_0 *) src0 + get_global_id(0); + global uchar * qs = (global uchar *) dst_qs + (QK5_0/2)*get_global_id(0); + global uint * qh = (global uint *) dst_qh + get_global_id(0); + global half * d = (global half *) dst_d + get_global_id(0); + + *d = b->d; + *qh = *((global uint *)(b->qh)); + + for (int i = 0; i < QK5_0/2; ++i) { + qs[i] = b->qs[i]; + } +} + +kernel void kernel_restore_block_q5_0( + global uchar * src_qs, + global uint * src_qh, + global half * src_d, + global struct block_q5_0 * dst +) { + global struct block_q5_0 * b = (global struct block_q5_0 *) dst + get_global_id(0); + global uchar * qs = (global uchar *) src_qs + (QK5_0/2)*get_global_id(0); + global uint * qh = (global uint *) src_qh + get_global_id(0); + global half * d = (global half *) src_d + get_global_id(0); + + b->d = *d; + *((global uint *)(b->qh)) = *qh; + for (int i = 0; i < QK5_0/2; ++i) { + b->qs[i] = qs[i]; + } +} + kernel void kernel_convert_block_q5_0_trans4_ns( __global struct block_q5_0 * src0, __global uint * dst_qs, @@ -636,6 +683,59 @@ kernel void kernel_restore_block_q5_0_trans4_ns( ((__global ushort8 *)(&(b->qs[0])))[0] = pre_block; } +//------------------------------------------------------------------------------ +// kernel_convert_block_q5_1 +// Convert the block_q5_1 format to 4 separate arrays (AOS -> SOA). +// This kernel does not deshuffle the bits. +//------------------------------------------------------------------------------ +kernel void kernel_convert_block_q5_1( + global struct block_q5_1 * src0, + global uchar * dst_qs, + global uint * dst_qh, + global half * dst_d, + global half * dst_m, + ulong n_blk +) { + if (get_global_id(0) >= n_blk) { + return; + } + + global struct block_q5_1 * b = (global struct block_q5_1 *) src0 + get_global_id(0); + global uchar * qs = (global uchar *) dst_qs + (QK5_1/2)*get_global_id(0); + global uint * qh = (global uint *) dst_qh + get_global_id(0); + global half * d = (global half *) dst_d + get_global_id(0); + global half * m = (global half *) dst_m + get_global_id(0); + + *d = b->d; + *m = b->m; + *qh = *((global uint *)(b->qh)); + + for (int i = 0; i < QK5_1/2; ++i) { + qs[i] = b->qs[i]; + } +} + +kernel void kernel_restore_block_q5_1( + global uchar * src_qs, + global uint * src_qh, + global half * src_d, + global half * src_m, + global struct block_q5_1 * dst +) { + global struct block_q5_1 * b = (global struct block_q5_1 *) dst + get_global_id(0); + global uchar * qs = (global uchar *) src_qs + (QK5_1/2)*get_global_id(0); + global uint * qh = (global uint *) src_qh + get_global_id(0); + global half * d = (global half *) src_d + get_global_id(0); + global half * m = (global half *) src_m + get_global_id(0); + + b->d = *d; + b->m = *m; + *((global uint *)(b->qh)) = *qh; + for (int i = 0; i < QK5_1/2; ++i) { + b->qs[i] = qs[i]; + } +} + kernel void kernel_convert_block_q5_1_trans4_ns( __global struct block_q5_1 * src0, __global uint * dst_qs, diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_q5_0_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_q5_0_f32_l4_lm.cl new file mode 100644 index 000000000..1e980a478 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mm_q5_0_f32_l4_lm.cl @@ -0,0 +1,173 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#define LOAD_VEC_A 8 +#define LOAD_VEC_B 4 + +#define BM 64 +#define BN 64 +#define BK 32 +#define TM 4 +#define TN 8 + +kernel void kernel_mul_mm_q5_0_f32_l4_lm( + global uchar4 * src0_qs, + global uint * src0_qh, + global half * src0_d, + global float4 * src1, + ulong offset1, + global float * dst, + ulong offsetd, + + int ne00, + int ne01, + int ne02, + int ne11, + int ne12, + + int stride_a, + int stride_b, + int stride_d, + + int batch_stride_a, + int batch_stride_b, + int batch_stride_d, + + int r2, + int r3 +) { + src1 = (global float4*)((global char*)src1 + offset1); + dst = (global float *)((global char*)dst + offsetd); + + local float buf_a[BM * BK]; + local float buf_b[BN * BK]; + + const int batch_idx = get_global_id(2); + + const int i13 = batch_idx / ne12; + const int i12 = batch_idx % ne12; + + const int i03 = i13 / r3; + const int i02 = i12 / r2; + + const int batch_idx_a = i03 * ne02 + i02; + + const int ir = get_group_id(0); + const int ic = get_group_id(1); + + const int tid = get_local_id(0); + const int th_r = tid % (BM / TM); + const int th_c = tid / (BM / TM); + + const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A); + const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A); + const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B); + const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B); + + const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK; + const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK; + + int pos_a = (batch_idx_a * batch_stride_a + ir * BM * stride_a) / LOAD_VEC_A; + int pos_b = (batch_idx * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B; + + float sums[TM * TN]; + float cache_a[TM]; + float cache_b[TN]; + + for (int i = 0; i < TM * TN; i++) { + sums[i] = 0.0f; + } + + for (int block = 0; block < ne00; block += BK) { + for (int l = 0; l < BM; l += loadstride_a) { + if (ir*BM + loadc_a + l < ne01) { + int idx = pos_a + (loadc_a + l) * stride_a / LOAD_VEC_A + loadr_a; + int ib = idx / 4; + int iqs = idx % 4; + + float d = (float)src0_d[ib]; + uint qh_val = src0_qh[ib]; + + global uchar4 * qs_ptr = src0_qs + ib*4 + iqs; + uchar4 q = *qs_ptr; + + uint qh_lo = qh_val >> (iqs * 4); + uint qh_hi = qh_val >> (iqs * 4 + 16); + + uchar4 b_lo = (uchar4)((uchar)qh_lo, (uchar)(qh_lo >> 1), (uchar)(qh_lo >> 2), (uchar)(qh_lo >> 3)) & (uchar)1; + uchar4 b_hi = (uchar4)((uchar)qh_hi, (uchar)(qh_hi >> 1), (uchar)(qh_hi >> 2), (uchar)(qh_hi >> 3)) & (uchar)1; + + float4 v1 = (convert_float4((q & (uchar)0x0F) | (b_lo << (uchar)4)) - 16.0f) * d; + float4 v2 = (convert_float4((q >> (uchar)4) | (b_hi << (uchar)4)) - 16.0f) * d; + + buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = v1.s0; + buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = v1.s1; + buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = v1.s2; + buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = v1.s3; + buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = v2.s0; + buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = v2.s1; + buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = v2.s2; + buf_a[(loadr_a * 4 + 19) * BM + loadc_a + l] = v2.s3; + } else { + buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 19) * BM + loadc_a + l] = 0.0f; + } + } + + for (int l = 0; l < BN; l += loadstride_b) { + if (ic*BN + loadc_b + l < ne11) { + int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b; + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3; + } else { + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f; + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + + pos_a += BK / LOAD_VEC_A; + pos_b += BK / LOAD_VEC_B; + + for (int i = 0; i < BK; i++) { + for (int j = 0; j < TM; j++) { + cache_a[j] = buf_a[(i) * BM + th_r * TM + j]; + } + + for (int j = 0; j < TN; j++) { + cache_b[j] = buf_b[(i) * BN + th_c * TN + j]; + } + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + const int sums_idx = cc*TM + cr; + sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]); + } + } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + const int dr = ir * BM + th_r * TM; + const int dc = ic * BN + th_c * TN; + + const int offsets = batch_idx * batch_stride_d; + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + if (dr + cr < ne01 && dc + cc < ne11) { + dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr]; + } + } + } +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_q5_1_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_q5_1_f32_l4_lm.cl new file mode 100644 index 000000000..ba06be546 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mm_q5_1_f32_l4_lm.cl @@ -0,0 +1,175 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#define LOAD_VEC_A 8 +#define LOAD_VEC_B 4 + +#define BM 64 +#define BN 64 +#define BK 32 +#define TM 4 +#define TN 8 + +kernel void kernel_mul_mm_q5_1_f32_l4_lm( + global uchar4 * src0_qs, + global uint * src0_qh, + global half * src0_d, + global half * src0_m, + global float4 * src1, + ulong offset1, + global float * dst, + ulong offsetd, + + int ne00, + int ne01, + int ne02, + int ne11, + int ne12, + + int stride_a, + int stride_b, + int stride_d, + + int batch_stride_a, + int batch_stride_b, + int batch_stride_d, + + int r2, + int r3 +) { + src1 = (global float4*)((global char*)src1 + offset1); + dst = (global float *)((global char*)dst + offsetd); + + local float buf_a[BM * BK]; + local float buf_b[BN * BK]; + + const int batch_idx = get_global_id(2); + + const int i13 = batch_idx / ne12; + const int i12 = batch_idx % ne12; + + const int i03 = i13 / r3; + const int i02 = i12 / r2; + + const int batch_idx_a = i03 * ne02 + i02; + + const int ir = get_group_id(0); + const int ic = get_group_id(1); + + const int tid = get_local_id(0); + const int th_r = tid % (BM / TM); + const int th_c = tid / (BM / TM); + + const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A); + const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A); + const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B); + const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B); + + const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK; + const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK; + + int pos_a = (batch_idx_a * batch_stride_a + ir * BM * stride_a) / LOAD_VEC_A; + int pos_b = (batch_idx * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B; + + float sums[TM * TN]; + float cache_a[TM]; + float cache_b[TN]; + + for (int i = 0; i < TM * TN; i++) { + sums[i] = 0.0f; + } + + for (int block = 0; block < ne00; block += BK) { + for (int l = 0; l < BM; l += loadstride_a) { + if (ir*BM + loadc_a + l < ne01) { + int idx = pos_a + (loadc_a + l) * stride_a / LOAD_VEC_A + loadr_a; + int ib = idx / 4; + int iqs = idx % 4; + + float d = (float)src0_d[ib]; + float m = (float)src0_m[ib]; + uint qh_val = src0_qh[ib]; + + global uchar4 * qs = src0_qs + ib*4 + iqs; + uchar4 q = *qs; + + uint qh_lo = qh_val >> (iqs * 4); + uint qh_hi = qh_val >> (iqs * 4 + 16); + + uchar4 b_lo = (uchar4)((uchar)qh_lo, (uchar)(qh_lo >> 1), (uchar)(qh_lo >> 2), (uchar)(qh_lo >> 3)) & (uchar)1; + uchar4 b_hi = (uchar4)((uchar)qh_hi, (uchar)(qh_hi >> 1), (uchar)(qh_hi >> 2), (uchar)(qh_hi >> 3)) & (uchar)1; + + float4 v1 = convert_float4((q & (uchar)0x0F) | (b_lo << (uchar)4)) * d + m; + float4 v2 = convert_float4((q >> (uchar)4) | (b_hi << (uchar)4)) * d + m; + + buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = v1.s0; + buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = v1.s1; + buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = v1.s2; + buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = v1.s3; + buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = v2.s0; + buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = v2.s1; + buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = v2.s2; + buf_a[(loadr_a * 4 + 19) * BM + loadc_a + l] = v2.s3; + } else { + buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = 0.0f; + buf_a[(loadr_a * 4 + 19) * BM + loadc_a + l] = 0.0f; + } + } + + for (int l = 0; l < BN; l += loadstride_b) { + if (ic*BN + loadc_b + l < ne11) { + int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b; + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3; + } else { + buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f; + buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f; + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + + pos_a += BK / LOAD_VEC_A; + pos_b += BK / LOAD_VEC_B; + + for (int i = 0; i < BK; i++) { + for (int j = 0; j < TM; j++) { + cache_a[j] = buf_a[(i) * BM + th_r * TM + j]; + } + + for (int j = 0; j < TN; j++) { + cache_b[j] = buf_b[(i) * BN + th_c * TN + j]; + } + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + const int sums_idx = cc*TM + cr; + sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]); + } + } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + const int dr = ir * BM + th_r * TM; + const int dc = ic * BN + th_c * TN; + + const int offsets = batch_idx * batch_stride_d; + + for (int cc = 0; cc < TN; cc++) { + for (int cr = 0; cr < TM; cr++) { + if (dr + cr < ne01 && dc + cc < ne11) { + dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr]; + } + } + } +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32.cl new file mode 100644 index 000000000..6d8c9e8f0 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32.cl @@ -0,0 +1,241 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_subgroups +#pragma OPENCL EXTENSION cl_intel_subgroups : enable +#else +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#endif + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK5_0 32 + +struct block_q5_0 { + half d; + uchar qh[4]; + uchar qs[QK5_0 / 2]; +}; + +inline float block_q5_0_dot_y( + global const struct block_q5_0 * qb_curr, + float sumy, + float16 yl, + int il, + global const float * yb +) { + float d = qb_curr->d; + + float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f); + + global const ushort * qs = ((global const ushort *)((global const uchar *) qb_curr + 6 + il)); + + acc.s0 += yl.s0 * (qs[0] & 0x000F); + acc.s0 += yl.s1 * (qs[0] & 0x0F00); + acc.s0 += yl.s8 * (qs[0] & 0x00F0); + acc.s3 += yl.s9 * (qs[0] & 0xF000); + + acc.s0 += yl.s2 * (qs[1] & 0x000F); + acc.s1 += yl.s3 * (qs[1] & 0x0F00); + acc.s2 += yl.sa * (qs[1] & 0x00F0); + acc.s3 += yl.sb * (qs[1] & 0xF000); + + acc.s0 += yl.s4 * (qs[2] & 0x000F); + acc.s1 += yl.s5 * (qs[2] & 0x0F00); + acc.s2 += yl.sc * (qs[2] & 0x00F0); + acc.s3 += yl.sd * (qs[2] & 0xF000); + + acc.s0 += yl.s6 * (qs[3] & 0x000F); + acc.s1 += yl.s7 * (qs[3] & 0x0F00); + acc.s2 += yl.se * (qs[3] & 0x00F0); + acc.s3 += yl.sf * (qs[3] & 0xF000); + + uint qh_val = *((global const uint *)((global const uchar *) qb_curr + 2)); + uchar qh_lo = (uchar)((qh_val >> il) & 0xFF); + uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF); + + float qh_sum = 0.0f; + qh_sum += yb[0] * (float)((qh_lo >> 0) & 1); + qh_sum += yb[1] * (float)((qh_lo >> 1) & 1); + qh_sum += yb[2] * (float)((qh_lo >> 2) & 1); + qh_sum += yb[3] * (float)((qh_lo >> 3) & 1); + qh_sum += yb[4] * (float)((qh_lo >> 4) & 1); + qh_sum += yb[5] * (float)((qh_lo >> 5) & 1); + qh_sum += yb[6] * (float)((qh_lo >> 6) & 1); + qh_sum += yb[7] * (float)((qh_lo >> 7) & 1); + qh_sum += yb[16] * (float)((qh_hi >> 0) & 1); + qh_sum += yb[17] * (float)((qh_hi >> 1) & 1); + qh_sum += yb[18] * (float)((qh_hi >> 2) & 1); + qh_sum += yb[19] * (float)((qh_hi >> 3) & 1); + qh_sum += yb[20] * (float)((qh_hi >> 4) & 1); + qh_sum += yb[21] * (float)((qh_hi >> 5) & 1); + qh_sum += yb[22] * (float)((qh_hi >> 6) & 1); + qh_sum += yb[23] * (float)((qh_hi >> 7) & 1); + + return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum - 16.0f * sumy); +} + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 // each subgroup works on 4 rows +#define N_SIMDGROUP 1 // number of subgroups in a thread group +#define N_SIMDWIDTH 16 // assuming subgroup size is 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +inline void mul_vec_q_n_f32( + global void * src0, + global float * src1, + global float * dst, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + const ulong nb = ne00/QK5_0; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02); + + global struct block_q5_0 * x = (global struct block_q5_0 *) src0 + offset0; + global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1; + + float16 yl; + float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f); + + int ix = get_sub_group_local_id()/2; + int il = 8*(get_sub_group_local_id()%2); + + global float * yb = y + ix * QK5_0 + il; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) { + float sumy = 0; + + sumy += yb[0]; + sumy += yb[1]; + sumy += yb[2]; + sumy += yb[3]; + sumy += yb[4]; + sumy += yb[5]; + sumy += yb[6]; + sumy += yb[7]; + + sumy += yb[16]; + sumy += yb[17]; + sumy += yb[18]; + sumy += yb[19]; + sumy += yb[20]; + sumy += yb[21]; + sumy += yb[22]; + sumy += yb[23]; + + + yl.s0 = yb[0]; + yl.s1 = yb[1]/256.f; + + yl.s2 = yb[2]; + yl.s3 = yb[3]/256.f; + + yl.s4 = yb[4]; + yl.s5 = yb[5]/256.f; + + yl.s6 = yb[6]; + yl.s7 = yb[7]/256.f; + + yl.s8 = yb[16]/16.f; + yl.s9 = yb[17]/4096.f; + + yl.sa = yb[18]/16.f; + yl.sb = yb[19]/4096.f; + + yl.sc = yb[20]/16.f; + yl.sd = yb[21]/4096.f; + + yl.se = yb[22]/16.f; + yl.sf = yb[23]/4096.f; + + sumf.s0 += block_q5_0_dot_y(x+ib+0*nb, sumy, yl, il, yb); + sumf.s1 += block_q5_0_dot_y(x+ib+1*nb, sumy, yl, il, yb); + sumf.s2 += block_q5_0_dot_y(x+ib+2*nb, sumy, yl, il, yb); + sumf.s3 += block_q5_0_dot_y(x+ib+3*nb, sumy, yl, il, yb); + + yb += QK5_0 * (N_SIMDWIDTH/2); + } + + float4 tot = (float4)( + sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1), + sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3) + ); + + if (get_sub_group_local_id() == 0) { + if (first_row + 0 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0; + } + if (first_row + 1 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1; + } + if (first_row + 2 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2; + } + if (first_row + 3 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3; + } + } +} + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_q5_0_f32( + global void * src0, + ulong offset0, + global float * src1, + ulong offset1, + global float * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + src0 = (global void*)((global char*)src0 + offset0); + src1 = (global float*)((global char*)src1 + offset1); + dst = (global float*)((global char*)dst + offsetd); + + mul_vec_q_n_f32(src0, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3); +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32_flat.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32_flat.cl new file mode 100644 index 000000000..34ec133d3 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_q5_0_f32_flat.cl @@ -0,0 +1,243 @@ + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_subgroups +#pragma OPENCL EXTENSION cl_intel_subgroups : enable +#else +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#endif + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK5_0 32 + +inline float block_q5_0_dot_y_flat( + global const uchar * x, + global const uint * qh_ptr, + global const half * dh, + float sumy, + float16 yl, + int il, + global const float * yb +) { + float d = *dh; + global const ushort * qs = ((global const ushort *)(x + il)); + + float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f); + + acc.s0 += yl.s0 * (qs[0] & 0x000F); + acc.s0 += yl.s1 * (qs[0] & 0x0F00); + acc.s0 += yl.s8 * (qs[0] & 0x00F0); + acc.s3 += yl.s9 * (qs[0] & 0xF000); + + acc.s0 += yl.s2 * (qs[1] & 0x000F); + acc.s1 += yl.s3 * (qs[1] & 0x0F00); + acc.s2 += yl.sa * (qs[1] & 0x00F0); + acc.s3 += yl.sb * (qs[1] & 0xF000); + + acc.s0 += yl.s4 * (qs[2] & 0x000F); + acc.s1 += yl.s5 * (qs[2] & 0x0F00); + acc.s2 += yl.sc * (qs[2] & 0x00F0); + acc.s3 += yl.sd * (qs[2] & 0xF000); + + acc.s0 += yl.s6 * (qs[3] & 0x000F); + acc.s1 += yl.s7 * (qs[3] & 0x0F00); + acc.s2 += yl.se * (qs[3] & 0x00F0); + acc.s3 += yl.sf * (qs[3] & 0xF000); + + uint qh_val = *qh_ptr; + uchar qh_lo = (uchar)((qh_val >> il) & 0xFF); + uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF); + + float qh_sum = 0.0f; + qh_sum += yb[0] * (float)((qh_lo >> 0) & 1); + qh_sum += yb[1] * (float)((qh_lo >> 1) & 1); + qh_sum += yb[2] * (float)((qh_lo >> 2) & 1); + qh_sum += yb[3] * (float)((qh_lo >> 3) & 1); + qh_sum += yb[4] * (float)((qh_lo >> 4) & 1); + qh_sum += yb[5] * (float)((qh_lo >> 5) & 1); + qh_sum += yb[6] * (float)((qh_lo >> 6) & 1); + qh_sum += yb[7] * (float)((qh_lo >> 7) & 1); + qh_sum += yb[16] * (float)((qh_hi >> 0) & 1); + qh_sum += yb[17] * (float)((qh_hi >> 1) & 1); + qh_sum += yb[18] * (float)((qh_hi >> 2) & 1); + qh_sum += yb[19] * (float)((qh_hi >> 3) & 1); + qh_sum += yb[20] * (float)((qh_hi >> 4) & 1); + qh_sum += yb[21] * (float)((qh_hi >> 5) & 1); + qh_sum += yb[22] * (float)((qh_hi >> 6) & 1); + qh_sum += yb[23] * (float)((qh_hi >> 7) & 1); + + return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum - 16.0f * sumy); +} + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 // each subgroup works on 4 rows +#define N_SIMDGROUP 1 // number of subgroups in a thread group +#define N_SIMDWIDTH 16 // assuming subgroup size is 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +inline void mul_vec_q_n_f32_flat( + global void * src0_qs, + global void * src0_qh, + global void * src0_d, + global float * src1, + global float * dst, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + const ulong nb = ne00/QK5_0; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02); + + ulong offset0_qs = offset0 * (QK5_0/2); + + global uchar * x = (global uchar *) src0_qs + offset0_qs; + global uint * qh = (global uint *) src0_qh + offset0; + global half * d = (global half *) src0_d + offset0; + global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1; + + float16 yl; + float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f); + + int ix = get_sub_group_local_id()/2; + int il = 8*(get_sub_group_local_id()%2); + + global float * yb = y + ix * QK5_0 + il; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) { + float sumy = 0; + + sumy += yb[0]; + sumy += yb[1]; + sumy += yb[2]; + sumy += yb[3]; + sumy += yb[4]; + sumy += yb[5]; + sumy += yb[6]; + sumy += yb[7]; + + sumy += yb[16]; + sumy += yb[17]; + sumy += yb[18]; + sumy += yb[19]; + sumy += yb[20]; + sumy += yb[21]; + sumy += yb[22]; + sumy += yb[23]; + + + yl.s0 = yb[0]; + yl.s1 = yb[1]/256.f; + + yl.s2 = yb[2]; + yl.s3 = yb[3]/256.f; + + yl.s4 = yb[4]; + yl.s5 = yb[5]/256.f; + + yl.s6 = yb[6]; + yl.s7 = yb[7]/256.f; + + yl.s8 = yb[16]/16.f; + yl.s9 = yb[17]/4096.f; + + yl.sa = yb[18]/16.f; + yl.sb = yb[19]/4096.f; + + yl.sc = yb[20]/16.f; + yl.sd = yb[21]/4096.f; + + yl.se = yb[22]/16.f; + yl.sf = yb[23]/4096.f; + + sumf.s0 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 0*nb*(QK5_0/2), qh + ib + 0*nb, d + ib + 0*nb, sumy, yl, il, yb); + sumf.s1 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 1*nb*(QK5_0/2), qh + ib + 1*nb, d + ib + 1*nb, sumy, yl, il, yb); + sumf.s2 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 2*nb*(QK5_0/2), qh + ib + 2*nb, d + ib + 2*nb, sumy, yl, il, yb); + sumf.s3 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 3*nb*(QK5_0/2), qh + ib + 3*nb, d + ib + 3*nb, sumy, yl, il, yb); + + yb += QK5_0 * (N_SIMDWIDTH/2); + } + + float4 tot = (float4)( + sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1), + sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3) + ); + + if (get_sub_group_local_id() == 0) { + if (first_row + 0 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0; + } + if (first_row + 1 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1; + } + if (first_row + 2 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2; + } + if (first_row + 3 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3; + } + } +} + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_q5_0_f32_flat( + global void * src0_qs, + global void * src0_qh, + global void * src0_d, + global float * src1, + ulong offset1, + global float * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + src1 = (global float*)((global char*)src1 + offset1); + dst = (global float*)((global char*)dst + offsetd); + + mul_vec_q_n_f32_flat(src0_qs, src0_qh, src0_d, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3); +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32.cl new file mode 100644 index 000000000..1480f6750 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32.cl @@ -0,0 +1,243 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_subgroups +#pragma OPENCL EXTENSION cl_intel_subgroups : enable +#else +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#endif + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK5_1 32 + +struct block_q5_1 { + half d; + half m; + uchar qh[4]; + uchar qs[QK5_1 / 2]; +}; + +inline float block_q5_1_dot_y( + global const struct block_q5_1 * qb_curr, + float sumy, + float16 yl, + int il, + global const float * yb +) { + float d = qb_curr->d; + float m = qb_curr->m; + + float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f); + + global const ushort * qs = ((global const ushort *)((global const uchar *) qb_curr + 8 + il)); + + acc.s0 += yl.s0 * (qs[0] & 0x000F); + acc.s0 += yl.s1 * (qs[0] & 0x0F00); + acc.s0 += yl.s8 * (qs[0] & 0x00F0); + acc.s3 += yl.s9 * (qs[0] & 0xF000); + + acc.s0 += yl.s2 * (qs[1] & 0x000F); + acc.s1 += yl.s3 * (qs[1] & 0x0F00); + acc.s2 += yl.sa * (qs[1] & 0x00F0); + acc.s3 += yl.sb * (qs[1] & 0xF000); + + acc.s0 += yl.s4 * (qs[2] & 0x000F); + acc.s1 += yl.s5 * (qs[2] & 0x0F00); + acc.s2 += yl.sc * (qs[2] & 0x00F0); + acc.s3 += yl.sd * (qs[2] & 0xF000); + + acc.s0 += yl.s6 * (qs[3] & 0x000F); + acc.s1 += yl.s7 * (qs[3] & 0x0F00); + acc.s2 += yl.se * (qs[3] & 0x00F0); + acc.s3 += yl.sf * (qs[3] & 0xF000); + + uint qh_val = *((global const uint *)((global const uchar *) qb_curr + 4)); + uchar qh_lo = (uchar)((qh_val >> il) & 0xFF); + uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF); + + float qh_sum = 0.0f; + qh_sum += yb[0] * (float)((qh_lo >> 0) & 1); + qh_sum += yb[1] * (float)((qh_lo >> 1) & 1); + qh_sum += yb[2] * (float)((qh_lo >> 2) & 1); + qh_sum += yb[3] * (float)((qh_lo >> 3) & 1); + qh_sum += yb[4] * (float)((qh_lo >> 4) & 1); + qh_sum += yb[5] * (float)((qh_lo >> 5) & 1); + qh_sum += yb[6] * (float)((qh_lo >> 6) & 1); + qh_sum += yb[7] * (float)((qh_lo >> 7) & 1); + qh_sum += yb[16] * (float)((qh_hi >> 0) & 1); + qh_sum += yb[17] * (float)((qh_hi >> 1) & 1); + qh_sum += yb[18] * (float)((qh_hi >> 2) & 1); + qh_sum += yb[19] * (float)((qh_hi >> 3) & 1); + qh_sum += yb[20] * (float)((qh_hi >> 4) & 1); + qh_sum += yb[21] * (float)((qh_hi >> 5) & 1); + qh_sum += yb[22] * (float)((qh_hi >> 6) & 1); + qh_sum += yb[23] * (float)((qh_hi >> 7) & 1); + + return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum) + sumy * m; +} + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 // each subgroup works on 4 rows +#define N_SIMDGROUP 1 // number of subgroups in a thread group +#define N_SIMDWIDTH 16 // assuming subgroup size is 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +inline void mul_vec_q_n_f32( + global void * src0, + global float * src1, + global float * dst, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + const ulong nb = ne00/QK5_1; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02); + + global struct block_q5_1 * x = (global struct block_q5_1 *) src0 + offset0; + global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1; + + float16 yl; + float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f); + + int ix = get_sub_group_local_id()/2; + int il = 8*(get_sub_group_local_id()%2); + + global float * yb = y + ix * QK5_1 + il; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) { + float sumy = 0; + + sumy += yb[0]; + sumy += yb[1]; + sumy += yb[2]; + sumy += yb[3]; + sumy += yb[4]; + sumy += yb[5]; + sumy += yb[6]; + sumy += yb[7]; + + sumy += yb[16]; + sumy += yb[17]; + sumy += yb[18]; + sumy += yb[19]; + sumy += yb[20]; + sumy += yb[21]; + sumy += yb[22]; + sumy += yb[23]; + + + yl.s0 = yb[0]; + yl.s1 = yb[1]/256.f; + + yl.s2 = yb[2]; + yl.s3 = yb[3]/256.f; + + yl.s4 = yb[4]; + yl.s5 = yb[5]/256.f; + + yl.s6 = yb[6]; + yl.s7 = yb[7]/256.f; + + yl.s8 = yb[16]/16.f; + yl.s9 = yb[17]/4096.f; + + yl.sa = yb[18]/16.f; + yl.sb = yb[19]/4096.f; + + yl.sc = yb[20]/16.f; + yl.sd = yb[21]/4096.f; + + yl.se = yb[22]/16.f; + yl.sf = yb[23]/4096.f; + + sumf.s0 += block_q5_1_dot_y(x+ib+0*nb, sumy, yl, il, yb); + sumf.s1 += block_q5_1_dot_y(x+ib+1*nb, sumy, yl, il, yb); + sumf.s2 += block_q5_1_dot_y(x+ib+2*nb, sumy, yl, il, yb); + sumf.s3 += block_q5_1_dot_y(x+ib+3*nb, sumy, yl, il, yb); + + yb += QK5_1 * (N_SIMDWIDTH/2); + } + + float4 tot = (float4)( + sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1), + sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3) + ); + + if (get_sub_group_local_id() == 0) { + if (first_row + 0 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0; + } + if (first_row + 1 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1; + } + if (first_row + 2 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2; + } + if (first_row + 3 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3; + } + } +} + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_q5_1_f32( + global void * src0, + ulong offset0, + global float * src1, + ulong offset1, + global float * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + src0 = (global void*)((global char*)src0 + offset0); + src1 = (global float*)((global char*)src1 + offset1); + dst = (global float*)((global char*)dst + offsetd); + + mul_vec_q_n_f32(src0, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3); +} diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32_flat.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32_flat.cl new file mode 100644 index 000000000..57c2f1409 --- /dev/null +++ b/ggml/src/ggml-opencl/kernels/mul_mv_q5_1_f32_flat.cl @@ -0,0 +1,247 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +#ifdef cl_intel_subgroups +#pragma OPENCL EXTENSION cl_intel_subgroups : enable +#else +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +#endif + +#ifdef cl_intel_required_subgroup_size +#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable +#define INTEL_GPU 1 +#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16))) +#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32))) +#elif defined(cl_qcom_reqd_sub_group_size) +#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable +#define ADRENO_GPU 1 +#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half"))) +#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full"))) +#endif + +#define QK5_1 32 + +inline float block_q5_1_dot_y_flat( + global const uchar * x, + global const uint * qh_ptr, + global const half * dh, + global const half * mh, + float sumy, + float16 yl, + int il, + global const float * yb +) { + float d = *dh; + float m = *mh; + global const ushort * qs = ((global const ushort *)(x + il)); + + float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f); + + acc.s0 += yl.s0 * (qs[0] & 0x000F); + acc.s0 += yl.s1 * (qs[0] & 0x0F00); + acc.s0 += yl.s8 * (qs[0] & 0x00F0); + acc.s3 += yl.s9 * (qs[0] & 0xF000); + + acc.s0 += yl.s2 * (qs[1] & 0x000F); + acc.s1 += yl.s3 * (qs[1] & 0x0F00); + acc.s2 += yl.sa * (qs[1] & 0x00F0); + acc.s3 += yl.sb * (qs[1] & 0xF000); + + acc.s0 += yl.s4 * (qs[2] & 0x000F); + acc.s1 += yl.s5 * (qs[2] & 0x0F00); + acc.s2 += yl.sc * (qs[2] & 0x00F0); + acc.s3 += yl.sd * (qs[2] & 0xF000); + + acc.s0 += yl.s6 * (qs[3] & 0x000F); + acc.s1 += yl.s7 * (qs[3] & 0x0F00); + acc.s2 += yl.se * (qs[3] & 0x00F0); + acc.s3 += yl.sf * (qs[3] & 0xF000); + + uint qh_val = *qh_ptr; + uchar qh_lo = (uchar)((qh_val >> il) & 0xFF); + uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF); + + float qh_sum = 0.0f; + qh_sum += yb[0] * (float)((qh_lo >> 0) & 1); + qh_sum += yb[1] * (float)((qh_lo >> 1) & 1); + qh_sum += yb[2] * (float)((qh_lo >> 2) & 1); + qh_sum += yb[3] * (float)((qh_lo >> 3) & 1); + qh_sum += yb[4] * (float)((qh_lo >> 4) & 1); + qh_sum += yb[5] * (float)((qh_lo >> 5) & 1); + qh_sum += yb[6] * (float)((qh_lo >> 6) & 1); + qh_sum += yb[7] * (float)((qh_lo >> 7) & 1); + qh_sum += yb[16] * (float)((qh_hi >> 0) & 1); + qh_sum += yb[17] * (float)((qh_hi >> 1) & 1); + qh_sum += yb[18] * (float)((qh_hi >> 2) & 1); + qh_sum += yb[19] * (float)((qh_hi >> 3) & 1); + qh_sum += yb[20] * (float)((qh_hi >> 4) & 1); + qh_sum += yb[21] * (float)((qh_hi >> 5) & 1); + qh_sum += yb[22] * (float)((qh_hi >> 6) & 1); + qh_sum += yb[23] * (float)((qh_hi >> 7) & 1); + + return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum) + sumy * m; +} + +#undef N_DST +#undef N_SIMDGROUP +#undef N_SIMDWIDTH + +#ifdef INTEL_GPU +#define N_DST 4 // each subgroup works on 4 rows +#define N_SIMDGROUP 1 // number of subgroups in a thread group +#define N_SIMDWIDTH 16 // assuming subgroup size is 16 +#elif defined (ADRENO_GPU) +#define N_DST 4 +#define N_SIMDGROUP 1 +#define N_SIMDWIDTH 64 +#endif + +inline void mul_vec_q_n_f32_flat( + global void * src0_qs, + global void * src0_qh, + global void * src0_d, + global void * src0_m, + global float * src1, + global float * dst, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + const ulong nb = ne00/QK5_1; + + int r0 = get_group_id(0); + int r1 = get_group_id(1); + int im = get_group_id(2); + + int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST; + + int i12 = im%ne12; + int i13 = im/ne12; + + ulong offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02); + + ulong offset0_qs = offset0 * (QK5_1/2); + + global uchar * x = (global uchar *) src0_qs + offset0_qs; + global uint * qh = (global uint *) src0_qh + offset0; + global half * d = (global half *) src0_d + offset0; + global half * ms = (global half *) src0_m + offset0; + global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1; + + float16 yl; + float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f); + + int ix = get_sub_group_local_id()/2; + int il = 8*(get_sub_group_local_id()%2); + + global float * yb = y + ix * QK5_1 + il; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) { + float sumy = 0; + + sumy += yb[0]; + sumy += yb[1]; + sumy += yb[2]; + sumy += yb[3]; + sumy += yb[4]; + sumy += yb[5]; + sumy += yb[6]; + sumy += yb[7]; + + sumy += yb[16]; + sumy += yb[17]; + sumy += yb[18]; + sumy += yb[19]; + sumy += yb[20]; + sumy += yb[21]; + sumy += yb[22]; + sumy += yb[23]; + + + yl.s0 = yb[0]; + yl.s1 = yb[1]/256.f; + + yl.s2 = yb[2]; + yl.s3 = yb[3]/256.f; + + yl.s4 = yb[4]; + yl.s5 = yb[5]/256.f; + + yl.s6 = yb[6]; + yl.s7 = yb[7]/256.f; + + yl.s8 = yb[16]/16.f; + yl.s9 = yb[17]/4096.f; + + yl.sa = yb[18]/16.f; + yl.sb = yb[19]/4096.f; + + yl.sc = yb[20]/16.f; + yl.sd = yb[21]/4096.f; + + yl.se = yb[22]/16.f; + yl.sf = yb[23]/4096.f; + + sumf.s0 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 0*nb*(QK5_1/2), qh + ib + 0*nb, d + ib + 0*nb, ms + ib + 0*nb, sumy, yl, il, yb); + sumf.s1 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 1*nb*(QK5_1/2), qh + ib + 1*nb, d + ib + 1*nb, ms + ib + 1*nb, sumy, yl, il, yb); + sumf.s2 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 2*nb*(QK5_1/2), qh + ib + 2*nb, d + ib + 2*nb, ms + ib + 2*nb, sumy, yl, il, yb); + sumf.s3 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 3*nb*(QK5_1/2), qh + ib + 3*nb, d + ib + 3*nb, ms + ib + 3*nb, sumy, yl, il, yb); + + yb += QK5_1 * (N_SIMDWIDTH/2); + } + + float4 tot = (float4)( + sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1), + sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3) + ); + + if (get_sub_group_local_id() == 0) { + if (first_row + 0 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0; + } + if (first_row + 1 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1; + } + if (first_row + 2 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2; + } + if (first_row + 3 < ne01) { + dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3; + } + } +} + +#ifdef INTEL_GPU +REQD_SUBGROUP_SIZE_16 +#elif defined (ADRENO_GPU) +REQD_SUBGROUP_SIZE_64 +#endif +kernel void kernel_mul_mv_q5_1_f32_flat( + global void * src0_qs, + global void * src0_qh, + global void * src0_d, + global void * src0_m, + global float * src1, + ulong offset1, + global float * dst, + ulong offsetd, + int ne00, + int ne01, + int ne02, + int ne10, + int ne12, + int ne0, + int ne1, + int r2, + int r3 +) { + src1 = (global float*)((global char*)src1 + offset1); + dst = (global float*)((global char*)dst + offsetd); + + mul_vec_q_n_f32_flat(src0_qs, src0_qh, src0_d, src0_m, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3); +} From 5aa3a64596a1dc67a5aeb55dbd3d743f4d84126c Mon Sep 17 00:00:00 2001 From: Christian Hoener zu Siederdissen Date: Mon, 1 Jun 2026 20:01:26 +0200 Subject: [PATCH 07/20] nix : add nix-nodejs facilities to build Web UI (#23846) * nix: add nix-nodejs facilities to build Web UI Build the Web UI locally using standard Nix systems for building NodeJS packages. - Create derivation for the web UI - npm dependencies are imported via buildNodeModules. Does not require setting any shasum. - Copy build artifacts to the correct folders. - Prevents having to download from huggingface.co Fixes #23067 * nix: simplify webui derivation using LLAMA_UI_OUT_DIR - Move npm build to installPhase with LLAMA_UI_OUT_DIR=$out to write output directly to the Nix store - Copy built assets to tools/ui/dist (source tree) instead of build/tools/ui/dist so CMake's copy_src_dist() finds them --- .devops/nix/package.nix | 29 ++++++++++++++++++++++++++++- 1 file changed, 28 insertions(+), 1 deletion(-) diff --git a/.devops/nix/package.nix b/.devops/nix/package.nix index 30355d2fc..86d9d589d 100644 --- a/.devops/nix/package.nix +++ b/.devops/nix/package.nix @@ -3,6 +3,7 @@ glibc, config, stdenv, + stdenvNoCC, runCommand, cmake, ninja, @@ -19,6 +20,8 @@ openssl, shaderc, spirv-headers, + nodejs, + importNpmLock, useBlas ? builtins.all (x: !x) [ useCuda @@ -130,7 +133,31 @@ effectiveStdenv.mkDerivation (finalAttrs: { src = lib.cleanSource ../../.; }; - postPatch = '' + # Builds the webui locally, taking care not to require updating any sha256 hash. + webui = stdenvNoCC.mkDerivation { + pname = "webui"; + version = llamaVersion; + src = lib.cleanSource ../../tools/ui; + + nativeBuildInputs = [ + nodejs + importNpmLock.linkNodeModulesHook + ]; + + # no sha256 required when using buildNodeModules + npmDeps = importNpmLock.buildNodeModules { + npmRoot = ../../tools/ui; + inherit nodejs; + }; + + installPhase = '' + LLAMA_UI_OUT_DIR=$out npm run build --offline + ''; + }; + + postPatch = lib.optionalString useWebUi '' + cp -r ${finalAttrs.webui} tools/ui/dist + chmod -R u+w tools/ui/dist ''; # With PR#6015 https://github.com/ggml-org/llama.cpp/pull/6015, From 5dcb71166686799f0d873eab7386234302d05ecf Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Mon, 1 Jun 2026 22:26:58 +0300 Subject: [PATCH 08/20] speculative : fix n_outputs_max and remove draft-simple auto-enable (#23988) * speculative : add common_speculative_n_max helper function Extract the speculative max-draft-size logic from server_n_outputs_max into a reusable common_speculative_n_max() function in common/speculative. Assisted-by: llama.cpp:local pi * cont : draft context always has n_parallel outputs * llama : log n_outputs_max * speculative : remove draft-simple auto-enable * ci : enable server tests on PRs --- .github/workflows/server.yml | 3 --- common/arg.cpp | 4 --- common/speculative.cpp | 46 ++++++++++++++++++++++++--------- common/speculative.h | 3 +++ src/llama-context.cpp | 1 + tools/server/server-context.cpp | 34 ++---------------------- 6 files changed, 40 insertions(+), 51 deletions(-) diff --git a/.github/workflows/server.yml b/.github/workflows/server.yml index f42b30d5c..5a02cc15a 100644 --- a/.github/workflows/server.yml +++ b/.github/workflows/server.yml @@ -102,7 +102,6 @@ jobs: - name: Tests id: server_integration_tests - if: ${{ !github.event.pull_request }} run: | cd tools/server/tests pytest -v -x -m "not slow" @@ -116,7 +115,6 @@ jobs: - name: Tests (Backend sampling) id: server_integration_tests_backend_sampling - if: ${{ !github.event.pull_request }} run: | cd tools/server/tests export LLAMA_ARG_BACKEND_SAMPLING=1 @@ -169,7 +167,6 @@ jobs: - name: Tests id: server_integration_tests - if: ${{ !github.event.pull_request }} run: | cd tools/server/tests $env:PYTHONIOENCODING = ":replace" diff --git a/common/arg.cpp b/common/arg.cpp index e0f6c6066..2df446d9d 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -1041,11 +1041,9 @@ common_params_context common_params_parser_init(common_params & params, llama_ex // we define here to make sure it's included in llama-gen-docs if (ex == LLAMA_EXAMPLE_COMPLETION) { params.use_jinja = false; // disable jinja by default - } else if (ex == LLAMA_EXAMPLE_MTMD) { params.use_jinja = false; // disable jinja by default params.sampling.temp = 0.2; // lower temp by default for better quality - } else if (ex == LLAMA_EXAMPLE_SERVER) { params.n_parallel = -1; // auto by default } @@ -1066,7 +1064,6 @@ common_params_context common_params_parser_init(common_params & params, llama_ex sampler_type_names.pop_back(); // remove last semicolon } - /** * filter options by example * rules: @@ -1080,7 +1077,6 @@ common_params_context common_params_parser_init(common_params & params, llama_ex } }; - add_opt(common_arg( {"-h", "--help", "--usage"}, "print usage and exit", diff --git a/common/speculative.cpp b/common/speculative.cpp index 253a5ecec..73830fda6 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -1317,6 +1317,40 @@ static uint32_t common_get_enabled_speculative_configs(const std::vectortypes) { + switch (type) { + case COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE: + case COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3: + case COMMON_SPECULATIVE_TYPE_DRAFT_MTP: + n_max = std::max(n_max, std::max(0, spec->draft.n_max)); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE: + n_max = std::max(n_max, (int32_t) spec->ngram_simple.size_m); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K: + n_max = std::max(n_max, (int32_t) spec->ngram_map_k.size_m); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V: + n_max = std::max(n_max, (int32_t) spec->ngram_map_k4v.size_m); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_MOD: + n_max = std::max(n_max, std::max(0, spec->ngram_mod.n_max)); + break; + case COMMON_SPECULATIVE_TYPE_NGRAM_CACHE: + n_max = std::max(n_max, (int32_t) 8); + break; + case COMMON_SPECULATIVE_TYPE_NONE: + case COMMON_SPECULATIVE_TYPE_COUNT: + break; + } + } + + return n_max; +} + // initialization of the speculative decoding system // common_speculative * common_speculative_init(common_params_speculative & params, uint32_t n_seq) { @@ -1325,8 +1359,6 @@ common_speculative * common_speculative_init(common_params_speculative & params, { uint32_t enabled_configs = common_get_enabled_speculative_configs(params.types); - bool has_draft_model_path = !params.draft.mparams.path.empty(); - bool has_draft_simple = (enabled_configs & (1u << COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE)); bool has_draft_eagle3 = false; // TODO PR-18039: if params.speculative.eagle3 bool has_mtp = (enabled_configs & (1u << COMMON_SPECULATIVE_TYPE_DRAFT_MTP)) && params.draft.ctx_dft != nullptr; @@ -1359,16 +1391,6 @@ common_speculative * common_speculative_init(common_params_speculative & params, if (has_ngram_cache) { configs.push_back(common_speculative_config(COMMON_SPECULATIVE_TYPE_NGRAM_CACHE, params)); } - if (has_draft_simple) { - if (!has_draft_model_path) { - LOG_WRN("%s: draft model is not specified - cannot use 'draft' type\n", __func__); - has_draft_simple = false; - } - } else if (has_draft_model_path && !has_mtp && !has_draft_eagle3) { - LOG_WRN("%s: draft model is specified but 'draft' speculative type is not explicitly enabled - enabling it\n", __func__); - has_draft_simple = true; - } - if (has_draft_simple) { configs.push_back(common_speculative_config(COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE, params)); } diff --git a/common/speculative.h b/common/speculative.h index f24bac79e..deba7dac7 100644 --- a/common/speculative.h +++ b/common/speculative.h @@ -20,6 +20,9 @@ enum common_speculative_type common_speculative_type_from_name(const std::string // convert type to string std::string common_speculative_type_to_str(enum common_speculative_type type); +// return the max number of draft tokens based on the speculative parameters +int32_t common_speculative_n_max(const common_params_speculative * spec); + common_speculative * common_speculative_init(common_params_speculative & params, uint32_t n_seq); void common_speculative_free(common_speculative * spec); diff --git a/src/llama-context.cpp b/src/llama-context.cpp index a7dfb2492..946c2d0ef 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -229,6 +229,7 @@ llama_context::llama_context( LLAMA_LOG_INFO("%s: freq_base = %.1f\n", __func__, cparams.rope_freq_base); LLAMA_LOG_INFO("%s: freq_scale = %g\n", __func__, cparams.rope_freq_scale); LLAMA_LOG_INFO("%s: n_rs_seq = %u\n", __func__, cparams.n_rs_seq); + LLAMA_LOG_INFO("%s: n_outputs_max = %u\n", __func__, cparams.n_outputs_max); if (cparams.n_ctx_seq < hparams.n_ctx_train) { LLAMA_LOG_WARN("%s: n_ctx_seq (%u) < n_ctx_train (%u) -- the full capacity of the model will not be utilized\n", diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index a5a688254..44fca83c6 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -45,35 +45,7 @@ static uint32_t server_n_outputs_max(const common_params & params) { return n_batch; } - uint32_t n_outputs_per_seq = 1; - - for (const auto type : params.speculative.types) { - switch (type) { - case COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE: - case COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3: - case COMMON_SPECULATIVE_TYPE_DRAFT_MTP: - n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + std::max(0, params.speculative.draft.n_max)); - break; - case COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE: - n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + params.speculative.ngram_simple.size_m); - break; - case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K: - n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + params.speculative.ngram_map_k.size_m); - break; - case COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V: - n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + params.speculative.ngram_map_k4v.size_m); - break; - case COMMON_SPECULATIVE_TYPE_NGRAM_MOD: - n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + std::max(0, params.speculative.ngram_mod.n_max)); - break; - case COMMON_SPECULATIVE_TYPE_NGRAM_CACHE: - n_outputs_per_seq = std::max(n_outputs_per_seq, 1 + 8); - break; - case COMMON_SPECULATIVE_TYPE_NONE: - case COMMON_SPECULATIVE_TYPE_COUNT: - break; - } - } + const uint32_t n_outputs_per_seq = 1 + common_speculative_n_max(¶ms.speculative); const uint64_t n_outputs = (uint64_t) params.n_parallel * n_outputs_per_seq; @@ -862,9 +834,7 @@ private: measure_model_bytes = false; } - if (!has_draft) { - params_dft.n_outputs_max = params_base.n_parallel; - } + params_dft.n_outputs_max = params_base.n_parallel; auto mparams_dft = common_model_params_to_llama(params_dft); auto cparams_dft = common_context_params_to_llama(params_dft); From b8275a8acc17b7b9f2842411f47c9af22405d5dd Mon Sep 17 00:00:00 2001 From: Masashi Yoshimura Date: Tue, 2 Jun 2026 08:59:06 +0900 Subject: [PATCH 09/20] revert to using global_invocation_id for cpy shader (#23955) --- ggml/src/ggml-webgpu/wgsl-shaders/cpy.wgsl | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-webgpu/wgsl-shaders/cpy.wgsl b/ggml/src/ggml-webgpu/wgsl-shaders/cpy.wgsl index e268adfb1..67f1dc092 100644 --- a/ggml/src/ggml-webgpu/wgsl-shaders/cpy.wgsl +++ b/ggml/src/ggml-webgpu/wgsl-shaders/cpy.wgsl @@ -50,13 +50,13 @@ var params: Params; @compute @workgroup_size(WG_SIZE) fn main( - @builtin(global_invocation_index) gindex: u32, + @builtin(global_invocation_id) gid: vec3, ) { - if (gindex >= params.ne) { + if (gid.x >= params.ne) { return; } - var i = gindex; + var i = gid.x; let i3 = i / (params.src_ne2 * params.src_ne1 * params.src_ne0); i = i % (params.src_ne2 * params.src_ne1 * params.src_ne0); let i2 = i / (params.src_ne1 * params.src_ne0); @@ -64,7 +64,7 @@ fn main( let i1 = i / params.src_ne0; let i0 = i % params.src_ne0; - var j = gindex; + var j = gid.x; let j3 = j / (params.dst_ne2 * params.dst_ne1 * params.dst_ne0); j = j % (params.dst_ne2 * params.dst_ne1 * params.dst_ne0); let j2 = j / (params.dst_ne1 * params.dst_ne0); @@ -80,4 +80,3 @@ fn main( dst[params.offset_dst + dst_idx] = DST_TYPE((src[params.offset_src + src_idx])); } - From 210a6570ceda20c5d6439172c09ada08c3754cc9 Mon Sep 17 00:00:00 2001 From: lhez Date: Mon, 1 Jun 2026 19:15:09 -0700 Subject: [PATCH 10/20] opencl: fix compiler warnings for non-adreno path (#23922) * opencl: fix compiler warnings for non-adreno path * opencl: fix const cast warning --- ggml/src/ggml-opencl/ggml-opencl.cpp | 20 ++++++++++++-------- 1 file changed, 12 insertions(+), 8 deletions(-) diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 7cafbe0cd..b67ea46bc 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -380,7 +380,7 @@ struct ggml_backend_opencl_device_context { ADRENO_GPU_GEN adreno_gen = ADRENO_GPU_GEN::ADRENO_UNKNOWN; std::regex *opfilter = nullptr; // regex of ops to not claim - std::string opfilter_str; // regex string for opfilter + std::string opfilter_str = ""; // regex string for opfilter size_t global_mem_size = 0; }; @@ -6822,9 +6822,6 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, cl_buffer_region region; - cl_uchar mask_0F = 0x0F; - cl_uchar mask_F0 = 0xF0; - #ifdef GGML_OPENCL_USE_ADRENO_KERNELS // Adreno MoE Q6_K kernel needs special transposed layout if (use_adreno_moe_kernels(backend_ctx, tensor)) { @@ -6858,6 +6855,9 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, cl_kernel kernel = backend_ctx->kernel_convert_block_q6_k_trans4_ns; + cl_uchar mask_0F = 0x0F; + cl_uchar mask_F0 = 0xF0; + int ne00 = tensor->ne[0]; int ne01 = tensor->ne[1]; int ne02 = tensor->ne[2]; @@ -6994,7 +6994,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer, cl_int err; cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, - size, (void *) data, &err); + size, const_cast(data), &err); CL_CHECK(err); cl_kernel kernel = backend_ctx->kernel_convert_bf16_to_f16; @@ -7782,9 +7782,6 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, if (tensor->type == GGML_TYPE_Q6_K) { ggml_tensor_extra_cl_q6_K * extra = (ggml_tensor_extra_cl_q6_K *)tensor->extra; - cl_uchar mask_0F = 0x0F; - cl_uchar mask_F0 = 0xF0; - #ifdef GGML_OPENCL_USE_ADRENO_KERNELS if (use_adreno_moe_kernels(backend_ctx, tensor)) { cl_int err; @@ -7794,6 +7791,9 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer, cl_kernel kernel = backend_ctx->kernel_restore_block_q6_k_trans4_ns; + cl_uchar mask_0F = 0x0F; + cl_uchar mask_F0 = 0xF0; + int ne00 = tensor->ne[0]; int ne01 = tensor->ne[1]; int ne02 = tensor->ne[2]; @@ -14888,6 +14888,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, const int ne1 = dst->ne[1]; const int ne2 = dst->ne[2]; + GGML_UNUSED(ne2); + const int r2 = ne12/ne02; const int r3 = ne13/ne03; const int dst_rows = ne20*ne21; // ne20 = n_used_experts, ne21 = n_rows @@ -14902,6 +14904,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, const int n_tile_size = 32; const int max_post_router_tile = (ne20 * ne21 / n_tile_size) + ne02; + GGML_UNUSED(max_post_router_tile); + cl_kernel kernel; // subgroup mat vec From 1fd5f4803713ea3e1eda326483c9cc71a572cf02 Mon Sep 17 00:00:00 2001 From: Anav Prasad Date: Mon, 1 Jun 2026 19:38:37 -0700 Subject: [PATCH 11/20] clean up unused variables warnings (#23975) --- ggml/src/ggml-cuda/fattn-mma-f16.cuh | 6 +++--- ggml/src/ggml-cuda/gated_delta_net.cu | 10 +++++----- ggml/src/ggml-cuda/mmf.cuh | 6 +++--- ggml/src/ggml-cuda/mmvf.cu | 13 ++++++------- ggml/src/ggml-cuda/mmvq.cu | 13 ++++--------- ggml/src/ggml-cuda/topk-moe.cu | 2 +- 6 files changed, 22 insertions(+), 28 deletions(-) diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index 3c8b6eaaf..ac5abb133 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -568,7 +568,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( constexpr bool Q_in_reg = ggml_cuda_fattn_mma_get_Q_in_reg (DKQ, DV, ncols); constexpr int nstages = ggml_cuda_fattn_mma_get_nstages (DKQ, DV, ncols1, ncols2); - constexpr int stride_tile_Q = DKQ/2 + 4; constexpr int stride_tile_K = nbatch_K2 + 4; constexpr int stride_tile_V = V_is_K_view ? stride_tile_K : nbatch_V2 + 4; @@ -604,9 +603,9 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( #pragma unroll for (int k0_start = (DKQ/2-1) - (DKQ/2-1) % nbatch_K2; k0_start >= 0; k0_start -= nbatch_K2) { const int k0_stop = k0_start + nbatch_K2 < DKQ/2 ? k0_start + nbatch_K2 : DKQ/2; - const int k0_diff = k0_stop - k0_start; if constexpr (nstages <= 1) { + const int k0_diff = k0_stop - k0_start; constexpr bool use_cp_async = nstages == 1; flash_attn_ext_f16_load_tile (K_h2 + int64_t(k_VKQ_0)*stride_K + k0_start, tile_K, k0_diff, stride_K, k_VKQ_sup); @@ -640,6 +639,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( } } } else { + constexpr int stride_tile_Q = DKQ/2 + 4; #pragma unroll for (int k_KQ_0 = k0_start; k_KQ_0 < k0_stop; k_KQ_0 += T_A_KQ::J) { load_ldmatrix(Q_B[0], tile_Q + (threadIdx.y / np)*(T_B_KQ::I*stride_tile_Q) + k_KQ_0, stride_tile_Q); @@ -954,9 +954,9 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter( for (int i0_start = 0; i0_start < DV; i0_start += 2*nbatch_V2) { static_assert(DV % (2*nbatch_V2) == 0, "bad loop size"); const int i0_stop = i0_start + 2*nbatch_V2; - const int i0_diff = i0_stop - i0_start; if constexpr (nstages <= 1) { + const int i0_diff = i0_stop - i0_start; if (!V_is_K_view || i0_stop > 2*nbatch_K2) { constexpr bool use_cp_async = nstages == 1; flash_attn_ext_f16_load_tile diff --git a/ggml/src/ggml-cuda/gated_delta_net.cu b/ggml/src/ggml-cuda/gated_delta_net.cu index 018d5d37d..7cfda6523 100644 --- a/ggml/src/ggml-cuda/gated_delta_net.cu +++ b/ggml/src/ggml-cuda/gated_delta_net.cu @@ -43,7 +43,6 @@ gated_delta_net_cuda(const float * q, // output state layout (per-slot D * n_seqs) — same per-(seq,head) offset as before. const int64_t state_in_offset = sequence * K * H * S_v * S_v + h_idx * S_v * S_v; const int64_t state_out_offset = (sequence * H + h_idx) * S_v * S_v; - const int64_t state_size_per_token = S_v * S_v * H * n_seqs; // per-slot stride in output state += state_out_offset; curr_state += state_in_offset + col * S_v; attn_data += (sequence * n_tokens * H + h_idx) * S_v; @@ -61,10 +60,6 @@ gated_delta_net_cuda(const float * q, s_shard[r] = curr_state[i]; } - // slot mapping: target_slot = t - shift. When n_tokens < K only the last n_tokens slots - // are written; earlier slots are left untouched (caller-owned). - const int shift = (int) n_tokens - K; - for (int t = 0; t < n_tokens; t++) { const float * q_t = q + iq3 * sq3 + t * sq2 + iq1 * sq1; const float * k_t = k + iq3 * sq3 + t * sq2 + iq1 * sq1; @@ -148,6 +143,11 @@ gated_delta_net_cuda(const float * q, attn_data += S_v * H; if constexpr (keep_rs_t) { + // slot mapping: target_slot = t - shift. When n_tokens < K only the last n_tokens slots + // are written; earlier slots are left untouched (caller-owned). + const int shift = (int) n_tokens - K; + + const int64_t state_size_per_token = S_v * S_v * H * n_seqs; // per-slot stride in output const int target_slot = t - shift; if (target_slot >= 0 && target_slot < K) { float * curr_state = (dst + attn_score_elems) + target_slot * state_size_per_token + state_out_offset; diff --git a/ggml/src/ggml-cuda/mmf.cuh b/ggml/src/ggml-cuda/mmf.cuh index c2a8d54c9..d55cc1ec7 100644 --- a/ggml/src/ggml-cuda/mmf.cuh +++ b/ggml/src/ggml-cuda/mmf.cuh @@ -91,7 +91,7 @@ static __global__ void mul_mat_f( const int row0 = blockIdx.x * rows_per_block; int expert_idx = 0; - int col_base = 0; + [[maybe_unused]] int col_base = 0; const int channel_dst = has_ids ? 0 : blockIdx.y; @@ -122,12 +122,12 @@ static __global__ void mul_mat_f( ids += col_offset * stride_row_id; } - const float2 * y2 = (const float2 *) y; + [[maybe_unused]] const float2 * y2 = (const float2 *) y; extern __shared__ char data_mmv[]; char * shmem_base = data_mmv; - int * slot_map = (int *) shmem_base; + [[maybe_unused]] int * slot_map = (int *) shmem_base; char * compute_base = has_ids ? (shmem_base + GGML_PAD(cols_per_block, 16) * sizeof(int)) : shmem_base; tile_C C[ntA][ntB]; diff --git a/ggml/src/ggml-cuda/mmvf.cu b/ggml/src/ggml-cuda/mmvf.cu index 09d95f309..3d6de64b7 100644 --- a/ggml/src/ggml-cuda/mmvf.cu +++ b/ggml/src/ggml-cuda/mmvf.cu @@ -80,9 +80,8 @@ static __global__ void mul_mat_vec_f( gate_x += int64_t(sample_x) *stride_sample_x + channel_x *stride_channel_x + row*stride_row; } - const int channel_bias = ids ? channel_x : channel_dst; - if constexpr (has_fusion) { + const int channel_bias = ids ? channel_x : channel_dst; if (use_bias) { x_bias += int64_t(sample_dst)*stride_sample_dst + channel_bias*stride_channel_dst; } @@ -95,7 +94,7 @@ static __global__ void mul_mat_vec_f( extern __shared__ char data_mmv[]; float * buf_iw = (float *) data_mmv; - float * buf_iw_gate = nullptr; + [[maybe_unused]] float * buf_iw_gate = nullptr; if constexpr (has_fusion) { buf_iw_gate = (float *) (data_mmv + warp_size*sizeof(float)); } @@ -123,7 +122,7 @@ static __global__ void mul_mat_vec_f( if constexpr (std::is_same_v) { const float2 * x2 = (const float2 *) x; - const float2 * gate_x2 = nullptr; + [[maybe_unused]] const float2 * gate_x2 = nullptr; if constexpr (has_fusion) { if (use_gate) { gate_x2 = (const float2 *) gate_x; @@ -155,7 +154,7 @@ static __global__ void mul_mat_vec_f( } } else if constexpr (std::is_same_v) { const half2 * x2 = (const half2 *) x; - const half2 * gate_x2 = nullptr; + [[maybe_unused]] const half2 * gate_x2 = nullptr; if constexpr (has_fusion) { if (use_gate) { gate_x2 = (const half2 *) gate_x; @@ -266,7 +265,7 @@ static __global__ void mul_mat_vec_f( } #else const nv_bfloat162 * x2 = (const nv_bfloat162 *) x; - const nv_bfloat162 * gate_x2 = nullptr; + [[maybe_unused]] const nv_bfloat162 * gate_x2 = nullptr; if constexpr (has_fusion) { if (use_gate) { gate_x2 = (const nv_bfloat162 *) gate_x; @@ -274,7 +273,7 @@ static __global__ void mul_mat_vec_f( } for (int col2 = tid; col2 < ncols2; col2 += block_size) { const nv_bfloat162 tmpx = x2[col2]; - nv_bfloat162 tmpx_gate; + [[maybe_unused]] nv_bfloat162 tmpx_gate; if constexpr (has_fusion) { if (use_gate) { tmpx_gate = gate_x2[col2]; diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index ecb6fdeda..86b4a4930 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -515,7 +515,7 @@ static __global__ void mul_mat_vec_q( bool use_gate = false; bool use_bias = false; bool use_gate_bias = false; - const void * vgate = nullptr; + [[maybe_unused]] const void * vgate = nullptr; const float * x_bias = nullptr; const float * gate_bias = nullptr; ggml_glu_op active_glu; @@ -531,8 +531,8 @@ static __global__ void mul_mat_vec_q( } - float x_biases[ncols_dst] = { 0.0f }; - float gate_biases[ncols_dst] = { 0.0f }; + [[maybe_unused]] float x_biases[ncols_dst] = { 0.0f }; + [[maybe_unused]] float gate_biases[ncols_dst] = { 0.0f }; if constexpr (has_fusion) { const uint32_t channel_bias = ids ? channel_x : channel_dst; if (use_bias) { @@ -589,12 +589,7 @@ static __global__ void mul_mat_vec_q( } __shared__ float tmp_shared[nwarps-1 > 0 ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; - __shared__ float tmp_shared_gate[(has_fusion && (nwarps-1 > 0)) ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; - if constexpr (!has_fusion) { - (void) tmp_shared_gate; - } else if (!use_gate) { - (void) tmp_shared_gate; - } + [[maybe_unused]] __shared__ float tmp_shared_gate[(has_fusion && (nwarps-1 > 0)) ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; if (threadIdx.y > 0) { #pragma unroll diff --git a/ggml/src/ggml-cuda/topk-moe.cu b/ggml/src/ggml-cuda/topk-moe.cu index da20c9aab..c4253bfa4 100644 --- a/ggml/src/ggml-cuda/topk-moe.cu +++ b/ggml/src/ggml-cuda/topk-moe.cu @@ -134,7 +134,7 @@ __launch_bounds__(4 * WARP_SIZE, 1) __global__ void topk_moe_cuda(const float * // selection_wt is only needed when bias is present (selection uses wt + bias) // when no bias, we use wt directly for both selection and weight values - float selection_wt[has_bias ? experts_per_thread : 1]; + [[maybe_unused]] float selection_wt[has_bias ? experts_per_thread : 1]; if constexpr (has_bias) { #pragma unroll From 354ebac8cb92e93eb6f22bd507d4249b6846b90d Mon Sep 17 00:00:00 2001 From: Pascal Date: Tue, 2 Jun 2026 07:26:20 +0200 Subject: [PATCH 12/20] server: real-time reasoning interruption via control endpoint (#23971) * server: real-time reasoning interruption via control endpoint Builds on the manual reasoning budget trigger from #23949. Adds a CONTROL task that mirrors the CANCEL path on the live slot and calls common_sampler_reasoning_budget_force to end thinking mid-generation. POST /v1/chat/completions/control with { id_slot, action }, opt-in reasoning_control arms the budget sampler on demand. Router and single model. Minimal WebUI button as a skeleton for further UI work. * ui: track reasoning phase via explicit streaming state Add isReasoning to the chat store, mirroring the isLoading pattern: per conversation map, private setter, public accessor and reactive export. Set from the stream callbacks, true on reasoning chunks, false on the first content chunk, reset on stream end and resynced on conversation switch. The skip button now keys off isReasoning so it shows only during the thinking phase, not the whole generation. * ui: extract control endpoint and action into constants Move the chat completion routes, the slots route and the reasoning control action out of chat.service into api-endpoints and a dedicated control-actions module. No behavior change, drops the magic strings so the control protocol has a single source of truth. * server: target reasoning control by completion id Address @ngxson review on the control endpoint. Switch from id_slot to the chat completion id to avoid a TOCTOU: the slot can be reassigned between the lookup and the control request, so matching the live completion (oaicompat_cmpl_id) is safe and a finished one simply matches nothing. Rename the action to reasoning_end, guard it on the reasoning_control flag of the target slot, and reduce the response to {success} with an optional message. * ui: target reasoning control by completion id Keep the streamed completion id on the message and post it back to the control endpoint instead of probing /slots. Drops the slot discovery and the TOCTOU that came with it. Action renamed to reasoning_end, response read as {success}. * server: address review from @ngxson Move the control fields into task_params and drop the redundant comments on the control path. * server: document the reasoning control endpoint * Update tools/ui/src/lib/types/database.d.ts Co-authored-by: Aleksander Grygier * ui: rename cmplId to completionId Per @allozaur review, clearer name for the streamed completion id. * ui: wire completion id capture through the agentic flow The webui streams through the agentic flow, which relayed onModel but not onCompletionId, so the completion id never reached the message and the control request was never sent. Relay it through the flow and its callbacks type, declare id on the chunk type, and log an explicit error when the button fires without a usable id. * ui: target reasoning control model from the message The model is a property of the completion, so read it from the streaming message like the id, not from the model dropdown which is unrelated UI state. Makes the request self-consistent by construction instead of just unlikely to drift. --------- Co-authored-by: Aleksander Grygier --- common/common.h | 1 + common/sampling.cpp | 2 +- tools/server/README.md | 18 ++++ tools/server/server-common.cpp | 1 + tools/server/server-context.cpp | 82 +++++++++++++++++++ tools/server/server-context.h | 1 + tools/server/server-task.cpp | 1 + tools/server/server-task.h | 18 ++++ tools/server/server.cpp | 2 + .../app/chat/ChatForm/ChatForm.svelte | 1 + .../ChatFormActions/ChatFormActions.svelte | 25 +++++- tools/ui/src/lib/constants/api-endpoints.ts | 11 +++ tools/ui/src/lib/constants/control-actions.ts | 7 ++ tools/ui/src/lib/constants/index.ts | 1 + tools/ui/src/lib/services/chat.service.ts | 67 ++++++++++++++- tools/ui/src/lib/stores/agentic.svelte.ts | 2 + tools/ui/src/lib/stores/chat.svelte.ts | 37 +++++++++ tools/ui/src/lib/types/agentic.d.ts | 1 + tools/ui/src/lib/types/api.d.ts | 1 + tools/ui/src/lib/types/chat.d.ts | 1 + tools/ui/src/lib/types/database.d.ts | 2 + tools/ui/src/lib/types/settings.d.ts | 1 + 22 files changed, 277 insertions(+), 6 deletions(-) create mode 100644 tools/ui/src/lib/constants/control-actions.ts diff --git a/common/common.h b/common/common.h index 92064a0e4..f2c7ee027 100644 --- a/common/common.h +++ b/common/common.h @@ -277,6 +277,7 @@ struct common_params_sampling { std::vector reasoning_budget_end; // end tag token sequence std::vector reasoning_budget_forced; // forced sequence (message + end tag) std::string reasoning_budget_message; // message injected before end tag when budget exhausted + bool reasoning_control = false; // create the budget sampler on demand so reasoning can be ended at runtime bool backend_sampling = false; diff --git a/common/sampling.cpp b/common/sampling.cpp index bbfa9a9ec..85f8ed50b 100644 --- a/common/sampling.cpp +++ b/common/sampling.cpp @@ -293,7 +293,7 @@ struct common_sampler * common_sampler_init(const struct llama_model * model, st } // reasoning budget sampler (skip when budget is unlimited unless a lazy grammar is active, which needs rbudget for thinking-block suppression) - if (!params.reasoning_budget_start.empty() && !params.reasoning_budget_end.empty() && (params.grammar_lazy || params.reasoning_budget_tokens >= 0)) { + if (!params.reasoning_budget_start.empty() && !params.reasoning_budget_end.empty() && (params.grammar_lazy || params.reasoning_budget_tokens >= 0 || params.reasoning_control)) { rbudget = common_reasoning_budget_init( vocab, params.reasoning_budget_start, diff --git a/tools/server/README.md b/tools/server/README.md index df30ca646..f1eeec36a 100644 --- a/tools/server/README.md +++ b/tools/server/README.md @@ -1244,6 +1244,8 @@ The `response_format` parameter supports both plain JSON output (e.g. `{"type": `reasoning_format`: The reasoning format to be parsed. If set to `none`, it will output the raw generated text. +`reasoning_control`: Arms realtime reasoning control for this completion so it can be ended early via `/v1/chat/completions/control`. Defaults to `false`. + `generation_prompt`: The generation prompt that was prefilled in by the template. Prepended to model output before parsing. `parse_tool_calls`: Whether to parse the generated tool call. @@ -1350,6 +1352,22 @@ The server supports parsing and returning reasoning via the `reasoning_content` Reasoning input (preserve reasoning in history) is also supported by some specific templates. For more details, please refer to [PR#18994](https://github.com/ggml-org/llama.cpp/pull/18994). +### POST `/v1/chat/completions/control`: Control a running chat completion in real time + +Acts on an in-flight completion identified by its `id` (the `id` field streamed back by `/v1/chat/completions`). The request is processed in parallel with the SSE stream, so the client sends it while still reading tokens. + +*Options:* + +`id`: (Required) The chat completion id to act on. A completion that has already finished matches nothing and the call is a no-op. + +`action`: (Required) The control action to perform. Currently the only supported value is `reasoning_end`, which forces the end of the current reasoning block so the model moves on to the final answer. Requires `reasoning_control: true` on the original completion request. + +`model`: (Required in router mode) The model name, used to route the request to the right instance. Ignored in single model mode. + +**Response format** + +Returns a JSON object with a boolean `success` field, and an optional `message` field describing the reason when `success` is `false`. + ### POST `/v1/responses`: OpenAI-compatible Responses API *Options:* diff --git a/tools/server/server-common.cpp b/tools/server/server-common.cpp index fb71792fe..4c3f16a0a 100644 --- a/tools/server/server-common.cpp +++ b/tools/server/server-common.cpp @@ -1132,6 +1132,7 @@ json oaicompat_chat_params_parse( llama_params["reasoning_budget_start_tag"] = chat_params.thinking_start_tag; llama_params["reasoning_budget_end_tag"] = chat_params.thinking_end_tag; llama_params["reasoning_budget_message"] = opt.reasoning_budget_message; + llama_params["reasoning_control"] = json_value(body, "reasoning_control", false); } } diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index 44fca83c6..fae73f09f 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -1263,6 +1263,20 @@ private: return nullptr; } + server_slot * get_slot_by_cmpl_id(const std::string & cmpl_id) { + if (cmpl_id.empty()) { + return nullptr; + } + + for (server_slot & slot : slots) { + if (slot.is_processing() && slot.task && slot.task->params.oaicompat_cmpl_id == cmpl_id) { + return &slot; + } + } + + return nullptr; + } + server_slot * get_available_slot(const server_task & task) { server_slot * ret = nullptr; @@ -2114,6 +2128,37 @@ private: } } } break; + case SERVER_TASK_TYPE_CONTROL: + { + auto res = std::make_unique(); + res->id = task.id; + + server_slot * slot = get_slot_by_cmpl_id(task.params.control_cmpl_id); + if (slot == nullptr) { + res->success = false; + res->message = "no active completion for this id"; + queue_results.send(std::move(res)); + break; + } + + if (task.params.control_action == "reasoning_end") { + // the budget sampler only exists when reasoning control was armed + if (!slot->task->params.sampling.reasoning_control) { + res->success = false; + res->message = "reasoning control not enabled for this completion"; + queue_results.send(std::move(res)); + break; + } + // act on the live slot mid generation, never defer + common_sampler_reasoning_budget_force(slot->smpl.get()); + res->success = true; + } else { + res->success = false; + res->message = "unknown control action"; + } + + queue_results.send(std::move(res)); + } break; case SERVER_TASK_TYPE_NEXT_RESPONSE: { // do nothing @@ -4266,6 +4311,43 @@ void server_routes::init_routes() { TASK_RESPONSE_TYPE_OAI_CHAT); }; + this->post_control = [this](const server_http_req & req) { + auto res = create_response(); + const json body = json::parse(req.body); + + const std::string cmpl_id = json_value(body, "id", std::string()); + const std::string action = json_value(body, "action", std::string()); + if (cmpl_id.empty()) { + res->error(format_error_response("missing completion id", ERROR_TYPE_INVALID_REQUEST)); + return res; + } + if (action != "reasoning_end") { + res->error(format_error_response("unknown control action", ERROR_TYPE_INVALID_REQUEST)); + return res; + } + + auto & rd = res->rd; + { + server_task task(SERVER_TASK_TYPE_CONTROL); + task.id = rd.get_new_id(); + task.params.control_cmpl_id = cmpl_id; + task.params.control_action = action; + rd.post_task(std::move(task)); + } + + auto result = rd.next(req.should_stop); + if (!result) { + GGML_ASSERT(req.should_stop()); + return res; + } + if (result->is_error()) { + res->error(result->to_json()); + return res; + } + res->ok(result->to_json()); + return res; + }; + this->post_responses_oai = [this](const server_http_req & req) { auto res = create_response(); std::vector files; diff --git a/tools/server/server-context.h b/tools/server/server-context.h index 65853438c..73caff54a 100644 --- a/tools/server/server-context.h +++ b/tools/server/server-context.h @@ -110,6 +110,7 @@ struct server_routes { server_http_context::handler_t post_completions; server_http_context::handler_t post_completions_oai; server_http_context::handler_t post_chat_completions; + server_http_context::handler_t post_control; server_http_context::handler_t post_responses_oai; server_http_context::handler_t post_transcriptions_oai; server_http_context::handler_t post_anthropic_messages; diff --git a/tools/server/server-task.cpp b/tools/server/server-task.cpp index ff80be6cc..33de2e4d9 100644 --- a/tools/server/server-task.cpp +++ b/tools/server/server-task.cpp @@ -499,6 +499,7 @@ task_params server_task::params_from_json_cmpl( const auto end_tag = json_value(data, "reasoning_budget_end_tag", std::string()); const auto message = json_value(data, "reasoning_budget_message", std::string()); params.sampling.reasoning_budget_tokens = budget; + params.sampling.reasoning_control = json_value(data, "reasoning_control", false); if (!start_tag.empty()) { params.sampling.reasoning_budget_start = common_tokenize(vocab, start_tag, false, true); diff --git a/tools/server/server-task.h b/tools/server/server-task.h index d47dc690c..bdadcff76 100644 --- a/tools/server/server-task.h +++ b/tools/server/server-task.h @@ -19,6 +19,7 @@ enum server_task_type { SERVER_TASK_TYPE_RERANK, SERVER_TASK_TYPE_INFILL, SERVER_TASK_TYPE_CANCEL, + SERVER_TASK_TYPE_CONTROL, SERVER_TASK_TYPE_NEXT_RESPONSE, SERVER_TASK_TYPE_METRICS, SERVER_TASK_TYPE_SLOT_SAVE, @@ -84,6 +85,10 @@ struct task_params { std::string oaicompat_model; std::string oaicompat_cmpl_id; + // realtime control (SERVER_TASK_TYPE_CONTROL) + std::string control_action; + std::string control_cmpl_id; + // per-request parameters for chat parsing common_chat_parser_params chat_parser_params; @@ -551,6 +556,19 @@ struct server_task_result_slot_erase : server_task_result { virtual json to_json() override; }; +struct server_task_result_control : server_task_result { + bool success = false; + std::string message; // optional detail when success is false + + virtual json to_json() override { + json out = json { { "success", success } }; + if (!message.empty()) { + out["message"] = message; + } + return out; + } +}; + struct server_task_result_get_lora : server_task_result { struct lora { common_adapter_lora_info info; diff --git a/tools/server/server.cpp b/tools/server/server.cpp index 4d56d45e8..769e80a80 100644 --- a/tools/server/server.cpp +++ b/tools/server/server.cpp @@ -149,6 +149,7 @@ int llama_server(int argc, char ** argv) { routes.post_completions = models_routes->proxy_post; routes.post_completions_oai = models_routes->proxy_post; routes.post_chat_completions = models_routes->proxy_post; + routes.post_control = models_routes->proxy_post; routes.post_responses_oai = models_routes->proxy_post; routes.post_transcriptions_oai = models_routes->proxy_post; routes.post_anthropic_messages = models_routes->proxy_post; @@ -185,6 +186,7 @@ int llama_server(int argc, char ** argv) { ctx_http.post("/v1/completions", ex_wrapper(routes.post_completions_oai)); ctx_http.post("/chat/completions", ex_wrapper(routes.post_chat_completions)); ctx_http.post("/v1/chat/completions", ex_wrapper(routes.post_chat_completions)); + ctx_http.post("/v1/chat/completions/control", ex_wrapper(routes.post_control)); ctx_http.post("/v1/responses", ex_wrapper(routes.post_responses_oai)); ctx_http.post("/responses", ex_wrapper(routes.post_responses_oai)); ctx_http.post("/v1/audio/transcriptions", ex_wrapper(routes.post_transcriptions_oai)); diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte index 46ac82334..99faecb6b 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatForm.svelte @@ -541,6 +541,7 @@ canSend={canSubmit} {disabled} {isLoading} + isReasoning={chatStore.isReasoning} {isRecording} {showAddButton} {showModelSelector} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte index a94293dd9..627107ef5 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte @@ -1,6 +1,7 @@
{/if} + {#if isReasoning} + + {/if} + {#if isLoading && !canSubmit} - {:else if item.disabledTooltip} - - - - + Add files + - -

{item.disabledTooltip}

- - - {/if} - {/each} + +
+ {#each ATTACHMENT_FILE_ITEMS as item (item.id)} + {@const enabled = attachmentMenu.isItemEnabled(item.enabledWhen)} + {#if enabled} + + {:else if item.disabledTooltip} + + + + + + +

{item.disabledTooltip}

+ + + {/if} + {/each} +
+ + + + (mcpExpanded = open)}> + + {#if mcpExpanded} + + {:else} + + {/if} + + + + MCP Servers + + + {getEnabledMcpServers().length} server{getEnabledMcpServers().length !== 1 ? 's' : ''} + + + + +
+ {#each getEnabledMcpServers() as server (server.id)} + {@const healthState = mcpStore.getHealthCheckState(server.id)} + {@const hasError = healthState.status === HealthCheckStatus.ERROR} + {@const displayName = mcpStore.getServerLabel(server)} + {@const faviconUrl = mcpStore.getServerFavicon(server.id)} + {@const isEnabled = conversationsStore.isMcpServerEnabledForChat(server.id)} - {#if !attachmentMenu.isItemEnabled('hasVisionModality')} - {@const pdfItem = ATTACHMENT_FILE_ITEMS.find((i) => i.id === AttachmentMenuItemId.PDF)} - {#if pdfItem} - - - + {/each} - -

PDFs will be converted to text. Image-based PDFs may not work properly.

- - - {/if} + {#if getEnabledMcpServers().length === 0} +
+ No MCP servers configured +
+ {/if} +
+ + + + {#if toolsPanel.totalToolCount > 0} + (toolsExpanded = open)}> + + {#if toolsExpanded} + + {:else} + + {/if} + + + + Tools + + + {toolsPanel.totalToolCount} tool{toolsPanel.totalToolCount !== 1 ? 's' : ''} + + + + +
+ {#each toolsPanel.activeGroups as group (group.label)} + {@const { checked, indeterminate } = toolsPanel.getGroupCheckedState(group)} + {@const enabledCount = toolsPanel.getEnabledToolCount(group)} + {@const favicon = toolsPanel.getFavicon(group)} + + + {/each} +
+ + {/if} - {#each ATTACHMENT_EXTRA_ITEMS as item (item.id)} - {#if item.id === AttachmentMenuItemId.SYSTEM_MESSAGE} - - - - + System Message + - -

{attachmentMenu.getSystemMessageTooltip()}

- - - {/if} - {/each} + {#if hasMcpPromptsSupport} + + {/if} - - + {#if hasMcpResourcesSupport} + - {/if} - {/each} + MCP Resources + + {/if}
diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte index b11467da8..813227fbc 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddToolsSubmenu.svelte @@ -107,7 +107,7 @@ toolsStore.toggleGroup(group)} + onCheckedChange={() => toolsPanel.toggleGroupByLabel(group.label)} class="mr-2 h-4 w-4 shrink-0" /> diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte index fd866b243..712326cba 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionModels.svelte @@ -1,6 +1,11 @@ + +{#if modelSupportsThinking} + + + {#if thinkingEnabled} + + {:else} + + {/if} + + Thinking + + {#if thinkingEnabled} + {currentEffort} + {:else} + off + {/if} + + + + {#each REASONING_EFFORT_LEVELS as level (level.value)} + + {/each} + + +{/if} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormReasoningToggle.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormReasoningToggle.svelte new file mode 100644 index 000000000..f6bcbcb09 --- /dev/null +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormReasoningToggle.svelte @@ -0,0 +1,145 @@ + + +{#if modelSupportsThinking} + + + + + {#if thinkingEnabled} + + {:else} + + {/if} + + + + +

{tooltipText}

+ + + + +
Reasoning effort
+ + {#each REASONING_EFFORT_LEVELS as level (level.value)} + + {/each} + + +{/if} diff --git a/tools/ui/src/lib/components/app/chat/index.ts b/tools/ui/src/lib/components/app/chat/index.ts index be5535960..8ed3cc65e 100644 --- a/tools/ui/src/lib/components/app/chat/index.ts +++ b/tools/ui/src/lib/components/app/chat/index.ts @@ -240,6 +240,15 @@ export { default as ChatFormActionAddToolsSubmenu } from './ChatForm/ChatFormAct */ export { default as ChatFormActionAddMcpServersSubmenu } from './ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddMcpServersSubmenu.svelte'; +/** + * **ChatFormReasoningToggle** - Thinking toggle button with effort dropdown + * + * A toggle button with lightbulb icon that indicates thinking status. + * Shows the reasoning effort dropdown when clicked. + * Only visible when the current model supports thinking. + */ +export { default as ChatFormReasoningToggle } from './ChatForm/ChatFormActions/ChatFormReasoningToggle.svelte'; + /** * Hidden file input element for programmatic file selection. */ diff --git a/tools/ui/src/lib/components/app/models/ModelId.svelte b/tools/ui/src/lib/components/app/models/ModelId.svelte index 2fe952fee..f566b55ee 100644 --- a/tools/ui/src/lib/components/app/models/ModelId.svelte +++ b/tools/ui/src/lib/components/app/models/ModelId.svelte @@ -8,6 +8,7 @@ hideOrgName?: boolean; showRaw?: boolean; hideQuantization?: boolean; + hideTags?: boolean; aliases?: string[]; tags?: string[]; class?: string; @@ -17,7 +18,8 @@ modelId, hideOrgName = false, showRaw = undefined, - hideQuantization = false, + hideQuantization, + hideTags, aliases, tags, class: className = '', @@ -31,6 +33,8 @@ let parsed = $derived(ModelsService.parseModelId(modelId)); let resolvedShowRaw = $derived(showRaw ?? (config().showRawModelNames as boolean) ?? false); + let resolvedHideQuantization = $derived(hideQuantization ?? !config().showModelQuantization); + let resolvedHideTags = $derived(hideTags ?? !config().showModelTags); let uniqueAliases = $derived([...new Set(aliases ?? [])]); let uniqueTags = $derived([...new Set([...(parsed.tags ?? []), ...(tags ?? [])])]); @@ -53,7 +57,7 @@ {/if} - {#if parsed.quantization && !hideQuantization} + {#if parsed.quantization && !resolvedHideQuantization} {parsed.quantization} @@ -69,7 +73,7 @@ {/each} {/if} - {#if uniqueTags.length > 0} + {#if uniqueTags.length > 0 && !resolvedHideTags} {#each uniqueTags as tag (tag)} {tag} {/each} diff --git a/tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte b/tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte index 998a6a0fa..0f1fba880 100644 --- a/tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte +++ b/tools/ui/src/lib/components/app/models/ModelsSelectorDropdown.svelte @@ -106,9 +106,7 @@ ]} style="max-width: min(calc(100cqw - 10rem), 20rem)" > - - - + {:else}

No models available.

@@ -120,7 +118,7 @@ - + {#if selectedOption} @@ -144,6 +142,7 @@ modelId={selectedOption.model} class="min-w-0 overflow-hidden" hideOrgName={false} + hideQuantization {...props} /> {/snippet} @@ -158,9 +157,9 @@ {/if} {#if ms.updating || ms.isLoadingModel} - + {:else} - + {/if} @@ -251,7 +250,7 @@ onclick={() => ms.handleOpenChange(true)} disabled={disabled || ms.updating} > - + {#if selectedOption} @@ -262,6 +261,7 @@ modelId={selectedOption.model} class="min-w-0 overflow-hidden" hideOrgName={false} + hideQuantization {...props} /> {/snippet} @@ -274,7 +274,7 @@ {/if} {#if ms.updating} - + {/if} {/if} diff --git a/tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte b/tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte index d38ed8c07..2ddbf2405 100644 --- a/tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte +++ b/tools/ui/src/lib/components/app/models/ModelsSelectorSheet.svelte @@ -6,8 +6,7 @@ DialogModelInformation, ModelId, ModelsSelectorList, - SearchInput, - TruncatedText + SearchInput } from '$lib/components/app'; interface Props { @@ -67,7 +66,7 @@ @@ -168,12 +168,12 @@ onclick={() => ms.handleOpenChange(true)} disabled={disabled || ms.updating} > - + - + {#if ms.updating} - + {/if} {/if} diff --git a/tools/ui/src/lib/components/app/models/index.ts b/tools/ui/src/lib/components/app/models/index.ts index a6ba6817f..3ac6ecb67 100644 --- a/tools/ui/src/lib/components/app/models/index.ts +++ b/tools/ui/src/lib/components/app/models/index.ts @@ -74,8 +74,7 @@ export { default as ModelsSelectorOption } from './ModelsSelectorOption.svelte'; */ export { default as ModelsSelectorSheet } from './ModelsSelectorSheet.svelte'; -/** - * **ModelBadge** - Model name display badge +/** * **ModelBadge** - Model name display badge * * Compact badge showing current model name with package icon. * Only visible in single model mode. Supports tooltip and copy functionality. diff --git a/tools/ui/src/lib/constants/attachment-menu.ts b/tools/ui/src/lib/constants/attachment-menu.ts index 966f46f41..3d7381812 100644 --- a/tools/ui/src/lib/constants/attachment-menu.ts +++ b/tools/ui/src/lib/constants/attachment-menu.ts @@ -79,7 +79,9 @@ export const ATTACHMENT_FILE_ITEMS: AttachmentMenuItem[] = [ } ]; -export const ATTACHMENT_EXTRA_ITEMS: AttachmentMenuItem[] = [ +export const ATTACHMENT_EXTRA_ITEMS: AttachmentMenuItem[] = []; + +export const ATTACHMENT_PROMPT_ITEMS: AttachmentMenuItem[] = [ { id: AttachmentMenuItemId.SYSTEM_MESSAGE, label: 'System Message', @@ -87,10 +89,7 @@ export const ATTACHMENT_EXTRA_ITEMS: AttachmentMenuItem[] = [ enabledWhen: AttachmentItemEnabledWhen.ALWAYS, hasEnabledTooltip: true, action: AttachmentAction.SYSTEM_PROMPT_CLICK - } -]; - -export const ATTACHMENT_MCP_ITEMS: AttachmentMenuItem[] = [ + }, { id: AttachmentMenuItemId.MCP_PROMPT, label: 'MCP Prompt', @@ -98,7 +97,10 @@ export const ATTACHMENT_MCP_ITEMS: AttachmentMenuItem[] = [ enabledWhen: AttachmentItemEnabledWhen.ALWAYS, action: AttachmentAction.MCP_PROMPT_CLICK, visibleWhen: AttachmentItemVisibleWhen.HAS_MCP_PROMPTS_SUPPORT - }, + } +]; + +export const ATTACHMENT_MCP_ITEMS: AttachmentMenuItem[] = [ { id: AttachmentMenuItemId.MCP_RESOURCES, label: 'MCP Resources', diff --git a/tools/ui/src/lib/constants/index.ts b/tools/ui/src/lib/constants/index.ts index eb8591037..54f39feee 100644 --- a/tools/ui/src/lib/constants/index.ts +++ b/tools/ui/src/lib/constants/index.ts @@ -5,6 +5,8 @@ export * from './agentic'; export * from './api-endpoints'; export * from './attachment-labels'; export * from './database'; +export * from './reasoning-effort'; +export * from './reasoning-effort-tokens'; export * from './storage'; export * from './attachment-menu'; export * from './auto-scroll'; diff --git a/tools/ui/src/lib/constants/reasoning-effort-tokens.ts b/tools/ui/src/lib/constants/reasoning-effort-tokens.ts new file mode 100644 index 000000000..059af71de --- /dev/null +++ b/tools/ui/src/lib/constants/reasoning-effort-tokens.ts @@ -0,0 +1,12 @@ +import { ReasoningEffort } from '$lib/enums'; + +/** + * Reasoning effort to token budget mapping. + * Maps the ReasoningEffort enum values to concrete token counts for the server. + */ +export const REASONING_EFFORT_TOKENS: Record = { + [ReasoningEffort.LOW]: 512, + [ReasoningEffort.MEDIUM]: 2048, + [ReasoningEffort.HIGH]: 8192, + [ReasoningEffort.MAX]: -1 // unlimited +}; diff --git a/tools/ui/src/lib/constants/reasoning-effort.ts b/tools/ui/src/lib/constants/reasoning-effort.ts new file mode 100644 index 000000000..d854e912a --- /dev/null +++ b/tools/ui/src/lib/constants/reasoning-effort.ts @@ -0,0 +1,21 @@ +import { ReasoningEffort } from '$lib/enums'; +import type { ReasoningEffortLevel } from '$lib/types'; + +/** + * Reasoning effort UI labels. + * Keys match the ReasoningEffort enum values for type-safe lookups. + */ +export const REASONING_EFFORT_LABELS: Record = { + [ReasoningEffort.LOW]: 'Low', + [ReasoningEffort.MEDIUM]: 'Medium', + [ReasoningEffort.HIGH]: 'High', + [ReasoningEffort.MAX]: 'Max' +}; + +export const REASONING_EFFORT_LEVELS: ReasoningEffortLevel[] = [ + { value: 'off', label: 'Off', isOff: true }, + { value: ReasoningEffort.LOW, label: 'Low' }, + { value: ReasoningEffort.MEDIUM, label: 'Medium' }, + { value: ReasoningEffort.HIGH, label: 'High' }, + { value: ReasoningEffort.MAX, label: 'Max', hasInfo: true } +]; diff --git a/tools/ui/src/lib/constants/settings-keys.ts b/tools/ui/src/lib/constants/settings-keys.ts index 7cdd2db7c..5fff9f94c 100644 --- a/tools/ui/src/lib/constants/settings-keys.ts +++ b/tools/ui/src/lib/constants/settings-keys.ts @@ -29,6 +29,8 @@ export const SETTINGS_KEYS = { ALWAYS_SHOW_SIDEBAR_ON_DESKTOP: 'alwaysShowSidebarOnDesktop', FULL_HEIGHT_CODE_BLOCKS: 'fullHeightCodeBlocks', SHOW_RAW_MODEL_NAMES: 'showRawModelNames', + SHOW_MODEL_QUANTIZATION: 'showModelQuantization', + SHOW_MODEL_TAGS: 'showModelTags', SHOW_SYSTEM_MESSAGE: 'showSystemMessage', // Sampling TEMPERATURE: 'temperature', @@ -64,6 +66,7 @@ export const SETTINGS_KEYS = { // Developer DISABLE_REASONING_PARSING: 'disableReasoningParsing', EXCLUDE_REASONING_FROM_CONTEXT: 'excludeReasoningFromContext', + ENABLE_THINKING: 'enableThinking', SHOW_RAW_OUTPUT_SWITCH: 'showRawOutputSwitch', // PY_INTERPRETER_ENABLED: 'pyInterpreterEnabled', CUSTOM_JSON: 'customJson', diff --git a/tools/ui/src/lib/constants/settings-registry.ts b/tools/ui/src/lib/constants/settings-registry.ts index 20ac33c85..9246b9703 100644 --- a/tools/ui/src/lib/constants/settings-registry.ts +++ b/tools/ui/src/lib/constants/settings-registry.ts @@ -330,6 +330,30 @@ const SETTINGS_REGISTRY: Record = { paramType: SyncableParameterType.BOOLEAN } }, + { + key: SETTINGS_KEYS.SHOW_MODEL_QUANTIZATION, + label: 'Show model quantization information', + help: 'Display quantization badges (e.g. Q8_0, Q4_K_M) next to model names throughout the interface.', + defaultValue: true, + type: SettingsFieldType.CHECKBOX, + section: SETTINGS_SECTION_SLUGS.DISPLAY, + sync: { + serverKey: SETTINGS_KEYS.SHOW_MODEL_QUANTIZATION, + paramType: SyncableParameterType.BOOLEAN + } + }, + { + key: SETTINGS_KEYS.SHOW_MODEL_TAGS, + label: 'Show model tags', + help: 'Display model tags (e.g. "vision", "reasoning") next to model names throughout the interface.', + defaultValue: true, + type: SettingsFieldType.CHECKBOX, + section: SETTINGS_SECTION_SLUGS.DISPLAY, + sync: { + serverKey: SETTINGS_KEYS.SHOW_MODEL_TAGS, + paramType: SyncableParameterType.BOOLEAN + } + }, { key: SETTINGS_KEYS.ALWAYS_SHOW_AGENTIC_TURNS, label: 'Always show agentic turns in conversation', @@ -646,6 +670,14 @@ const SETTINGS_REGISTRY: Record = { paramType: SyncableParameterType.BOOLEAN } }, + { + key: SETTINGS_KEYS.ENABLE_THINKING, + label: 'Enable thinking', + help: 'Enable model thinking/reasoning for each request. When off, the model will skip the thinking phase and go straight to the response.', + defaultValue: false, + type: SettingsFieldType.CHECKBOX, + section: SETTINGS_SECTION_SLUGS.DEVELOPER + }, { key: SETTINGS_KEYS.SHOW_RAW_OUTPUT_SWITCH, label: 'Enable raw output toggle', diff --git a/tools/ui/src/lib/constants/storage.ts b/tools/ui/src/lib/constants/storage.ts index d03254b9c..5d33e82f3 100644 --- a/tools/ui/src/lib/constants/storage.ts +++ b/tools/ui/src/lib/constants/storage.ts @@ -19,6 +19,8 @@ export const CONFIG_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.config`; export const DISABLED_TOOLS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.disabledTools`; export const FAVORITE_MODELS_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.favoriteModels`; export const MCP_DEFAULT_ENABLED_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.mcpDefaultEnabled`; +export const THINKING_ENABLED_DEFAULT_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.thinkingEnabledDefault`; +export const REASONING_EFFORT_DEFAULT_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.reasoningEffortDefault`; export const USER_OVERRIDES_LOCALSTORAGE_KEY = `${STORAGE_APP_NAME}.userOverrides`; // Deprecated old key names (kept for backward compat while users migrate) diff --git a/tools/ui/src/lib/enums/index.ts b/tools/ui/src/lib/enums/index.ts index b80b5b61e..7438a7642 100644 --- a/tools/ui/src/lib/enums/index.ts +++ b/tools/ui/src/lib/enums/index.ts @@ -19,6 +19,8 @@ export { ReasoningFormat } from './chat.enums'; +export { ReasoningEffort } from './reasoning-effort.enums'; + export { FileTypeCategory, FileTypeImage, diff --git a/tools/ui/src/lib/enums/reasoning-effort.enums.ts b/tools/ui/src/lib/enums/reasoning-effort.enums.ts new file mode 100644 index 000000000..dadb0c726 --- /dev/null +++ b/tools/ui/src/lib/enums/reasoning-effort.enums.ts @@ -0,0 +1,10 @@ +/** + * Reasoning effort levels for thinking models. + * These values are sent to the server and mapped to token budgets. + */ +export enum ReasoningEffort { + LOW = 'low', + MEDIUM = 'medium', + HIGH = 'high', + MAX = 'max' +} diff --git a/tools/ui/src/lib/hooks/use-tools-panel.svelte.ts b/tools/ui/src/lib/hooks/use-tools-panel.svelte.ts index 911af322a..9a8acec0f 100644 --- a/tools/ui/src/lib/hooks/use-tools-panel.svelte.ts +++ b/tools/ui/src/lib/hooks/use-tools-panel.svelte.ts @@ -17,6 +17,8 @@ export interface UseToolsPanelReturn { getFavicon(group: { source: ToolSource; label: string }): string | null; isGroupDisabled(group: ToolGroup): boolean; toggleGroupExpanded(label: string): void; + /** Toggle all tools in a group by label (avoids stale group object references). */ + toggleGroupByLabel(label: string): void; handleOpen(): void; } @@ -91,6 +93,13 @@ export function useToolsPanel(): UseToolsPanelReturn { } } + function toggleGroupByLabel(label: string): void { + // Find current group by label to get up-to-date tool references + const group = activeGroups.find((g) => g.label === label); + if (!group) return; + toolsStore.toggleGroup(group); + } + function handleOpen(): void { if (toolsStore.builtinTools.length === 0 && !toolsStore.loading) { toolsStore.fetchBuiltinTools(); @@ -117,6 +126,7 @@ export function useToolsPanel(): UseToolsPanelReturn { getFavicon, isGroupDisabled, toggleGroupExpanded, + toggleGroupByLabel, handleOpen }; } diff --git a/tools/ui/src/lib/services/chat.service.ts b/tools/ui/src/lib/services/chat.service.ts index 09d616bd7..468009a54 100644 --- a/tools/ui/src/lib/services/chat.service.ts +++ b/tools/ui/src/lib/services/chat.service.ts @@ -6,6 +6,7 @@ import { ATTACHMENT_LABEL_MCP_PROMPT, ATTACHMENT_LABEL_MCP_RESOURCE, LEGACY_AGENTIC_REGEX, + REASONING_EFFORT_TOKENS, SETTINGS_KEYS, API_CHAT, API_SLOTS, @@ -162,6 +163,8 @@ export class ChatService { // Config options disableReasoningParsing, excludeReasoningFromContext, + enableThinking, + reasoningEffort, continueFinalMessage } = options; @@ -243,6 +246,18 @@ export class ChatService { ? ReasoningFormat.NONE : ReasoningFormat.AUTO; + const reasoningBudgetTokens = + enableThinking && reasoningEffort ? (REASONING_EFFORT_TOKENS[reasoningEffort] ?? -1) : -1; + + requestBody.chat_template_kwargs = { + ...(requestBody.chat_template_kwargs ?? {}), + enable_thinking: enableThinking + }; + + if (reasoningBudgetTokens >= 0) { + requestBody.thinking_budget_tokens = reasoningBudgetTokens; + } + // arms the budget sampler so reasoning can be ended at runtime via the control endpoint requestBody.reasoning_control = true; diff --git a/tools/ui/src/lib/stores/chat.svelte.ts b/tools/ui/src/lib/stores/chat.svelte.ts index 667fb417c..b899130e5 100644 --- a/tools/ui/src/lib/stores/chat.svelte.ts +++ b/tools/ui/src/lib/stores/chat.svelte.ts @@ -1852,6 +1852,9 @@ class ChatStore { if (currentConfig.excludeReasoningFromContext) apiOptions.excludeReasoningFromContext = true; + apiOptions.enableThinking = conversationsStore.getThinkingEnabled(); + apiOptions.reasoningEffort = conversationsStore.getReasoningEffort(); + if (hasValue(currentConfig.temperature)) apiOptions.temperature = Number(currentConfig.temperature); diff --git a/tools/ui/src/lib/stores/conversations.svelte.ts b/tools/ui/src/lib/stores/conversations.svelte.ts index d6589232f..cab6f59fa 100644 --- a/tools/ui/src/lib/stores/conversations.svelte.ts +++ b/tools/ui/src/lib/stores/conversations.svelte.ts @@ -26,7 +26,7 @@ import { MigrationService } from '$lib/services/migration.service'; import { config } from '$lib/stores/settings.svelte'; import { filterByLeafNodeId, findLeafNode, generateConversationTitle } from '$lib/utils'; import type { McpServerOverride } from '$lib/types/database'; -import { MessageRole, HtmlInputType, FileExtensionText } from '$lib/enums'; +import { MessageRole, HtmlInputType, FileExtensionText, ReasoningEffort } from '$lib/enums'; import { ISO_DATE_TIME_SEPARATOR, ISO_DATE_TIME_SEPARATOR_REPLACEMENT, @@ -38,7 +38,9 @@ import { ISO_TIME_SEPARATOR_REPLACEMENT, NON_ALPHANUMERIC_REGEX, MULTIPLE_UNDERSCORE_REGEX, - MCP_DEFAULT_ENABLED_LOCALSTORAGE_KEY + MCP_DEFAULT_ENABLED_LOCALSTORAGE_KEY, + THINKING_ENABLED_DEFAULT_LOCALSTORAGE_KEY, + REASONING_EFFORT_DEFAULT_LOCALSTORAGE_KEY } from '$lib/constants'; import { ROUTES } from '$lib/constants/routes'; @@ -74,6 +76,12 @@ class ConversationsStore { /** Pending MCP server overrides for new conversations (before first message) */ pendingMcpServerOverrides = $state(ConversationsStore.loadMcpDefaults()); + /** Global (non-conversation-specific) thinking toggle default */ + pendingThinkingEnabled = $state(ConversationsStore.loadThinkingDefaults()); + + /** Global (non-conversation-specific) reasoning effort default */ + pendingReasoningEffort = $state(ConversationsStore.loadReasoningEffortDefault()); + /** Load MCP default overrides from localStorage */ private static loadMcpDefaults(): McpServerOverride[] { if (typeof globalThis.localStorage === 'undefined') return []; @@ -104,6 +112,45 @@ class ConversationsStore { } } + /** Load thinking-enabled default from localStorage */ + private static loadThinkingDefaults(): boolean { + if (typeof globalThis.localStorage === 'undefined') return false; + try { + const raw = localStorage.getItem(THINKING_ENABLED_DEFAULT_LOCALSTORAGE_KEY); + if (!raw) return false; + const parsed = raw === 'true'; + return typeof parsed === 'boolean' ? parsed : false; + } catch { + return false; + } + } + + /** Persist thinking-enabled default to localStorage */ + private saveThinkingDefaults(): void { + if (typeof globalThis.localStorage === 'undefined') return; + localStorage.setItem( + THINKING_ENABLED_DEFAULT_LOCALSTORAGE_KEY, + this.pendingThinkingEnabled ? 'true' : 'false' + ); + } + + /** Load reasoning effort default from localStorage */ + private static loadReasoningEffortDefault(): ReasoningEffort { + if (typeof globalThis.localStorage === 'undefined') return ReasoningEffort.MEDIUM; + try { + const raw = localStorage.getItem(REASONING_EFFORT_DEFAULT_LOCALSTORAGE_KEY); + return (raw as ReasoningEffort) || ReasoningEffort.MEDIUM; + } catch { + return ReasoningEffort.MEDIUM; + } + } + + /** Persist reasoning effort default to localStorage */ + private saveReasoningEffortDefaults(): void { + if (typeof globalThis.localStorage === 'undefined') return; + localStorage.setItem(REASONING_EFFORT_DEFAULT_LOCALSTORAGE_KEY, this.pendingReasoningEffort); + } + /** Callback for title update confirmation dialog */ titleUpdateConfirmationCallback?: (currentTitle: string, newTitle: string) => Promise; @@ -253,6 +300,12 @@ class ConversationsStore { this.pendingMcpServerOverrides = []; } + // Inherit global thinking default into the new conversation + conversation.thinkingEnabled = this.pendingThinkingEnabled; + await DatabaseService.updateConversation(conversation.id, { + thinkingEnabled: this.pendingThinkingEnabled + }); + this.conversations = [conversation, ...this.conversations]; this.activeConversation = conversation; this.activeMessages = []; @@ -276,6 +329,7 @@ class ConversationsStore { } this.pendingMcpServerOverrides = []; + this.pendingThinkingEnabled = false; this.activeConversation = conversation; if (conversation.currNode) { @@ -304,8 +358,9 @@ class ConversationsStore { clearActiveConversation(): void { this.activeConversation = null; this.activeMessages = []; - // reload MCP defaults so new chats inherit persisted state + // reload defaults so new chats inherit persisted state this.pendingMcpServerOverrides = ConversationsStore.loadMcpDefaults(); + this.pendingThinkingEnabled = ConversationsStore.loadThinkingDefaults(); } /** @@ -703,6 +758,84 @@ class ConversationsStore { this.saveMcpDefaults(); } + /** + * Gets the effective thinking-enabled state for the active conversation. + * Returns the conversation override if set, otherwise the global default. + */ + getThinkingEnabled(): boolean { + if (this.activeConversation) { + return this.activeConversation.thinkingEnabled ?? this.pendingThinkingEnabled; + } + return this.pendingThinkingEnabled; + } + + /** + * Sets the thinking-enabled state for the active conversation. + * If no conversation exists, stores the global default. + * @param enabled - The enabled state + */ + async setThinkingEnabled(enabled: boolean): Promise { + if (!this.activeConversation) { + this.pendingThinkingEnabled = enabled; + this.saveThinkingDefaults(); + return; + } + + this.activeConversation = { + ...this.activeConversation, + thinkingEnabled: enabled + }; + + await DatabaseService.updateConversation(this.activeConversation.id, { + thinkingEnabled: enabled + }); + + const convIndex = this.conversations.findIndex((c) => c.id === this.activeConversation!.id); + if (convIndex !== -1) { + this.conversations[convIndex].thinkingEnabled = enabled; + this.conversations = [...this.conversations]; + } + } + + /** + * Gets the effective reasoning effort for the active conversation. + * Returns the conversation override if set, otherwise the global default. + */ + getReasoningEffort(): ReasoningEffort { + if (this.activeConversation) { + return this.activeConversation.reasoningEffort ?? this.pendingReasoningEffort; + } + return this.pendingReasoningEffort; + } + + /** + * Sets the reasoning effort for the active conversation. + * If no conversation exists, stores the global default. + * @param effort - The effort level ('low' | 'medium' | 'high' | 'max') + */ + async setReasoningEffort(effort: ReasoningEffort): Promise { + if (!this.activeConversation) { + this.pendingReasoningEffort = effort; + this.saveReasoningEffortDefaults(); + return; + } + + this.activeConversation = { + ...this.activeConversation, + reasoningEffort: effort + }; + + await DatabaseService.updateConversation(this.activeConversation.id, { + reasoningEffort: effort + }); + + const convIndex = this.conversations.findIndex((c) => c.id === this.activeConversation!.id); + if (convIndex !== -1) { + this.conversations[convIndex].reasoningEffort = effort; + this.conversations = [...this.conversations]; + } + } + /** * Forks a conversation at a specific message, creating a new conversation * containing messages from root up to the target message, then navigates to it. diff --git a/tools/ui/src/lib/stores/models.svelte.ts b/tools/ui/src/lib/stores/models.svelte.ts index bc99d7412..1990ba604 100644 --- a/tools/ui/src/lib/stores/models.svelte.ts +++ b/tools/ui/src/lib/stores/models.svelte.ts @@ -4,6 +4,10 @@ import { ServerModelStatus, ModelModality } from '$lib/enums'; import { ModelsService } from '$lib/services/models.service'; import { PropsService } from '$lib/services/props.service'; import { serverStore, isRouterMode } from '$lib/stores/server.svelte'; +import { + detectThinkingSupport, + detectThinkingSupportWithReason +} from '$lib/utils/chat-template-thinking-detector'; import { TTLCache } from '$lib/utils'; import { MODEL_PROPS_CACHE_TTL_MS, @@ -215,6 +219,67 @@ class ModelsStore { return usage !== undefined && usage.size > 0; } + // + // Thinking Support Detection + // + + /** + * 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. + */ + get supportsThinking(): boolean { + const modelId = this.selectedModelName; + if (!modelId) { + if (!isRouterMode()) { + return detectThinkingSupport(serverStore.props?.chat_template ?? ''); + } + return false; + } + + if (isRouterMode() && !this.modelPropsCache.get(modelId)) { + this.fetchModelProps(modelId); + } + const props = this.getModelProps(modelId); + return detectThinkingSupport(props?.chat_template ?? ''); + } + + /** + * Check if a specific model supports thinking. + * Fetches model props if not cached (in router mode). + */ + checkModelSupportsThinking(modelId: string): boolean { + if (!modelId) return false; + + if (isRouterMode() && !this.modelPropsCache.get(modelId)) { + this.fetchModelProps(modelId); + } + + const props = this.getModelProps(modelId); + return detectThinkingSupport(props?.chat_template ?? ''); + } + + /** + * Detailed thinking support detection result with reason for debugging/UI. + */ + get thinkingSupportDetails(): { supported: boolean; reason: string } { + 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)) { + this.fetchModelProps(modelId); + } + const props = this.getModelProps(modelId); + return detectThinkingSupportWithReason(props?.chat_template ?? ''); + } /** * @@ -362,6 +427,7 @@ class ModelsStore { try { const props = await PropsService.fetchForModel(modelId); this.modelPropsCache.set(modelId, props); + this.propsCacheVersion++; return props; } catch (error) { console.warn(`Failed to fetch props for model ${modelId}:`, error); @@ -755,3 +821,7 @@ export const propsCacheVersion = () => modelsStore.propsCacheVersion; export const singleModelName = () => modelsStore.singleModelName; export const selectedModelContextSize = () => modelsStore.selectedModelContextSize; export const favoriteModelIds = () => modelsStore.favoriteModelIds; +export const supportsThinking = () => modelsStore.supportsThinking; +export const checkModelSupportsThinking = (modelId: string) => + modelsStore.checkModelSupportsThinking(modelId); +export const thinkingSupportDetails = () => modelsStore.thinkingSupportDetails; diff --git a/tools/ui/src/lib/types/database.d.ts b/tools/ui/src/lib/types/database.d.ts index aecc250e8..add628476 100644 --- a/tools/ui/src/lib/types/database.d.ts +++ b/tools/ui/src/lib/types/database.d.ts @@ -1,5 +1,5 @@ import type { ChatMessageTimings, ChatRole, ChatMessageType } from '$lib/types/chat'; -import { AttachmentType } from '$lib/enums'; +import { AttachmentType, ReasoningEffort } from '$lib/enums'; export interface McpServerOverride { serverId: string; @@ -12,6 +12,8 @@ export interface DatabaseConversation { lastModified: number; name: string; mcpServerOverrides?: McpServerOverride[]; + thinkingEnabled?: boolean; + reasoningEffort?: ReasoningEffort; forkedFromConversationId?: string; } diff --git a/tools/ui/src/lib/types/index.ts b/tools/ui/src/lib/types/index.ts index 0eb1e6701..b88de5d66 100644 --- a/tools/ui/src/lib/types/index.ts +++ b/tools/ui/src/lib/types/index.ts @@ -162,3 +162,6 @@ export type { // Tools types export type { ToolEntry, ToolGroup } from './tools'; + +// Reasoning +export type { ReasoningEffortLevel } from './reasoning'; diff --git a/tools/ui/src/lib/types/reasoning.ts b/tools/ui/src/lib/types/reasoning.ts new file mode 100644 index 000000000..b72a3b77f --- /dev/null +++ b/tools/ui/src/lib/types/reasoning.ts @@ -0,0 +1,6 @@ +export interface ReasoningEffortLevel { + value: string; + label: string; + isOff?: boolean; + hasInfo?: boolean; +} diff --git a/tools/ui/src/lib/types/settings.d.ts b/tools/ui/src/lib/types/settings.d.ts index 44e83f8f3..d1cdca957 100644 --- a/tools/ui/src/lib/types/settings.d.ts +++ b/tools/ui/src/lib/types/settings.d.ts @@ -2,7 +2,12 @@ import type { SETTING_CONFIG_DEFAULT, SETTINGS_SECTION_TITLES } from '$lib/const import type { ChatMessagePromptProgress, ChatMessageTimings } from './chat'; import type { OpenAIToolDefinition } from './mcp'; import type { DatabaseMessageExtra } from './database'; -import type { ParameterSource, SyncableParameterType, SettingsFieldType } from '$lib/enums'; +import type { + ParameterSource, + ReasoningEffort, + SyncableParameterType, + SettingsFieldType +} from '$lib/enums'; import type { Icon } from '@lucide/svelte'; import type { Component } from 'svelte'; @@ -65,6 +70,10 @@ export interface SettingsChatServiceOptions { disableReasoningParsing?: boolean; // Strip reasoning content from context before sending excludeReasoningFromContext?: boolean; + // Enable model thinking/reasoning via chat_template_kwargs + enableThinking?: boolean; + // Reasoning effort level (low/medium/high/max) for thinking models + reasoningEffort?: ReasoningEffort; tools?: OpenAIToolDefinition[]; // Generation parameters temperature?: number; diff --git a/tools/ui/src/lib/utils/chat-template-thinking-detector.ts b/tools/ui/src/lib/utils/chat-template-thinking-detector.ts new file mode 100644 index 000000000..da6382f5c --- /dev/null +++ b/tools/ui/src/lib/utils/chat-template-thinking-detector.ts @@ -0,0 +1,86 @@ +/** + * Detects whether a model's chat template supports thinking/reasoning control. + * + * The server "/props" endpoint does NOT expose a supports_thinking flag. + * It is computed internally by common_chat_templates_support_enable_thinking + * in common/chat.cpp. A proper server flag would make this unnecessary. + * + * Detection order (most reliable first): + * 1. Thinking-control Jinja2 variables === pass-through via chat_template_kwargs + * 2. Thinking-control Jinja2 conditionals === template-native on/off logic + * 3. Paired thinking-content tag pairs === models that output special tags + */ + +const THINKING_KWARG_VARS = ['enable_thinking', 'reasoning_effort', 'thinking_budget']; + +/** + * Paired thinking-content tag patterns. + * + * Inspected: llama-cpp-deepseek-r1/v3, nim-nemotron-{3,4}-nano, qwen-qwq-32b, + * qwen-3-32b, google-gemma-4-31b-it, kimikimi-k2-thinking, apertus-8b-instruct, + * mistralai-Mistral-Small-3.2-24B, ByteDance-Seed-OSS. + * + * The self-closing entry is Kimi-K2, Gemma4 fixed-length pair, + * where both tags always appear adjacent with no content between. + */ +const THINKING_TAG_PATTERNS: Array<[string, string | null]> = [ + ['', ''], + ['<|channel>thought', '<|channel|>'], + ['<|think|>', ''], + ['', ''], + ['', null] +]; + +const JINJA_THINKING_CONDITIONALS: RegExp[] = [ + // Matches: {% if enable thinking %}, {% if enable_thinking %}, {% if (enable_thinking is defined) %} + // Handles: underscore-separated (enable_thinking), space-separated (enable thinking), + // and optional parens/brackets before enable (if (enable_thinking ) + /\{%-?\s*if\s+\(?\s*\w*enable[\s_]+\w*(thinking|think|reasoning)/i, + /\{%-?\s*if\s+\w*(thinking|reasoning)\s*(is not|==|!=)/i, + /\{%-?\s*if\s+not\s+\w*enable/i, + /\{%-?\s*if\s+ns\.enable_thinking/i +]; + +/** Guards against false positives: + * - Generic thought keyword (tool descriptions say "chain of thought") + * - Qwen vertical-bar token (used for ALL tool calls, not thinking) + */ +export function detectThinkingSupport(t: string): boolean { + if (!t) return false; + for (const kwarg of THINKING_KWARG_VARS) { + const regex = new RegExp( + `(\\{\\{[^{}]*\\b${kwarg}\\b[^{}]*\\}\\}|\\{%[^{}]*\\b${kwarg}\\b[^{}]*%\\})`, + 'i' + ); + if (regex.test(t)) return true; + } + for (const p of JINJA_THINKING_CONDITIONALS) { + if (p.test(t)) return true; + } + for (const [s, e] of THINKING_TAG_PATTERNS) { + if (t.includes(s) && (!e || t.includes(e))) return true; + } + return false; +} + +export function detectThinkingSupportWithReason(t: string): { supported: boolean; reason: string } { + if (!t) return { supported: false, reason: 'No chat template available' }; + for (const kwarg of THINKING_KWARG_VARS) { + const regex = new RegExp( + `(\\{\\{[^{}]*\\b${kwarg}\\b[^{}]*\\}\\}|\\{%[^{}]*\\b${kwarg}\\b[^{}]*%\\})`, + 'i' + ); + if (regex.test(t)) { + return { supported: true, reason: 'Found: ' + kwarg }; + } + } + for (const p of JINJA_THINKING_CONDITIONALS) { + if (p.test(t)) return { supported: true, reason: 'Found: thinking conditional' }; + } + for (const [s, e] of THINKING_TAG_PATTERNS) { + if (t.includes(s) && (!e || t.includes(e))) { + return { supported: true, reason: 'Found: ' + s + (e ? ' .. ' + e : ' (self)') }; + } + } + return { supported: false, reason: 'No thinking patterns found' }; +} diff --git a/tools/ui/tests/e2e/demo.test.ts b/tools/ui/tests/e2e/demo.test.ts index b7b4bac33..99a3e86da 100644 --- a/tools/ui/tests/e2e/demo.test.ts +++ b/tools/ui/tests/e2e/demo.test.ts @@ -1,6 +1,7 @@ import { expect, test } from '@playwright/test'; -test('home page has expected h1', async ({ page }) => { +test('home page loads correctly', async ({ page }) => { await page.goto('/'); - await expect(page.locator('h1').first()).toBeVisible(); + // Wait for the greeting to become visible (stores need time to initialize) + await expect(page.locator('h1', { hasText: /Hello there/ })).toBeVisible(); }); diff --git a/tools/ui/tests/stories/ChatScreenForm.a11y.stories.svelte b/tools/ui/tests/stories/ChatScreenForm.a11y.stories.svelte index 5d28f420e..46f852c06 100644 --- a/tools/ui/tests/stories/ChatScreenForm.a11y.stories.svelte +++ b/tools/ui/tests/stories/ChatScreenForm.a11y.stories.svelte @@ -44,7 +44,7 @@ await screen.findByRole('menu'); await waitFor(() => { - expect(document.activeElement).toHaveTextContent('Text Files'); + expect(document.activeElement).toHaveTextContent('Add files'); }); }} /> diff --git a/tools/ui/vite.config.ts b/tools/ui/vite.config.ts index f89a689d5..5b57eae3a 100644 --- a/tools/ui/vite.config.ts +++ b/tools/ui/vite.config.ts @@ -10,6 +10,8 @@ import { llamaCppBuildPlugin } from './scripts/vite-plugin-llama-cpp-build'; const __dirname = dirname(fileURLToPath(import.meta.url)); +const SERVER_ORIGIN = import.meta.env?.VITE_PUBLIC_SERVER_ORIGIN || 'http://localhost:8080'; + export default defineConfig({ resolve: { alias: { @@ -75,12 +77,12 @@ export default defineConfig({ server: { proxy: { - '/v1': 'http://localhost:8080', - '/props': 'http://localhost:8080', - '/models': 'http://localhost:8080', - '/tools': 'http://localhost:8080', - '/slots': 'http://localhost:8080', - '/cors-proxy': 'http://localhost:8080' + '/v1': SERVER_ORIGIN, + '/props': SERVER_ORIGIN, + '/models': SERVER_ORIGIN, + '/tools': SERVER_ORIGIN, + '/slots': SERVER_ORIGIN, + '/cors-proxy': SERVER_ORIGIN }, headers: { 'Cross-Origin-Embedder-Policy': 'require-corp', From 69cea5b6692341a51e9705b83743f4be96fcb182 Mon Sep 17 00:00:00 2001 From: Marcos Del Sol Vives Date: Tue, 2 Jun 2026 10:45:25 +0200 Subject: [PATCH 19/20] ui: simplify network error handling (#23431) Previously error to string conversion was split in two different files, with one converting errors into strings, and another function analyzing those strings to generate yet another string. Now the the error handling for network fetches has been centralised and uses directly HTTP error codes whereas possible to generate the human-readable error strings. It also fixes an issue where all JSON errors reported from the backend, such as "Invalid API key", would get turned incorrectly in to "Failed to connect to server" due to poor matching logic in the now-gone getErrorMessage function. --- tools/ui/src/lib/constants/error.ts | 23 +++++++++ tools/ui/src/lib/stores/server.svelte.ts | 30 +----------- tools/ui/src/lib/utils/api-fetch.ts | 61 ++++++++++++++++++++---- 3 files changed, 77 insertions(+), 37 deletions(-) create mode 100644 tools/ui/src/lib/constants/error.ts diff --git a/tools/ui/src/lib/constants/error.ts b/tools/ui/src/lib/constants/error.ts new file mode 100644 index 000000000..4339bd25d --- /dev/null +++ b/tools/ui/src/lib/constants/error.ts @@ -0,0 +1,23 @@ +export const ERROR_MESSAGES = { + NETWORK: { + GENERIC: 'Failed to connect to server', + NXDOMAIN: 'Server not found - check server address', + REFUSED: 'Connection refused - server may be offline', + TIMEOUT: 'Request timed out', + UNREACHABLE: 'Server is not running or unreachable' + }, + HTTP: { + GENERIC: 'Request failed', + ACCESS_DENIED: 'Access denied', + INTERNAL_ERROR: 'Server error - check server logs', + NOT_FOUND: 'Not found', + TEMPORARILY_UNAVAILABLE: 'Server temporarily unavailable' + } +}; + +export const HTTP_CODE_TO_STRING: Record = { + 401: ERROR_MESSAGES.HTTP.ACCESS_DENIED, + 403: ERROR_MESSAGES.HTTP.ACCESS_DENIED, + 500: ERROR_MESSAGES.HTTP.INTERNAL_ERROR, + 503: ERROR_MESSAGES.HTTP.TEMPORARILY_UNAVAILABLE +}; diff --git a/tools/ui/src/lib/stores/server.svelte.ts b/tools/ui/src/lib/stores/server.svelte.ts index dfcb9b2bb..d9a9f855a 100644 --- a/tools/ui/src/lib/stores/server.svelte.ts +++ b/tools/ui/src/lib/stores/server.svelte.ts @@ -82,8 +82,8 @@ class ServerStore { this.props = props; this.error = null; this.detectRole(props); - } catch (error) { - this.error = this.getErrorMessage(error); + } catch (error: unknown) { + this.error = error instanceof Error ? error.message : String(error); console.error('Error fetching server properties:', error); } finally { this.loading = false; @@ -95,32 +95,6 @@ class ServerStore { await fetchPromise; } - private getErrorMessage(error: unknown): string { - if (error instanceof Error) { - const message = error.message || ''; - - if (error.name === 'TypeError' && message.includes('fetch')) { - return 'Server is not running or unreachable'; - } else if (message.includes('ECONNREFUSED')) { - return 'Connection refused - server may be offline'; - } else if (message.includes('ENOTFOUND')) { - return 'Server not found - check server address'; - } else if (message.includes('ETIMEDOUT')) { - return 'Request timed out'; - } else if (message.includes('503')) { - return 'Server temporarily unavailable'; - } else if (message.includes('500')) { - return 'Server error - check server logs'; - } else if (message.includes('404')) { - return 'Server endpoint not found'; - } else if (message.includes('403') || message.includes('401')) { - return 'Access denied'; - } - } - - return 'Failed to connect to server'; - } - clear(): void { this.props = null; this.error = null; diff --git a/tools/ui/src/lib/utils/api-fetch.ts b/tools/ui/src/lib/utils/api-fetch.ts index 80781b98e..82a9383dd 100644 --- a/tools/ui/src/lib/utils/api-fetch.ts +++ b/tools/ui/src/lib/utils/api-fetch.ts @@ -1,6 +1,7 @@ import { base } from '$app/paths'; import { getJsonHeaders, getAuthHeaders } from './api-headers'; import { UrlProtocol } from '$lib/enums'; +import { ERROR_MESSAGES, HTTP_CODE_TO_STRING } from '$lib/constants/error'; /** * API Fetch Utilities @@ -54,10 +55,15 @@ export async function apiFetch(path: string, options: ApiFetchOptions = {}): ? path : `${base}${path}`; - const response = await fetch(url, { - ...fetchOptions, - headers - }); + let response; + try { + response = await fetch(url, { + ...fetchOptions, + headers + }); + } catch (e) { + throw new Error(beautifyNetworkError(e)); + } if (!response.ok) { const errorMessage = await parseErrorMessage(response); @@ -101,10 +107,15 @@ export async function apiFetchWithParams( const baseHeaders = authOnly ? getAuthHeaders() : getJsonHeaders(); const headers = { ...baseHeaders, ...customHeaders }; - const response = await fetch(url.toString(), { - ...fetchOptions, - headers - }); + let response; + try { + response = await fetch(url.toString(), { + ...fetchOptions, + headers + }); + } catch (e) { + throw new Error(beautifyNetworkError(e)); + } if (!response.ok) { const errorMessage = await parseErrorMessage(response); @@ -154,5 +165,37 @@ async function parseErrorMessage(response: Response): Promise { // JSON parsing failed, use status text } - return `Request failed: ${response.status} ${response.statusText}`; + const httpErrorStr = HTTP_CODE_TO_STRING[response.status]; + if (httpErrorStr) { + return httpErrorStr; + } + + return `${ERROR_MESSAGES.HTTP.GENERIC}: ${response.status} ${response.statusText}`; +} + +/** + * Converts a network issue into a human-readable message. + * @param throwable - The throwable raised during fetch operation + * @returns Error in an human-readable format + */ +function beautifyNetworkError(throwable: unknown): string { + let message; + if (throwable instanceof Error) { + message = throwable.message; + if (throwable.name === 'TypeError' && message.includes('fetch')) { + return ERROR_MESSAGES.NETWORK.UNREACHABLE; + } + } else { + message = String(throwable); + } + + if (message.includes('ECONNREFUSED')) { + return ERROR_MESSAGES.NETWORK.REFUSED; + } else if (message.includes('ENOTFOUND')) { + return ERROR_MESSAGES.NETWORK.NXDOMAIN; + } else if (message.includes('ETIMEDOUT')) { + return ERROR_MESSAGES.NETWORK.TIMEOUT; + } + + return `${ERROR_MESSAGES.NETWORK.GENERIC} (${message})`; } From d5ab0834ab0960272a2230a6ec2e884c7a65d192 Mon Sep 17 00:00:00 2001 From: Mikhail Podvitskii Date: Tue, 2 Jun 2026 11:40:22 +0200 Subject: [PATCH 20/20] docs : update HOWTO-add-model.md (#23883) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * docs: update HOWTO-add-model.md with new model registration and graph-building instructions * docs: improve formatting in HOWTO-add-model.md * Update docs/development/HOWTO-add-model.md Co-authored-by: Sigbjørn Skjæret --------- Co-authored-by: Sigbjørn Skjæret --- docs/development/HOWTO-add-model.md | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/docs/development/HOWTO-add-model.md b/docs/development/HOWTO-add-model.md index 5390e42ba..ef2b37088 100644 --- a/docs/development/HOWTO-add-model.md +++ b/docs/development/HOWTO-add-model.md @@ -25,7 +25,7 @@ The convert script reads the model configuration, tokenizer, tensor names+data a The required steps to implement for an HF model are: -1. Define the model `ModelBase.register` annotation in a new `TextModel` or `MmprojModel` subclass, example: +1. Define the model `ModelBase.register` annotation in a new `TextModel` or `MmprojModel` subclass in the [conversion](/conversion) folder, example: ```python @ModelBase.register("MyModelForCausalLM") @@ -98,7 +98,7 @@ The model params and tensors layout must be defined in `llama.cpp` source files: 1. Define a new `llm_arch` enum value in `src/llama-arch.h`. 2. In `src/llama-arch.cpp`: - Add the architecture name to the `LLM_ARCH_NAMES` map. - - Add the list of model tensors to `llm_get_tensor_names` (you may also need to update `LLM_TENSOR_NAMES`) + - You may also need to update `LLM_KV_NAMES`, `LLM_TENSOR_NAMES` and `LLM_TENSOR_INFOS` 3. Add any non-standard metadata loading in the `llama_model_loader` constructor in `src/llama-model-loader.cpp`. 4. If the model has a RoPE operation, add a case for the architecture in `llama_model_rope_type` function in `src/llama-model.cpp`. @@ -106,10 +106,11 @@ NOTE: The dimensions in `ggml` are typically in the reverse order of the `pytorc ### 3. Build the GGML graph implementation -This is the funniest part, you have to provide the inference graph implementation of the new model architecture in `src/llama-model.cpp`. -Create a new struct that inherits from `llm_graph_context` and implement the graph-building logic in its constructor. -Have a look at existing implementations like `llm_build_llama`, `llm_build_dbrx` or `llm_build_bert`. -Then, in the `llama_model::build_graph` method, add a case for your architecture to instantiate your new graph-building struct. +This is the funniest part, you have to provide the inference graph implementation of the new model architecture in `src/llama-model.cpp`: +1. Create a new struct that inherits from `llama_model_base`. +2. Implement the graph-building logic in its `build_arch_graph` method. +3. The `build_arch_graph` method should return a constructed graph (inherited from `llm_graph_context`). Have a look at existing implementations like `llama_model_llama`, `llama_model_dbrx` or `llama_model_bert`. +4. Then, in the `llama_model_mapping` function, add a case for your architecture to instantiate your new graph-building struct. Some `ggml` backends do not support all operations. Backend implementations can be added in a separate PR.