diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index 0a36f0990..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)); @@ -1245,9 +1250,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 +1364,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); 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/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: 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) 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/conv-transpose-1d.cu b/ggml/src/ggml-cuda/conv-transpose-1d.cu index 7dad163a5..ebf2aa804 100644 --- a/ggml/src/ggml-cuda/conv-transpose-1d.cu +++ b/ggml/src/ggml-cuda/conv-transpose-1d.cu @@ -11,31 +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 k = 0; k < src0_ne0; k++) { + int input_numer = out_t + p0 - k*d0; + if (input_numer < 0 || input_numer % s0 != 0) { + continue; + } - int i_min = (idx >= src0_ne0) ? ((idx - src0_ne0 + s0) / s0) : 0; - int i_max_val = idx / s0; - int i_max = (i_max_val < src1_ne0) ? i_max_val : (src1_ne0 - 1); + int input_t = input_numer / s0; + if (input_t >= src1_ne0) { + continue; + } - for (int i = i_min; i <= i_max; i++) { - int weight_idx = idx - i*s0; - - float kernel_weight = src0[kernel_offset + weight_idx]; - float input_value = src1[input_offset+i]; - - 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( diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index c185ad5da..c47aa157d 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; @@ -5408,12 +5438,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: { diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 33f67732f..4e4d2bc53 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -135,7 +135,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 diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 74f909f96..ae7f48c1a 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -1150,9 +1150,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; } 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 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 && 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('{}', {