From 78d2f524682d9fee790a6460c93d018dafeb5229 Mon Sep 17 00:00:00 2001 From: fairydreaming <166155368+fairydreaming@users.noreply.github.com> Date: Sun, 5 Jul 2026 17:26:24 +0200 Subject: [PATCH 01/14] cuda : concat implementation for quantized types (#25303) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * cuda : concat implementation for quantized types * chore : apply am17an clever suggestion to shorten the code --------- Co-authored-by: Stanisław Szymczyk --- ggml/src/ggml-cuda/concat.cu | 52 ++++++++++++++++++++------------- ggml/src/ggml-cuda/ggml-cuda.cu | 24 +++++++++++---- 2 files changed, 50 insertions(+), 26 deletions(-) diff --git a/ggml/src/ggml-cuda/concat.cu b/ggml/src/ggml-cuda/concat.cu index 8d557092b..276ee64e8 100644 --- a/ggml/src/ggml-cuda/concat.cu +++ b/ggml/src/ggml-cuda/concat.cu @@ -152,8 +152,8 @@ static void concat_cuda(const ggml_tensor * src0, const ggml_tensor * src1, ggml src0_d + i3*(src0->nb[3] / sizeof(T)), src1_d + i3*(src1->nb[3] / sizeof(T)), dst_d + i3*( dst->nb[3] / sizeof(T)), - src0->ne[0], src0->ne[1], src0->ne[2], - dst->ne[0], dst->ne[1], dst->ne[2], dim, stream); + ggml_row_size(src0->type, src0->ne[0])/sizeof(T), src0->ne[1], src0->ne[2], + ggml_row_size(dst->type, dst->ne[0])/sizeof(T), dst->ne[1], dst->ne[2], dim, stream); } } else { const size_t size0 = ggml_nbytes(src0); @@ -163,6 +163,8 @@ static void concat_cuda(const ggml_tensor * src0, const ggml_tensor * src1, ggml CUDA_CHECK(cudaMemcpyAsync((char *) dst->data + size0, src1->data, size1, cudaMemcpyDeviceToDevice, stream)); } } else { + GGML_ASSERT(!ggml_is_quantized(src0->type)); + dim3 grid_dim(dst->ne[1], dst->ne[2], dst->ne[3]); auto launch_kernel = [&](auto dim) { concat_non_cont<<>>( @@ -204,24 +206,34 @@ void ggml_cuda_op_concat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { GGML_ASSERT(src0->type == src1->type); GGML_ASSERT(dst->type == src0->type); - GGML_ASSERT(!ggml_is_quantized(src0->type)); - GGML_ASSERT(ggml_blck_size(src0->type) == 1); - switch (ggml_type_size(src0->type)) { - case 1: - concat_cuda(src0, src1, dst, dim, stream); - break; - case 2: - concat_cuda(src0, src1, dst, dim, stream); - break; - case 4: - concat_cuda(src0, src1, dst, dim, stream); - break; - case 8: - concat_cuda(src0, src1, dst, dim, stream); - break; - default: - GGML_ABORT("Unsupported type size: %zu", ggml_type_size(src0->type)); - break; + if (ggml_is_quantized(src0->type)) { + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(src1)); + GGML_ASSERT(src0->ne[0] % ggml_blck_size(src0->type) == 0); + GGML_ASSERT(src1->ne[0] % ggml_blck_size(src1->type) == 0); + + // if tensors are contiguous and ne[0] is multiple of the block size we can concat both tensors as byte tensors + concat_cuda(src0, src1, dst, dim, stream); + } else { + GGML_ASSERT(ggml_blck_size(src0->type) == 1); + + switch (ggml_type_size(src0->type)) { + case 1: + concat_cuda(src0, src1, dst, dim, stream); + break; + case 2: + concat_cuda(src0, src1, dst, dim, stream); + break; + case 4: + concat_cuda(src0, src1, dst, dim, stream); + break; + case 8: + concat_cuda(src0, src1, dst, dim, stream); + break; + default: + GGML_ABORT("Unsupported type size: %zu", ggml_type_size(src0->type)); + break; + } } } diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 78d2218e5..83749f094 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5387,12 +5387,24 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g ggml_type src1_type = op->src[1]->type; return src0_type == src1_type && src0_type == op->type && - !ggml_is_quantized(src0_type) && - ggml_blck_size(src0_type) == 1 && - (ggml_type_size(src0_type) == 1 || - ggml_type_size(src0_type) == 2 || - ggml_type_size(src0_type) == 4 || - ggml_type_size(src0_type) == 8); + ( + ( + ggml_is_quantized(src0_type) && + ggml_is_contiguous(op->src[0]) && + ggml_is_contiguous(op->src[1]) && + op->src[0]->ne[0] % ggml_blck_size(src0_type) == 0 && + op->src[1]->ne[0] % ggml_blck_size(src0_type) == 0 + ) || ( + !ggml_is_quantized(src0_type) && + ggml_blck_size(src0_type) == 1 && + ( + ggml_type_size(src0_type) == 1 || + ggml_type_size(src0_type) == 2 || + ggml_type_size(src0_type) == 4 || + ggml_type_size(src0_type) == 8 + ) + ) + ); } break; case GGML_OP_CONV_TRANSPOSE_1D: { From 7a63fdede1aca8b29e64146f33ae03af6c3ee3cd Mon Sep 17 00:00:00 2001 From: Vexxie Date: Sun, 5 Jul 2026 18:10:09 +0100 Subject: [PATCH 02/14] ggml: Update VMM Pool allocation ggml-cuda.cu - Turing P2P access fix (fixes #24489) (#24491) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Update ggml-cuda.cu - Turing P2P access fix. * Add original code as fallback behaviour when NCCL or P2P is not set/true. * Update ggml/src/ggml-cuda/ggml-cuda.cu to add comment as per suggestion Co-authored-by: Johannes Gäßler --------- Co-authored-by: Johannes Gäßler --- ggml/src/ggml-cuda/ggml-cuda.cu | 42 ++++++++++++++++++++++++++++----- 1 file changed, 36 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 83749f094..cda31bbfb 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -543,12 +543,42 @@ struct ggml_cuda_pool_vmm : public ggml_cuda_pool { // the memory allocation handle is no longer needed after mapping CU_CHECK(cuMemRelease(handle)); - // set access - CUmemAccessDesc access = {}; - access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; - access.location.id = device; - access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; - CU_CHECK(cuMemSetAccess((CUdeviceptr)((char *)(pool_addr) + pool_size), reserve_size, &access, 1)); + // VMM Bug fix for P2P access if GGML_CUDA_P2P is set, or if NCCL build + bool use_peer_access = getenv("GGML_CUDA_P2P") != nullptr; +#if defined(GGML_USE_NCCL) + use_peer_access = true; +#endif // defined(GGML_USE_NCCL) + + if (use_peer_access) { + // NCCL implicitly enables peer access (cudaDeviceEnablePeerAccess), and + // GGML_CUDA_P2P enables it explicitly. Unlike cudaMalloc buffers, VMM + // allocations do not become peer-accessible from that alone, so access + // must be granted explicitly here. + std::vector access_descs; + const int device_count = ggml_cuda_info().device_count; + for (int id = 0; id < device_count; ++id) { + if (id != device) { + int can_access_peer = 0; + CUDA_CHECK(cudaDeviceCanAccessPeer(&can_access_peer, id, device)); + if (!can_access_peer) { + continue; + } + } + CUmemAccessDesc access = {}; + access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; + access.location.id = id; + access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; + access_descs.push_back(access); + } + CU_CHECK(cuMemSetAccess(start_ptr, reserve_size, access_descs.data(), access_descs.size())); + } else { + // set access for non P2P + CUmemAccessDesc access = {}; + access.location.type = CU_MEM_LOCATION_TYPE_DEVICE; + access.location.id = device; + access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE; + CU_CHECK(cuMemSetAccess(start_ptr, reserve_size, &access, 1)); + } // add to the pool pool_size += reserve_size; From 4b2a0cdee141d7906f661ba573e4b455f684bbc3 Mon Sep 17 00:00:00 2001 From: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com> Date: Mon, 6 Jul 2026 01:56:11 +0800 Subject: [PATCH 03/14] ggml : fix tensor-parallel + -ncmoe crash on MoE models (#25028) Tensor parallelism (-sm tensor) combined with -ncmoe (CPU-offloaded MoE experts) aborts during warm-up on MoE models with GGML_ASSERT(ggml_is_contiguous(tensor)) in ggml-backend-meta.cpp. The failing tensor is the MoE router output (ffn_moe_topk): it is mirrored (GGML_BACKEND_SPLIT_AXIS_MIRRORED, replicated across backends since routing must be identical) and happens to be a non-contiguous view. ggml_backend_meta_buffer_{get,set}_tensor asserted contiguity before consulting the split state, so a mirrored non-contiguous tensor tripped the assert even though the GGML_BACKEND_SPLIT_AXIS_MIRRORED case right below already handles it. Move the split-state lookup above the assert and allow the mirrored case in both get_tensor and set_tensor. Diagnosis credit to the reporter (@nathanmp). Fixes #24886 Signed-off-by: liminfei-amd <91481003+liminfei-amd@users.noreply.github.com> --- ggml/src/ggml-backend-meta.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index 0a36f0990..7bd329164 100644 --- a/ggml/src/ggml-backend-meta.cpp +++ b/ggml/src/ggml-backend-meta.cpp @@ -1245,9 +1245,8 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor(ggml_backend_buffer static void ggml_backend_meta_buffer_set_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset, size_t size) { const size_t n_bufs = ggml_backend_meta_buffer_n_bufs(buffer); - GGML_ASSERT(ggml_is_contiguous(tensor)); - const ggml_backend_meta_split_state split_state = ggml_backend_meta_get_split_state(tensor, /*assume_sync =*/ false); + GGML_ASSERT(ggml_is_contiguous(tensor) || split_state.axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED); if (split_state.n_segments != 1 || split_state.nr[0] != 1) { GGML_ASSERT(split_state.axis >= 0 && split_state.axis < GGML_MAX_DIMS); @@ -1360,9 +1359,8 @@ static void ggml_backend_meta_buffer_set_tensor(ggml_backend_buffer_t buffer, gg static void ggml_backend_meta_buffer_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size) { const size_t n_bufs = ggml_backend_meta_buffer_n_bufs(buffer); - GGML_ASSERT(ggml_is_contiguous(tensor)); - const ggml_backend_meta_split_state split_state = ggml_backend_meta_get_split_state(tensor, /*assume_sync =*/ false); + GGML_ASSERT(ggml_is_contiguous(tensor) || split_state.axis == GGML_BACKEND_SPLIT_AXIS_MIRRORED); if (split_state.n_segments != 1 || split_state.nr[0] != 1) { GGML_ASSERT(split_state.axis >= 0 && split_state.axis < GGML_MAX_DIMS); From 3e5036fbfb1613dc19034844c1989aa645183164 Mon Sep 17 00:00:00 2001 From: Eve <139727413+netrunnereve@users.noreply.github.com> Date: Sun, 5 Jul 2026 18:38:47 +0000 Subject: [PATCH 04/14] abort if we see a multi buffer (#25276) --- ggml/src/ggml-backend-meta.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index 7bd329164..1f29ec867 100644 --- a/ggml/src/ggml-backend-meta.cpp +++ b/ggml/src/ggml-backend-meta.cpp @@ -1144,6 +1144,11 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor_impl(ggml_backend_m ggml_context * simple_ctx = stc.ctxs[j].get(); ggml_backend_buffer_t simple_buf = buf_ctx->bufs[j].get(); + if ((simple_buf != nullptr) && ggml_backend_buffer_is_multi_buffer(simple_buf)) { + // see https://github.com/ggml-org/llama.cpp/issues/22197 + GGML_ABORT("multi buffers are not supported by the meta backend"); + } + if (split_dim >= 0 && split_dim < GGML_MAX_DIMS) { // TODO: the following assert fails for llama-parallel even though the results are correct: // GGML_ASSERT(ggml_is_contiguously_allocated(tensor)); From 2da668617612d2df773f966e3b0ee22dc2beef7b Mon Sep 17 00:00:00 2001 From: Al G Date: Sun, 5 Jul 2026 19:39:36 +0100 Subject: [PATCH 05/14] Fix stale tensor-split params for draft models (#24814) * meta: fix tensor split metadata for GQA attention * Tidied the code a bit to match existing style * Revert "Tidied the code a bit to match existing style" This reverts commit b90c6c6300091fe09e2350a3d4edcfcf15db8d2e. * Reverted the ggml-backend-meta asset hack. --- src/llama-model.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/llama-model.cpp b/src/llama-model.cpp index e07f6e986..8d68cff45 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1012,9 +1012,17 @@ struct llama_model::impl { std::vector dev_layer; bool has_tensor_overrides; + + std::vector tensor_split_owned; }; llama_model::llama_model(const llama_model_params & params) : params(params), pimpl(std::make_unique()) { + if (params.tensor_split != nullptr) { + // llama_model_params stores tensor_split as a borrowed pointer, but the model + // may need it later for tensor-parallel KV-cache split metadata. + pimpl->tensor_split_owned.assign(params.tensor_split, params.tensor_split + llama_max_devices()); + this->params.tensor_split = pimpl->tensor_split_owned.data(); + } pimpl->has_tensor_overrides = params.tensor_buft_overrides && params.tensor_buft_overrides[0].pattern; } From 72874f559c598b8f89fbb24864868337cf5afb4c Mon Sep 17 00:00:00 2001 From: adavyas <121313528+adavyas@users.noreply.github.com> Date: Sun, 5 Jul 2026 20:49:06 -0700 Subject: [PATCH 06/14] ggml-cuda: optimize conv_transpose_1d indexing (#25310) --- ggml/src/ggml-cuda/conv-transpose-1d.cu | 26 +++++++++++++------------ 1 file changed, 14 insertions(+), 12 deletions(-) diff --git a/ggml/src/ggml-cuda/conv-transpose-1d.cu b/ggml/src/ggml-cuda/conv-transpose-1d.cu index 8418ba667..ebf2aa804 100644 --- a/ggml/src/ggml-cuda/conv-transpose-1d.cu +++ b/ggml/src/ggml-cuda/conv-transpose-1d.cu @@ -11,30 +11,32 @@ static __global__ void conv_transpose_1d_kernel( return; } - int out_index = global_index / dst_ne0; + int out_t = global_index % dst_ne0; + int out_ch = (global_index / dst_ne0) % dst_ne1; + int plane = global_index / (dst_ne0 * dst_ne1); float accumulator = 0; for (int c = 0; c < src0_ne2; c++) { - int idx = global_index % dst_ne0; + int kernel_offset = src0_ne0 * (out_ch + src0_ne1 * c); + int input_offset = src1_ne0 * (c + src1_ne1 * plane); - int kernel_offset = (src0_ne0 * src0_ne1 * c) + (out_index * src0_ne0); - int input_offset = src1_ne0 * c; - - for (int i = 0; i < src1_ne0; i++) { - if (!(idx >= i*s0 && idx < i*s0 + src0_ne0)) { + for (int k = 0; k < src0_ne0; k++) { + int input_numer = out_t + p0 - k*d0; + if (input_numer < 0 || input_numer % s0 != 0) { continue; } - int weight_idx = idx - i*s0; - float kernel_weight = src0[kernel_offset + weight_idx]; - float input_value = src1[input_offset+i]; + int input_t = input_numer / s0; + if (input_t >= src1_ne0) { + continue; + } - accumulator += kernel_weight * input_value; + accumulator += src0[kernel_offset + k] * src1[input_offset + input_t]; } } dst[global_index] = accumulator; - GGML_UNUSED_VARS(p0, d0, src0_ne3, src1_ne3, dst_ne3, src1_ne1, dst_ne1, src1_ne2, dst_ne2); + GGML_UNUSED_VARS(src0_ne3, src1_ne2, src1_ne3, dst_ne2, dst_ne3); } static void conv_transpose_1d_f32_f32_cuda( From 898b08854d64befbbb7f6e22a5e3052d65e94f4b Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Mon, 6 Jul 2026 08:41:39 +0200 Subject: [PATCH 07/14] ui: fake 200 for proxy DELETE req (#25298) --- tools/ui/src/lib/services/mcp.service.ts | 24 ++++++++++++++++++++++ tools/ui/tests/unit/mcp-service.test.ts | 26 ++++++++++++++++++++++++ 2 files changed, 50 insertions(+) diff --git a/tools/ui/src/lib/services/mcp.service.ts b/tools/ui/src/lib/services/mcp.service.ts index 443298917..ae98632a6 100644 --- a/tools/ui/src/lib/services/mcp.service.ts +++ b/tools/ui/src/lib/services/mcp.service.ts @@ -314,6 +314,30 @@ export class MCPService { ) ); + if (method === 'DELETE' && url.includes(CORS_PROXY_ENDPOINT)) { + const response = new Response(null, { status: 200, statusText: 'OK' }); + + logIfEnabled( + this.createLog( + MCPConnectionPhase.INITIALIZING, + `HTTP 200 ${method} ${url} (fake response)`, + MCPLogLevel.INFO, + { + response: { + url, + status: response.status, + statusText: response.statusText, + durationMs: 0, + isFake: true + } + } + ) + ); + + // fake response, bypass real fetch() + return response; + } + try { const response = await fetch(input, { ...baseInit, diff --git a/tools/ui/tests/unit/mcp-service.test.ts b/tools/ui/tests/unit/mcp-service.test.ts index 1f6fdda37..d3af90681 100644 --- a/tools/ui/tests/unit/mcp-service.test.ts +++ b/tools/ui/tests/unit/mcp-service.test.ts @@ -154,6 +154,32 @@ describe('MCPService', () => { }); }); + it('DELETE request with CORS proxy should return a fake 200 response', async () => { + const logs: MCPConnectionLog[] = []; + const fetchMock = vi.fn(); + + vi.stubGlobal('fetch', fetchMock); + + const config: MCPServerConfig = { + url: 'https://example.com/mcp', + transport: MCPTransportType.STREAMABLE_HTTP, + useProxy: true + }; + + const controller = createDiagnosticFetch(config, (log) => logs.push(log), {}, true); + + const response = await controller.fetch( + 'http://localhost:8080/cors-proxy?url=https%3A%2F%2Fexample.com%2Fmcp', + { method: 'DELETE' } + ); + + expect(fetchMock).not.toHaveBeenCalled(); + expect(response.status).toBe(200); + expect(logs.at(-1)?.details).toMatchObject({ + response: { status: 200, isFake: true } + }); + }); + it('partially redacts mcp-session-id in diagnostic request and response logs', async () => { const logs: MCPConnectionLog[] = []; const response = new Response('{}', { From d06ddd35896e2748fc95ccdb46fb260c5af2c029 Mon Sep 17 00:00:00 2001 From: a-huk <56552991+a-huk@users.noreply.github.com> Date: Mon, 6 Jul 2026 09:02:26 +0200 Subject: [PATCH 08/14] ggml-hip: enable -ffast-math for HIP builds (#23862) --- ggml/src/ggml-hip/CMakeLists.txt | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt index a7d4e0ea2..72850f26d 100644 --- a/ggml/src/ggml-hip/CMakeLists.txt +++ b/ggml/src/ggml-hip/CMakeLists.txt @@ -155,3 +155,5 @@ if (GGML_HIP_RCCL) endif() target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas) + +target_compile_options(ggml-hip PRIVATE "$<$:-ffast-math>") From 48719618e8321e64d4dddaaa56249671295f8a72 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adrien=20Gallou=C3=ABt?= Date: Mon, 6 Jul 2026 09:53:35 +0200 Subject: [PATCH 09/14] scripts : use HF_TOKEN when downloading UI assets (#25280) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Adrien Gallouët --- scripts/ui-assets.cmake | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/scripts/ui-assets.cmake b/scripts/ui-assets.cmake index 349fa9bf8..dc0417ea0 100644 --- a/scripts/ui-assets.cmake +++ b/scripts/ui-assets.cmake @@ -186,6 +186,12 @@ function(hf_download version out_var out_resolved) set(archive "${UI_BINARY_DIR}/dist.tar.gz") + # Use HF_TOKEN to benefit from higher rate limits + set(auth_headers "") + if(DEFINED ENV{HF_TOKEN} AND NOT "$ENV{HF_TOKEN}" STREQUAL "") + list(APPEND auth_headers "HTTPHEADER" "Authorization: Bearer $ENV{HF_TOKEN}") + endif() + set(candidates "") if(NOT "${version}" STREQUAL "") list(APPEND candidates "${version}") @@ -198,7 +204,7 @@ function(hf_download version out_var out_resolved) message(STATUS "UI: downloading from ${resolved}: ${base}/dist.tar.gz") file(DOWNLOAD "${base}/dist.tar.gz?download=true" "${archive}" - STATUS status TIMEOUT 300 + STATUS status TIMEOUT 300 ${auth_headers} ) list(GET status 0 rc) if(NOT rc EQUAL 0) @@ -208,7 +214,7 @@ function(hf_download version out_var out_resolved) endif() file(DOWNLOAD "${base}/dist.tar.gz.sha256?download=true" "${archive}.sha256" - STATUS status TIMEOUT 30 + STATUS status TIMEOUT 30 ${auth_headers} ) list(GET status 0 rc) if(NOT rc EQUAL 0) From d80e87850173a674ea2673fd003ac54ac93774c0 Mon Sep 17 00:00:00 2001 From: Pascal Date: Mon, 6 Jul 2026 10:30:07 +0200 Subject: [PATCH 10/14] ui: restore Ctrl+B sidebar toggle shortcut (#25307) --- .../navigation/SidebarNavigation/SidebarNavigation.svelte | 5 ++++- tools/ui/src/lib/enums/keyboard.enums.ts | 1 + tools/ui/src/lib/hooks/use-keyboard-shortcuts.svelte.ts | 6 ++++++ 3 files changed, 11 insertions(+), 1 deletion(-) diff --git a/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte b/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte index fe503f53b..a23f4682e 100644 --- a/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte +++ b/tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte @@ -27,7 +27,10 @@ let { onSearchClick = () => {} }: Props = $props(); - const { handleKeydown } = useKeyboardShortcuts({ activateSearchMode: () => onSearchClick() }); + const { handleKeydown } = useKeyboardShortcuts({ + activateSearchMode: () => onSearchClick(), + toggleSidebar: () => toggleExpandedMode() + }); let isExpandedMode = $state(false); let hoveredTooltip = $state(null); diff --git a/tools/ui/src/lib/enums/keyboard.enums.ts b/tools/ui/src/lib/enums/keyboard.enums.ts index 46cd4a776..735d3e4b4 100644 --- a/tools/ui/src/lib/enums/keyboard.enums.ts +++ b/tools/ui/src/lib/enums/keyboard.enums.ts @@ -9,6 +9,7 @@ export enum KeyboardKey { ARROW_LEFT = 'ArrowLeft', ARROW_RIGHT = 'ArrowRight', TAB = 'Tab', + B_LOWER = 'b', D_LOWER = 'd', D_UPPER = 'D', E_UPPER = 'E', diff --git a/tools/ui/src/lib/hooks/use-keyboard-shortcuts.svelte.ts b/tools/ui/src/lib/hooks/use-keyboard-shortcuts.svelte.ts index 05966a1a1..61df30b79 100644 --- a/tools/ui/src/lib/hooks/use-keyboard-shortcuts.svelte.ts +++ b/tools/ui/src/lib/hooks/use-keyboard-shortcuts.svelte.ts @@ -9,6 +9,7 @@ interface KeyboardShortcutsCallbacks { deleteActiveConversation?: () => void; navigateToPrevConversation?: () => void; navigateToNextConversation?: () => void; + toggleSidebar?: () => void; } export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) { @@ -21,6 +22,11 @@ export function useKeyboardShortcuts(callbacks: KeyboardShortcutsCallbacks) { callbacks.onSearchActivated?.(); } + if (isCmdOrCtrl && event.key === KeyboardKey.B_LOWER) { + event.preventDefault(); + callbacks.toggleSidebar?.(); + } + if ( isCmdOrCtrl && event.shiftKey && From 86961efd5675279a4ec2f3d1a3c7d6e3803c1935 Mon Sep 17 00:00:00 2001 From: hokanosekai <69720899+hokanosekai@users.noreply.github.com> Date: Mon, 6 Jul 2026 10:35:57 +0200 Subject: [PATCH 11/14] vulkan: fix 32-bit integer overflow in CEIL_DIV (#25245) --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index c0ab9f1c6..3e38c17de 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -129,7 +129,7 @@ typedef struct VkPhysicalDeviceShaderMixedFloatDotProductFeaturesVALVE { #endif #define ROUNDUP_POW2(M, N) (((M) + (N) - 1) & ~((N) - 1)) -#define CEIL_DIV(M, N) (((M) + (N)-1) / (N)) +#define CEIL_DIV(M, N) (((M) / (N)) + (((M) % (N)) != 0)) static bool is_pow2(uint32_t x) { return x > 1 && (x & (x-1)) == 0; } #define VK_VENDOR_ID_AMD 0x1002 From 3b4fca11ac1b0ded8fabbfad7d1f386d2d24e852 Mon Sep 17 00:00:00 2001 From: shalinib-ibm Date: Mon, 6 Jul 2026 15:48:17 +0530 Subject: [PATCH 12/14] ggml-cpu: Enable tiled matmul on AIX (#25199) The matmul_tiled path uses large local stack buffers for A_pack and B_pack. On AIX this can trigger a segmentation fault, so reduce the buffer footprint there to keep the tiled path usable. Performance Impact: ~ 2x gains in PP_Speed for FP32, Q4_0 and Q8_0 models tested with llama-bench, llama-batched-bench and llama-cli. Models used: Llama3.2 3b Instruct F32, qwen 2.5 3b Q4_0 and Q8_0 --- ggml/src/ggml-cpu/llamafile/sgemm.cpp | 34 ++++++++++++++++----------- 1 file changed, 20 insertions(+), 14 deletions(-) diff --git a/ggml/src/ggml-cpu/llamafile/sgemm.cpp b/ggml/src/ggml-cpu/llamafile/sgemm.cpp index 0b8323e60..5efaaa5b2 100644 --- a/ggml/src/ggml-cpu/llamafile/sgemm.cpp +++ b/ggml/src/ggml-cpu/llamafile/sgemm.cpp @@ -2321,24 +2321,28 @@ class tinyBLAS_Q0_PPC { } void matmul(int64_t m, int64_t n) { - #if defined(_AIX) || defined(__BIG_ENDIAN__) - mnpack(0, m, 0, n); - #else - const int64_t mc = 64; - const int64_t kc = 64; + int64_t mc = 64; int64_t nc = 64; + int64_t kc = 64; + int64_t n_chunk = 64; + #if defined(_AIX) || defined(__BIG_ENDIAN__) + mc = 32; + nc = 32; + kc = 32; + n_chunk = 32 + #endif int64_t n_aligned = 0; - if (n % 64 == 0) { + if (n % n_chunk == 0) { n_aligned = n; } else if (n == 4) { n_aligned = 4; - } else if (n < 64) { + } else if (n < n_chunk) { n_aligned = (n / 8) * 8; } else { - n_aligned = (n / 64) * 64; + n_aligned = (n / n_chunk) * n_chunk; } if (n_aligned > 0) { - if (n_aligned % 64 == 0) nc = 64; + if (n_aligned % n_chunk == 0) nc = n_chunk; else if (n_aligned == n) nc = n; else if (n_aligned % 32 == 0) nc = 32; else if (n_aligned % 24 == 0) nc = 24; @@ -2354,7 +2358,6 @@ class tinyBLAS_Q0_PPC { } else { mnpack(0, m, 0, n); } - #endif } private: @@ -3195,16 +3198,19 @@ class tinyBLAS_PPC { } void matmul(int64_t m, int64_t n) { + int64_t mc = 256; + int64_t nc = 256; + int64_t kc = 256; #if defined(_AIX) || defined(__BIG_ENDIAN__) - mnpack(0, m, 0, n); - #else - int64_t mc = 256; int64_t nc = 256; int64_t kc = 256; + mc = 128; + nc = 128; + kc = 128; + #endif if (m % mc == 0 && n % nc == 0 && k % kc == 0) { matmul_tiled(m, n, mc, nc, kc); } else { mnpack(0, m, 0, n); } - #endif } private: From 20a04b22063020cd0f29b7781f5352d7a6abf786 Mon Sep 17 00:00:00 2001 From: ragz4125 <65285549+ragz4125@users.noreply.github.com> Date: Mon, 6 Jul 2026 16:36:40 +0530 Subject: [PATCH 13/14] ggml-cpu: use UE4M3 LUT in ARM NVFP4 dot product (#25331) --- ggml/src/ggml-cpu/arch/arm/quants.c | 8 ++++---- ggml/src/ggml-cpu/simd-mappings.h | 4 ++-- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-cpu/arch/arm/quants.c b/ggml/src/ggml-cpu/arch/arm/quants.c index fe6213329..445366797 100644 --- a/ggml/src/ggml-cpu/arch/arm/quants.c +++ b/ggml/src/ggml-cpu/arch/arm/quants.c @@ -812,10 +812,10 @@ void ggml_vec_dot_nvfp4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const vo const float dy0 = GGML_CPU_FP16_TO_FP32(y[2*ib].d); const float dy1 = GGML_CPU_FP16_TO_FP32(y[2*ib+1].d); const float32x4_t nvsc = { - ggml_ue4m3_to_fp32(x[ib].d[0]), - ggml_ue4m3_to_fp32(x[ib].d[1]), - ggml_ue4m3_to_fp32(x[ib].d[2]), - ggml_ue4m3_to_fp32(x[ib].d[3]) + GGML_CPU_UE4M3_TO_FP32(x[ib].d[0]), + GGML_CPU_UE4M3_TO_FP32(x[ib].d[1]), + GGML_CPU_UE4M3_TO_FP32(x[ib].d[2]), + GGML_CPU_UE4M3_TO_FP32(x[ib].d[3]) }; const float32x4_t scales = vmulq_f32(nvsc, (float32x4_t){dy0, dy0, dy1, dy1}); diff --git a/ggml/src/ggml-cpu/simd-mappings.h b/ggml/src/ggml-cpu/simd-mappings.h index be50c25c0..fca5119e1 100644 --- a/ggml/src/ggml-cpu/simd-mappings.h +++ b/ggml/src/ggml-cpu/simd-mappings.h @@ -131,8 +131,8 @@ extern float ggml_table_f32_ue4m3[1 << 8]; #define GGML_CPU_E8M0_TO_FP32_HALF(x) GGML_E8M0_TO_FP32_HALF(x) #endif -// Use lookup table for UE4M3 on x86 (faster than bit manipulation) -#if defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) +// Use lookup table for UE4M3 on x86 and ARM (faster than bit manipulation) +#if defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) || defined(__ARM_NEON) #define GGML_CPU_UE4M3_TO_FP32(x) ggml_table_f32_ue4m3[(uint8_t)(x)] #else #define GGML_CPU_UE4M3_TO_FP32(x) ggml_ue4m3_to_fp32(x) From bfdf581b8b2a3c8e999227a44dfa3e890f1038bd Mon Sep 17 00:00:00 2001 From: Xuan-Son Nguyen Date: Mon, 6 Jul 2026 16:10:04 +0200 Subject: [PATCH 14/14] server: temporary skip model downloading API test (#25355) --- tools/server/tests/unit/test_router.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tools/server/tests/unit/test_router.py b/tools/server/tests/unit/test_router.py index 94165e520..cdd6f5314 100644 --- a/tools/server/tests/unit/test_router.py +++ b/tools/server/tests/unit/test_router.py @@ -314,6 +314,7 @@ def _wait_for_sse_event(collected: list, event_type: str, model: str, timeout: i return False +@pytest.mark.skip(reason="sse_thread sometimes hangs on GH actions, to be investigated") def test_router_download_model(): """Case 1: download a model, verify SSE events and GET /models.""" global server @@ -357,6 +358,7 @@ def test_router_download_model(): assert MODEL_DOWNLOAD_ID in ids, f"{MODEL_DOWNLOAD_ID} not found in /models after download" +@pytest.mark.skip(reason="sse_thread sometimes hangs on GH actions, to be investigated") def test_router_delete_model(): """Case 2: delete the downloaded model, verify it disappears from GET /models.""" global server