Public/koboldcpp
1
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-09-19 01:05:09 +02:00
Code Issues Packages Projects Releases Wiki Activity

Merge commit 'bfdf581b8b2a3c8e999227a44dfa3e890f1038bd' into concedo_experimental

Browse Source 
# Conflicts:
#	ggml/src/ggml-cuda/conv-transpose-1d.cu
#	ggml/src/ggml-hip/CMakeLists.txt
#	scripts/ui-assets.cmake
This commit is contained in:
Concedo
2026-07-07 20:53:03 +08:00
parent 838f14d652 bfdf581b8b
commit 0c56ceb613
15 changed files with 207 additions and 73 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-backend-meta.cpp
+7 -4
View File
@@ -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);
ggml/src/ggml-cpu/arch/arm/quants.c
+4 -4
View File
@@ -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});
ggml/src/ggml-cpu/llamafile/sgemm.cpp
+20 -14
View File
@@ -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:
ggml/src/ggml-cpu/simd-mappings.h
+2 -2
View File
@@ -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)
ggml/src/ggml-cuda/concat.cu
+32 -20
View File
@@ -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<T, dim><<<grid_dim, CUDA_CONCAT_BLOCK_SIZE, 0, stream>>>(
@@ -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<uint8_t>(src0, src1, dst, dim, stream);
break;
case 2:
concat_cuda<uint16_t>(src0, src1, dst, dim, stream);
break;
case 4:
concat_cuda<uint32_t>(src0, src1, dst, dim, stream);
break;
case 8:
concat_cuda<uint64_t>(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<uint8_t>(src0, src1, dst, dim, stream);
} else {
GGML_ASSERT(ggml_blck_size(src0->type) == 1);
switch (ggml_type_size(src0->type)) {
case 1:
concat_cuda<uint8_t>(src0, src1, dst, dim, stream);
break;
case 2:
concat_cuda<uint16_t>(src0, src1, dst, dim, stream);
break;
case 4:
concat_cuda<uint32_t>(src0, src1, dst, dim, stream);
break;
case 8:
concat_cuda<uint64_t>(src0, src1, dst, dim, stream);
break;
default:
GGML_ABORT("Unsupported type size: %zu", ggml_type_size(src0->type));
break;
}
}
}
ggml/src/ggml-cuda/conv-transpose-1d.cu
+16 -15
View File
@@ -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(
ggml/src/ggml-cuda/ggml-cuda.cu
+54 -12
View File
@@ -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<CUmemAccessDesc> 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:
{
ggml/src/ggml-vulkan/ggml-vulkan.cpp
+1 -1
View File
@@ -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
src/llama-model.cpp
+8
View File
@@ -1150,9 +1150,17 @@ struct llama_model::impl {
std::vector<layer_dev> dev_layer;
bool has_tensor_overrides;
std::vector<float> tensor_split_owned;
};
llama_model::llama_model(const llama_model_params & params) : params(params), pimpl(std::make_unique<impl>()) {
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;
}
tools/server/tests/unit/test_router.py
+2
View File
@@ -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
tools/ui/src/lib/components/app/navigation/SidebarNavigation/SidebarNavigation.svelte
+4 -1
View File
@@ -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<string | null>(null);
tools/ui/src/lib/enums/keyboard.enums.ts
+1
View File
@@ -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',
tools/ui/src/lib/hooks/use-keyboard-shortcuts.svelte.ts
+6
View File
@@ -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 &&
tools/ui/src/lib/services/mcp.service.ts
+24
View File
@@ -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,
tools/ui/tests/unit/mcp-service.test.ts
+26
View File
@@ -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('{}', {