tools/ui: Respect configured MCP transport instead of always using Streamable HTTP - #97
Open
joydolma wants to merge 48 commits into
Open
tools/ui: Respect configured MCP transport instead of always using Streamable HTTP#97joydolma wants to merge 48 commits into
joydolma wants to merge 48 commits into
Conversation
* Optimized arm NEON(+DOTPROD) q1 dot * Implemented arm I8MM nrc==2 for q1 dot * Applied copilot advice about feature guards for Q1 Arm LUTs
* Implemented ARM NEON DP q1 4x4 repack * Hoisted out scaling by b_d in gemm * Added 4x8 NEON I8MM repack kernels * Cleanup for q1 arm repack * Added missing aliases for arch fallback * Corrected unused var statements * Extended table guard condition to account for i8mm w/o dp build Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Q1_0/Q2_0 had no x86 vec_dot path (arch-fallback routed the generic functions to a scalar loop). Add an AVX-512-VNNI/AVX-512VL fast path guarded by __AVX512VNNI__ && __AVX512VL__, scalar fallback otherwise: - helper ggml_hsum_i32_8_vnni to reduce _mm256_dpbusd_epi32 accumulators - Q1_0: build a sign mask from the bit field, blend +qy/-qy, accumulate with dpbusd(ones, sel) - Q2_0: vectorized 2-bit unpack (replicate-4 + 16-bit shift/mask + pack), then dpbusd(codes, qy) - dpbusd(ones, qy) = sum((code-1)*qy) Q2_0 prefill ~3.9x / decode ~3.0x vs scalar on EPYC 9655; Q1_0 ~parity (the +/-1 scalar loop already auto-vectorizes). Bit-exact vs scalar (test-quantize-fns + standalone unit test); KL-divergence vs FP16 unchanged between scalar and VNNI builds. Co-authored-by: Brian <brian@Brians-MacBook-Pro.local>
The Q1_0/Q2_0 VNNI work landed in the generic quants.c, but on x86 the generic Q1_0 path is dead (arch/x86 has an AVX2 ggml_vec_dot_q1_0_q8_0 that wins), and Q2_0 only reached x86 via an arch-fallback alias. - arch/x86/quants.c: add ggml_vec_dot_q2_0_q8_0 (AVX-512-VNNI + scalar fallback), reusing the existing hsum_i32_8 helper. Math is unchanged. - arch-fallback.h: drop the x86 ggml_vec_dot_q2_0_q8_0 alias so x86 uses the arch impl, mirroring how q1_0 is already wired. - quants.c: restore portable scalar for q1_0/q2_0 generic (the VNNI in the generic file was x86-only and is now in arch/x86).
Adds an experimental tensor-core path for Q1_0 mul_mat at batch >= 128: activations quantized to int8 with per-128 absmax scales, weights repacked once per tensor to dense sign-bit words, dequant-in-SMEM via branchless SIMD unpack feeding 64x64x32 int8 wgmma, exact per-block fp32 scaling on the accumulator drain. Hybrid dispatch: persistent stream-K grid for starved shapes, fixed tile grid otherwise. Opt-in at build time (-DGGML_CUDA_HOPPER_Q1=ON -DGGML_CUDA_CUTLASS_DIR=...) and at runtime (env GGML_HOPPER_Q1); falls through to stock MMQ otherwise. Measured on H100 SXM (1-bit test model): pp512 +8.3%, pp2048 +8.6% vs stock MMQ; test-backend-ops MUL_MAT q1_0 43/43; logit-KLD vs stock path 0.0048 mean (noise-level).
Kernels templated on weight width (1- or 2-bit dense fields); Q2_0 adds a per-tensor dense repack of the 2-bit (q-1) fields and a branchless SIMD unpack (per-byte q - 1 via __vsub4, all four field values handled). Same gating, dispatch, and activation-quant path as Q1_0. Measured on H100 SXM (ternary test model): pp512 +7.7%, pp2048 +8.1% vs stock MMQ; test-backend-ops MUL_MAT q1_0+q2_0 86/86; logit-KLD vs stock path 0.0013 mean (noise-level).
- repack cache: key on (device, wdata, N, K, wbits) with a mutex (ggml may dispatch from one host thread per device), and publish entries only after a one-time stream sync so consumers on other streams cannot observe uninitialized dense buffers - per-device attr_set / SM count (a second GPU previously skipped the dynamic-SMEM opt-in and inherited device 0's SM count) - CUDA_CHECK on allocations and attribute calls - defensive int8 clamp in the activation quantizer (unreachable in exact arithmetic; insurance against fp rounding at the boundary) - CMake: target-scoped compile definitions/includes/options instead of directory-wide; option help text covers Q2_0
The path's kernels are sm_90a-only fatbins (wgmma does not exist on Blackwell); cc >= 900 alone admits cc 1200+ where the launch fails. Blackwell support is a separate tcgen05 path.
… version (PrismML-Eng#50) The windows-cpu pack step hardcoded VC\Redist\MSVC\14.44.35112, which broke when the windows-2025 runner image moved to a newer MSVC (arm64 job failed, fail-fast cancelled x64). Glob the VS product and redist version directories and pick the newest match so runner-image updates stop breaking the release. Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
The iOS job configures with LLAMA_BUILD_TOOLS=OFF, but LLAMA_BUILD_APP defaults to ON, so the llama-app target builds without the tools include paths and fails on '#include "build-info.h"'. Upstream's build-apple.yml passes -DLLAMA_BUILD_APP=OFF for the same reason. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ML-Eng#51) * kv-cache: add optional per-channel K-cache mean-centering (Q4_0 only) GGML_TYPE_Q4_0 is a symmetric quantizer, so a K channel with a real, consistent nonzero mean across tokens wastes dynamic range encoding that constant bias. This adds an opt-in mechanism that subtracts a fixed per-(kv-head, channel) bias from K right before it is written into the cache in llama_kv_cache::cpy_k(), gated strictly to k->type == GGML_TYPE_Q4_0. Subtracting the same bias from every cached key is exactly softmax-invariant: it adds the same constant (q . k_bar) to every logit in a query's row, which softmax does not see. Nothing else in attention needs to change, so this is a zero decode-time-cost quantization-fidelity improvement. llama_context_params gets a new path_kv_mean_center field (default NULL); llama_init_from_model() hard-rejects it when the K cache type isn't Q4_0, matching the existing convention for other cache-type-gated mismatches (e.g. "V cache quantization requires flash_attn"). llama_kv_cache::load_kv_mean_center() loads the bias tensors from a GGUF file and applies them; a require_q4_0 escape hatch (used only by tests) exists to validate the underlying math against an unquantized cache without confounding it with real quantization error. Also tags the K tensor right before cpy_k() with the existing cb() graph-build hook ("k_cache_in"), so calibration tooling can capture exactly the tensor that gets written into the cache regardless of what RoPE/rotation preprocessing a given architecture applies upstream. * common: add --kv-mean-center flag and GGUF bias-file writer Adds the CLI-facing side of K-cache mean-centering: common_params gains kv_mean_center_path (plumbed into llama_context_params via common_context_params_to_llama), and --kv-mean-center takes a path to a bias file generated by tools/kv-mean-center. The bias file format is a small GGUF file with one F32 tensor per layer, named "kv_bar.blk.<il>.k", holding n_embd_head_k(il) * n_head_kv(il) values laid out as [n_embd_head_k, n_head_kv]. The writer lives in common/kv-mean-center.* so it can be shared between the calibration tool and the test suite (which needs to synthesize a bias file to check the underlying math). * tools: add llama-kv-mean-center calibration tool New tool, following the tools/imatrix convention: loads a model, runs a plain text calibration corpus through it in chunks, and captures the "k_cache_in" tensor tagged in llm_graph_context::build_attn() via the same backend-scheduler eval-callback mechanism llama-imatrix uses to capture activations (params.cb_eval). The per-(head,channel) mean across all calibration tokens is written out as a bias file consumable by --kv-mean-center. * tests: add K-cache mean-centering regression + invariance tests Uses the same tiny-synthetic-model machinery as test-llama-archs.cpp (llama_model_saver + llama_model_init_from_user with a deterministic random tensor initializer), trimmed to plain LLM_ARCH_LLAMA, to check: - regression safety: two independent contexts with centering disabled produce bit-for-bit identical logits, and a Q4_0 K cache with no bias file loaded still decodes normally (cpy_k()'s new code path is a true no-op when k_bar is empty). - the --kv-mean-center gate: a non-Q4_0 K cache with a bias file is hard-rejected by llama_init_from_model(), while Q4_0 succeeds end-to-end through a real decode. - the softmax-invariance argument itself, against an unquantized F32 K cache with a synthetic nonzero bias applied through the exact same cpy_k() code path (bypassing the Q4_0 gate via load_kv_mean_center()'s require_q4_0=false test seam): output logits match the uncentered baseline to fp32 rounding (nmse ~5e-10 in practice), confirming the math without confounding it with real quantization error. * docs: document K-cache mean-centering Explains the technique, the softmax-invariance argument, usage, the bias file format, and the current scope/limitations (Q4_0-only, plain KV cache only, standard dense/GQA attention path only). * kv-mean-center: address review feedback - kv-mean-center.cpp: the K tensor captured at "k_cache_in" can be F16 or BF16 depending on backend/compute settings, not just F32. The collector was asserting F32-only and reinterpreting raw bytes as float*, which either aborts or silently computes a garbage mean on non-F32 backends. Now accepts F32/F16/BF16 and converts to F32 via ggml_fp16_to_fp32_row/ggml_bf16_to_fp32_row before accumulating. - common/arg.cpp: wire --chunks into the LLAMA_EXAMPLE_KV_MEAN_CENTER example set so the flag the tool's own README documents is actually available, instead of being silently filtered out by the shared arg parser. - docs/kv-mean-center.md: replace non-ASCII "~=" and "." characters (was U+2248 and U+00B7) with ASCII equivalents, matching the project's ASCII-only docs convention.
Generates a calibration corpus from the model itself via a temporary llama-server, removing the need for an external calibration text file. Includes a degenerate-output guard based on gzip compression ratio. Validated end to end: the resulting bias agrees with one calibrated on a standard multi-domain calibration set to within sampling noise.
Print only the header comment as help text instead of grepping every comment line, exit nonzero on unknown options, add gzip to the dependency preflight and the Requires line, and stop forcing -ngl 99: the server's own --n-gpu-layers default now applies unless -g is given.
…r-selfgen-corpus kv-mean-center: add make-calib-corpus.sh self-generated corpus helper
…composes with K rotation) (PrismML-Eng#53) * kv-cache: support mean-centering on hybrid-memory models Hybrid (recurrent + attention) models keep a standard llama_kv_cache for their attention sublayers, but llama_init_from_model only accepted path_kv_mean_center when the whole memory module was that cache, so any hybrid model failed to load a bias file. Route the load through get_mem_attn() for hybrid memory; bias tensors are matched by model layer id and layers absent from the attention cache are skipped, which the loader already handles. * kv-mean-center: document that centering must not be combined with the K-cache rotation Measured end to end, the pre-rotation bias applied post-rotation is worse than either feature alone; strengthen the README note into a warning with the measured numbers. * kv-mean-center: record the calibration basis and reject a rotation mismatch at load The bias lives in the basis the calibration run's K cache used: the collector taps the exact tensor cpy_k() writes, after any Hadamard rotation, and the rotation is gated on the K cache being quantized. A calibration run with the default F16 cache therefore measures the unrotated basis, and applying that bias to a rotated cache measurably degrades quality instead of improving it (KLD vs F16 cache 0.00144 rotation alone vs 0.0020 with the mismatched combination), while a bias calibrated with -ctk q4_0 composes (0.00111, the best of all measured configurations). The tool now detects whether the rotation was active from the captured tensor's ancestry, records it in the output file as kv_mean_center.k_rot, and load_kv_mean_center() rejects a bias whose basis does not match the inference-time rotation state (files predating the flag load with a warning). Docs updated with the calibrate-with-matching-settings rule and the measured numbers. * kv-mean-center: address review feedback Cover SWA memory layouts: the bias load now also routes through llama_kv_cache_iswa and llama_memory_hybrid_iswa (base and SWA sub-caches), so hybrid models with sliding-window attention are no longer rejected. Validate the kv_mean_center.k_rot metadata type before reading it, so a malformed bias file produces a loader error instead of an assertion abort. Run the Q4_0 cache-type validation before the basis check, so an unsupported cache type keeps its actionable error even when the file's basis would also mismatch, and hoist it out of the tensor loop. The gate test now uses a basis-matching file for the F16-cache case so it exercises the cache-type gate specifically.
The Q2_0 2-bit quant type had no entry in test-quantize-fns's per-type error-threshold tables, so it fell through to the default thresholds (MAX_QUANTIZATION_TOTAL_ERROR = 0.002, MAX_DOT_PRODUCT_ERROR = 0.02) that are tuned for 4-bit-and-up formats. A 2-bit format cannot meet those, so the test failed deterministically on every platform (absolute error 0.008678 > 0.002, dot-product error 0.141111 > 0.02). Q2_0 stores one fp16 scale per 128-element block with no zero-point, which puts its error in the same band as the ternary formats (tq1_0/tq2_0 measure 0.008681 / 0.141345 and pass at 0.01 / 0.15). Add matching Q2_0 thresholds (0.01 absolute, 0.15 dot product) rather than loosening the shared 2-bit k-quant constant, so the Q2_K / IQ2_S checks are unaffected.
…fns-q2_0-thresholds tests: add Q2_0 error thresholds to test-quantize-fns
…1_0/Q2_0 dp4a fix (PrismML-Eng#55) * speculative: dspark block-diffusion drafter + CUDA Markov resample; Q1_0/Q2_0 dp4a fix Adds the dspark speculative-decoding drafter and two low-bit/decode improvements. dspark drafter (common/speculative.cpp, src/models/dspark.cpp): - EAGLE-style block-diffusion drafter that reuses a multi-layer target-hidden-state tap (reusable capture path, also useful for EAGLE3-proper) and drafts a block of tokens per round. - Per-round sequential Markov resample: step_logits[k] = base_logits[k] + markov_w2(markov_w1(prev_token)), argmax, chaining the sampled token forward (never batched over the block). Host scalar path by default, optional host BLAS path (LLAMA_DSPARK_MARKOV_BLAS). - GGUF arch scaffolding, converter stub, and forward-graph/loop tests. See docs/dspark-scope.md for scope and the gating rationale. CUDA device-side Markov resample (common/dspark-markov.cu/.h): - Moves the sequential per-position resample onto the GPU: one H2D of the round's base logits, then a fused GEMV + add-base + argmax kernel per position that chains through a device-resident prev token, plus a final reduction. Self-contained (CUDA runtime only). Token-identical to the host scalar/BLAS path. Default when built with CUDA; opt out with LLAMA_DSPARK_MARKOV_CUDA=0. cuda: defer Q1_0/Q2_0 dp4a symbol correction (ggml/src/ggml-cuda/vecdotq.cuh): - vec_dot_q1_0_q8_1 / vec_dot_q2_0_q8_1 (the mul_mat_vec_q decode path) built a signed symbol per element before dp4a. Both now dp4a on the raw unsigned code/bit and apply one deferred affine correction at the end using Q8_1's stored real-valued activation sum (ds.y), matching the pattern vec_dot_q4_0_q8_1_impl already uses: Q1_0 dot = d*(2*sumi*ds.x - ds.y), Q2_0 dot = d*(sumi*ds.x - ds.y). Correctness: test-backend-ops MUL_MAT passes for both types. * cuda: forward-declare ggml_cuda_mul_mat_q1_hopper (fix -Werror=missing-declarations) Pre-existing on prism: the Hopper Q1 entry point is defined in mmq-hopper-q1.cu but only forward-declared locally in ggml-cuda.cu, so the definition's translation unit has no prior declaration and -Werror= missing-declarations breaks the cuda build. The full build matrix only runs on PRs (not prism pushes), so this stayed latent. One-line forward declaration; no behavior change. * tests: fix dspark test cross-platform builds (macos -Werror, windows DLL link) Surfaced by the full CI matrix (only runs on PRs, not prism pushes): - macos clang -Werror,-Wmissing-noreturn: the test-local fail() helpers never return; mark them [[noreturn]] (test-dspark-forward/loop/real-eval). - x64-windows-llvm link error: test-dspark-forward used the llama_model::get_tensor MEMBER function, which is not reliably exported across the Windows DLL boundary. Switch to the exported free function llama_internal_get_tensor_map (same pattern test-quantize-stats already uses cross-platform). No behavior change. Verified locally: all three targets build clean under clang (CPU build). * tests: use public vocab API for dspark n_vocab (windows DLL link) The prior fix swapped llama_model::get_tensor for llama_internal_get_tensor_map, but that free function is also not LLAMA_API-exported (its only other user, test-quantize-stats, is gated NOT WIN32, so it was never windows-linked) and still fails to resolve across the Windows DLL boundary. Use the exported public API instead: llama_vocab_n_tokens(llama_model_get_vocab(model)). The dspark converter's set_vocab() fills the 'none' tokenizer with dummy entries sized to the target's real vocab width, so n_tokens() reports the correct value (the old code comment claiming it is 0 predates that converter behavior). Verified locally: all three dspark test targets build clean under clang. * dspark: widen capture-copy size math to size_t (CodeQL overflow) CodeQL flagged the two dspark capture-copy sites in llama-context.cpp: 'row' was uint32_t, so n_tokens*row and n_outputs*row (feeding the byte size n_*row*sizeof(float)) were evaluated in 32-bit before widening to size_t. No overflow at real capture configs (row=n_capture_layers*n_embd is small), but latent for large captures. Make 'row' size_t so all downstream size math is 64-bit. No behavior change. libllama builds clean. * llama-context: bound capture-layer writes and zero unpopulated capture rows llama_set_capture_layers() wrote one capture_layer_idx[] slot per accepted layer id but only range-checked the id value, so a caller repeating layer ids could advance the write index past the fixed LLAMA_MAX_LAYERS array and corrupt adjacent cparams. Stop once the array is full. On any architecture whose graph does not build a capture tensor (currently every arch except qwen35), the capture copy was skipped while output_reserve() had already allocated embd_capture, so llama_get_embeddings_capture*() returned uninitialized memory. Zero the destination rows and warn once instead. * speculative: harden dspark drafter/target contract and recovery paths Validate the drafter against the target at construction: the drafter consumes the target's hidden states (each capture row is target_hidden wide, copied verbatim) and resamples over the target's vocabulary, so a drafter trained against a different target would over-read capture rows or index the wrong vocab. Fail loudly here rather than corrupt every round. When the per-round drafter cache tail cannot be cropped, do not advance n_cache past a tail that is still physically present; reset the drafter sequence and rebuild context next round. The markov-resample mask_token_id checks aborted via GGML_ASSERT, but a sampled token can legitimately equal the mask sentinel (a real vocab id) -- that only makes a poor draft the target rejects. Warn once instead of aborting a valid run; the sequential chaining is guaranteed structurally by construction. * convert: map dspark per-layer tensors under drafter./dspark. wrappers The per-layer fallthrough passed the original tensor name to map_tensor_name, so a decoder tensor nested under a drafter. or dspark. wrapper kept that unsupported prefix and failed to map (only the model. prefix map_tensor_name understands worked). Strip the dspark-specific wrapper for the fallthrough while preserving any standard model. prefix. * common: build the dspark Markov CUDA TU for ggml's architectures, not SM80 only The resample TU pinned CUDA_ARCHITECTURES to 80 whenever CMAKE_CUDA_ARCHITECTURES was not defined in this scope -- which is the common case, since ggml resolves it inside its own subdirectory and it does not propagate up. That forces a PTX JIT on Hopper/H100 and fails to build for the pre-Ampere GPUs the rest of ggml supports. Inherit the ggml-cuda target's resolved architecture list instead, falling back to native detection. * tests: gate dspark tier-2 on agreement and actually run the rs-ring test Tier 2 only failed on non-finite logits, so it passed on any finite output even with zero argmax matches. Gate on argmax-match-rate and top-5 overlap (both scale-invariant, defaulted high, CLI-overridable), and scan logits from token 0 so a bad first logit is no longer excluded from the diff/non-finite metrics. test-rs-ring-rotation was registered against the non-recurrent stories model, so it always took the self-skip path and never exercised the ring. Generate a tiny recurrent qwen35 fixture with the pure-Python generator when numpy/gguf are importable and require it; otherwise fall back to the stories model unchanged. * docs: correct dspark Markov CUDA default and drop stale scaffolding claims The Markov CUDA path is the default at runtime when the drafter has a Markov head (opt out with LLAMA_DSPARK_MARKOV_CUDA=0), not opt-in via =1 -- fix the root CMake option comment and dspark-markov.h to match the runtime behavior. Note that the device warp tree-reduction changes floating-point accumulation order versus the host scalar/BLAS paths, so the resample is functionally equivalent rather than bit-identical: a near-tie argmax can differ, which only changes the speculative proposal the target verify still arbitrates. Drop the 'no forward graph / scaffolding only' claims in the converter docstring and logging, the GGUF tensor-registry comment, and docs/dspark-scope.md -- the forward graph and block-diffusion draft loop are implemented in this branch. * tests: run rs-ring rotation on CPU (-ngl 0) The ring bit-identity check compares logits from a ring and a no-ring context and requires both to run the identical GDN compute path. On GPU backends without a fused GDN op the op falls back per context and the two diverge (observed on Metal: 'fused Gated Delta Net not supported, set to disabled'). The invariant and this test are defined on CPU, per the test header, so pin the run to CPU. * dspark: load and run drafters with GIDD log-SNR conditioning Some drafters ship a LogSnrEmbed module -- a sinusoidal featurization of a per-position log-SNR value run through a 2-layer SiLU MLP, added to the draft noise embedding before the backbone. Without loader support these drafters fail to load: their four log_snr_fc tensors are unmapped (wrong number of tensors). Add the optional GGUF metadata (dspark.log_snr_conditioning, min/max_log_snr) and the dspark.log_snr_fc1/fc2 tensors, gated on log_snr_conditioning so drafters without it load and run exactly as before. The per-position log-SNR pattern (anchor of each block at max_log_snr, mask positions at min_log_snr) and its featurization are a pure function of n_draft/block_size/min/max_log_snr, all known at graph-build time, so the feature matrix is precomputed host-side and staged via a new llm_graph_input_dspark_logsnr input; only the learned fc1/fc2 weights go through ggml. The featurization divides by (max_log_snr - min_log_snr), so when conditioning is enabled the bounds are required and validated finite and strictly ordered at load time rather than silently producing NaN embeddings. * dspark: drop unused pos from llama_set_dspark_ctx; document draft-dspark contract The pos argument to llama_set_dspark_ctx (and the v_ctx_pos it filled) was never consumed: the drafter graph only uploads the tap features, and each context row's decode position comes from the batch. Shipping a public-API parameter that has no effect is misleading, so drop it (this API is new in this branch, no external callers) along with the dead v_ctx_pos storage. Also document, at the speculative-type registry, that draft-dspark requires the driver to engage multi-layer capture (llama_set_capture_layers plus per-row logits) before drafting: the reference driver is tests/test-dspark-real-eval.cpp, and the generic CLI/server paths do not yet engage capture, so selecting it there fails at the first draft with a clear error. * speculative: validate the target's configured capture-layer count in dspark process() The dspark row copy reads n_embd_cap = n_capture * n_embd floats from each target capture row, but capture layers are engaged by the driver after the impl is constructed, so the ctor's n_embd/n_vocab checks cannot see the configured layer count. A driver that engaged fewer layers than the drafter was trained on would over-read past the end of the capture row; more layers would feed misaligned features. Check llama_get_n_capture(ctx_tgt) against the drafter's n_capture at process() time and fail loudly on mismatch. Also correct a stale ctor comment that still described the CUDA Markov resample as token-identical to the host path; it is functionally equivalent but not bit-identical (see common/dspark-markov.h).
The C++ loader requires dspark.log_snr_fc1/fc2 tensors and the log_snr_conditioning/min_log_snr/max_log_snr KV when a drafter uses log-SNR conditioning, but the Python side never mapped them, so converting any log-SNR-conditioned drafter failed with: Can not map tensor 'log_snr_embed.fc1.bias'.
…rismML-Eng#59) * metal: fix M5 device creation + add Q2_0 multi-column mul_mv kernels Two independent Metal changes: - The AGX_RELAX_CDM_CTXSTORE_TIMEOUT override (added for the long-context command-buffer timeout on M1/M2, ggml-org#20141) prevents MTLCreateSystemDefaultDevice() from returning a device on M5 / current macOS. Keep it on by default and disable it only when sysctl reports an M5 chip; GGML_METAL_RELAX_CDM_CTXSTORE_TIMEOUT=0/1 forces either way. - Add Q1_0-style multi-column mul_mv variants for Q2_0 (nr1 2/3/4): read the streamed weights once for nr1 output columns via a 2-bit weight expansion and an FMA inner loop, instead of re-reading them per column on the mul_mv_ext path. Opt-in via GGML_METAL_Q2_0_NR1 (default routing unchanged). Measured [4096,14336] on M5 Pro: nr1_2 93.2 us at ne11=2 vs 122 for the ext route. 41/41 test-backend-ops MUL_MAT q2_0 on both routings. * speculative: Metal DSpark Markov resample + quantized markov heads Adds a Metal device path for the block Markov resample, alongside the existing CUDA path. It builds one dependency-chain graph for the whole draft block (each step's GPU argmax feeds the next step's get_rows) and submits it once, reading the drafter's still-device-resident logits, so the sequential Markov dependency stays exact with a single sync. Falls back to the host path when the head type or backend is unsupported, or when DSPARK_MARKOV_CPU=1. Also teaches llama_model_dspark_get_markov to dequantize quantized markov head tensors (Q4_0/Q5_0/Q8_0) for the host/CUDA path, instead of rejecting everything but f32/f16/bf16 -- drafters that ship quantized heads previously had the correction silently disabled (has_markov=0). Validated on Metal: accept counts byte-identical to the CPU-forced path; the Metal resample runs ~1.8x the single-thread CPU Markov path.
…rismML-Eng#58) Two independent Metal changes: - The AGX_RELAX_CDM_CTXSTORE_TIMEOUT override (added for the long-context command-buffer timeout on M1/M2, ggml-org#20141) prevents MTLCreateSystemDefaultDevice() from returning a device on M5 / current macOS. Keep it on by default and disable it only when sysctl reports an M5 chip; GGML_METAL_RELAX_CDM_CTXSTORE_TIMEOUT=0/1 forces either way. - Add Q1_0-style multi-column mul_mv variants for Q2_0 (nr1 2/3/4): read the streamed weights once for nr1 output columns via a 2-bit weight expansion and an FMA inner loop, instead of re-reading them per column on the mul_mv_ext path. Opt-in via GGML_METAL_Q2_0_NR1 (default routing unchanged). Measured [4096,14336] on M5 Pro: nr1_2 93.2 us at ne11=2 vs 122 for the ext route. 41/41 test-backend-ops MUL_MAT q2_0 on both routings.
… path) (PrismML-Eng#61) * ggml: rows-indexed state read for the fused GDN op (ring decode path) On the ring-enabled decode path every GDN layer paid two extra dispatches per token just to feed the fused op its input state: a get_rows gather of the per-seq live states into a contiguous scratch, then a cpy of that gather into slot 0 of the (D, K, n_seqs) state input. Both are pure reads of the recurrent cache -- ~786k floats each way per layer on the 27B target -- serialized into a launch-bound decode graph, 96 dispatches and ~300 MB of scratch traffic per token across 48 layers. Add ggml_gated_delta_net_rows: the op takes the 2D cache view plus the per-seq row indices (inp->s_copy_main) as src[6] and reads each sequence's live state directly at cache row rows[seq]. K moves to op_params so both variants share one backend code path. The graph side gains build_rs_cache_view (rs_zero clear + extra-states relocation, no main gather) and qwen35 wires it on the ring path, with GGML_GDN_STATE_GATHER=1 restoring the legacy gathered path for A/B. Implemented on CPU and Metal (function-constant-gated read base, no kargs change). All other backends that support GATED_DELTA_NET reject src[6] in supports_op so rows-mode ops fall back instead of silently reading src[5] as a scratch. test-backend-ops gains rows-mode cases (single/multi-token, multi-seq, snapshot overflow, KDA): 38/38 OK on MTL0, CPU leg green. Real-eval gate: accept counts bit-identical to the gathered path at n_max 1..4 (alpaca x24). Measured on M5 Pro (cont6k Q1_0 x bin6l1 q4_0, ring 4): harness AR 32.0 -> 35.3 tok/s (+10.5%), spec@n3 33.5 -> 35.3 (+5.5%); ring-free llama-bench unchanged (~42), as expected. * ggml: fold recurrent GDN snapshot writes on Metal * metal gdn: always populate snapshot tail on write-fold, handle K==1 rows Two correctness fixes to the folded rows-mode GDN epilogue: - The write-fold followed the SET_ROWS view chain to prove the scatter consumes the GDN result, but not that it is the snapshot tail's sole consumer. The kernel now always writes the op's own documented output tail AND additionally scatters into the cache row, so a second consumer or an output/eval callback never observes an uninitialized region. - WRITE_ROWS scatter existed only in the K>1 branch; a rows-mode graph with K==1 suppressed the SET_ROWS but wrote only the output tail, losing the cache update. The K==1 final-state branch now scatters to the cache row as well. Gate: test-backend-ops GATED_DELTA_NET 39/39 on MTL0; e2e accept invariant 76/116 tau 2.3103 unchanged (default / fold-disabled / gathered). * qwen35: gate GDN rows mode to Metal-only GPU device sets rows mode uses the src[6] GDN variant, implemented on CPU and Metal only; other GPU backends reject it in supports_op, which would move the recurrent op (and its state traffic) to CPU. Select rows mode only when every GPU device in the model is Metal (ACCEL/BLAS devices are skipped). * ggml: disable OpenMP for Emscripten/WASM builds The WASM CI build enables OpenMP (-DGGML_USE_OPENMP -fopenmp=libomp), but Emscripten cannot emit the common symbols libomp's reduction helpers need (.gomp_critical_user_.reduction.var), so ggml-quants.c fails to compile. WASM has no host threads to benefit from OpenMP -- force it off for the Emscripten target instead of failing the build.
…e overflow, write-fold guards, K==1 tests) (PrismML-Eng#62) * address review: CPU workspace sizing, write-fold guards, K==1 rows tests - ggml-cpu: size the GDN scratch from the op-param K (snapshot slots), not src[5]->ne[1] -- in rows mode that dim is the cache row count, so a 1-row cache with K>1 (batch-1 block decode) undersized the scratch and overflowed the work buffer. - metal write-fold: honor ctx->use_fusion (GGML_METAL_FUSION_DISABLE), and verify the SET_ROWS target is exactly the snapshot tail (per-row state width, index count, dest row width) before suppressing it -- descent from the GDN output alone let a mis-sized view be fused, reading row indices out of bounds. - rows-mode state view: document + assert the main/extra row-range disjointness invariant that makes the deferred (read-after-relocate) main read safe. - tests: add rows-mode K==1 cases to exercise the K==1 final-state branch. * review round 2: byte-offset write-fold check, honest rows-mode ordering note, 1-row-cache K>1 test - write-fold: also verify the SET_ROWS view begins at the snapshot-tail byte offset (attn_size + (K-min(T,K))*state_size_per_snap), not just matching size/counts -- a same-sized view at another offset no longer folds. - rows-mode state view: drop the incorrect disjointness assert (s_copy returns idx*size+src0, an arbitrary slot, so it did not establish disjointness). Document the real read-before-relocation hazard (multi-seq; not reachable on the single-seq decode path) as tracked follow-up. - tests: add a rows-mode 1-row-cache K>1 case that reproduces the CPU workspace under-size the planner fix prevents. * write-fold: require compact snapshot-row stride before folding ggml_set_rows only requires contiguous rows (nb[0]); it permits an arbitrary row stride nb[1] that its kernel honors, but the fused GDN epilogue scatters the contiguous snapshot tail. Require the compact [D, n_write] layout (ne[0]==D, nb[0]==type_size, nb[1]==D*type_size) so a strided view falls through to the real SET_ROWS instead of being mis-scattered.
…n every prompt row Gives DSpark's multi-layer hidden-state tap capture its own masked flag, separate from embeddings_nextn_masked which it was previously reusing. Opting into masked=false keeps capture dense (every prompt position) regardless of batch.logits, so callers can request logits=false on context rows (as the plain AR path already does) while still getting a full per-position capture buffer for the drafter. Mirrors the existing embeddings_nextn unmasked path (llama_context.cpp) at every layer: cparams flag, output_reserve sizing, per-decode readback offset/size, and get_embeddings_capture_ith row resolution. test-dspark-real-eval.cpp now engages capture with masked=false and drops the speculative-path prefill's logits back to false on context rows, matching the AR baseline. Fixes the harness-side third of the PP slowdown reported in PrismML-Eng#33: capture previously needed logits=true on every row just to populate a capture row for it, which forced the full-vocab lm_head projection to run on every prompt position instead of one.
- restore the default masked=true on llama_set_capture_layers's public declaration -- it was mandatory there, breaking source compat for any existing 3-arg caller. - set_capture_layers() no longer stomps embeddings_nextn_masked=true as a side effect; that assignment predated the independent capture flag and made the assert below unreachable in the exact case it exists to catch (dense capture silently overriding a caller's masked=false nextn config instead of tripping the guard). - narrow_before_last_layer's capture-side deferral now only applies when the LAST layer is actually one of the requested capture layers; taps at any earlier layer already branched off cur before this point in the loop, so deferring the last layer's own narrowing for them was an unnecessary regression (recovers a little more speed on today's real checkpoints, whose taps never include the last layer). - guard dense (unmasked) capture to single-sequence ubatches: its rows are indexed/reordered assuming raw-token order, which split_equal()'s per-sequence interleaving on a multi-sequence ubatch would violate. No current consumer is multi-sequence; fail loudly instead of silently returning another sequence's capture if that changes.
…d-capture dspark: give tap capture its own unmasked path, avoid full-vocab lm_head on every prompt row
…PrismML-Eng#64) Mirror the Q1_0 default to Q2_0: for ne11>=2 use the nr1=2 multi-column variant that reads each streamed q2_0 weight group once per 2 src1 columns, instead of the mul_mv_ext route. Measured (M5 Pro, in-code microbench): nr1_2 = 93.2us vs 122us ext at ne11=2 (+31%); ne11=4 via 2 passes 171 vs 183. ne11==3 is carved out -- that is the nr1_3 occupancy cliff (195 vs 152 ext, tpb=16) that kept this path opt-in; three columns stay on ext. Output is identical (pure matvec routing); base decode (ne11=1) is unaffected. GGML_METAL_Q2_0_NR1=1 restores the old routing.
* server, speculative-simple: wire dspark tap capture (draft-dspark support) * dspark server/cli: review fixes -- batch headroom, n-max validation, ctx-shift and child-slot gating
Two cpp/wrong-type-format-argument defects (high severity): - the vision_feature_layer/proj_spatial_offsets size-mismatch throw passed size_t values to %d and had no argument for its leading %s; use %zu and pass __func__. - the qwen-flamingo projector-block loop used size_t bid, passed to the %d in the TN_QF_* tensor-name formats; make bid int so the format type matches.
…rismML-Eng#65) The download tests make live HTTP requests to http://ggml.ai/ and assert on the response. Network-restricted CI runners (self-hosted, windows-vulkan) can't reach it, so the good-URL GET fails and takes the whole arg-parser suite down on an unrelated connectivity issue. Probe the endpoint once and assert the download semantics only when it is actually reachable; otherwise print a notice and skip. No behavior change when network is present.
…g#80) common_remote_get_content() returns the HTTP status as long, but the skip-message printf used %d. Format it with %ld so the build does not fail under -Werror=format on the CUDA CI toolchain.
* cuda: speed up Q1_0 extraction with byte permutes * cuda: use unsigned halfword for Q1_0 byte-perm selectors unpack_q1_0_bytes() built the __byte_perm selectors by right-shifting a signed int16_t, so a packed halfword with the top bit set sign-extends through the shift and corrupts the selector nibbles. Take the packed word as uint16_t (widened through uint32_t before the shift) at the helper and at both the MMVQ and MMQ call sites.
…fork note (PrismML-Eng#78) * ci(release): ship full self-contained Windows Vulkan/HIP bundles readme: add Prism fork note (start with Bonsai-demo, main caveats) * readme: fix Q2_0 model-file guidance (fork=Q2_0, mainline=Q2_0_g64, PQ2_0 future) + link demo status; ASCII punctuation
ggml_vec_dot_q2_0_q8_0 gated its VNNI fast path on __AVX512VNNI__ && __AVX512VL__ with no AVX2/AVX-VNNI fallback, so x86 CPUs without AVX-512 silently took the scalar loop. This excludes all Intel 12th-14th gen consumer CPUs (Alder/Raptor Lake), where AVX-512 is fused off for the P/E hybrid design but AVX-VNNI is present. The fast-path body is already entirely 256-bit AVX2; the only AVX-512 dependency is the _mm256_dpbusd_epi32 intrinsic. AVX-VNNI exposes the identical operation as _mm256_dpbusd_avx_epi32, so alias the intrinsic and widen the guard. No algorithmic change, and no behavior change on AVX-512-VNNI hosts. Measured on Intel i5-13400 (Raptor Lake, AVX-VNNI, no AVX-512), 12 threads, CPU-only, same model and prompt (temperature=0): Ternary-Bonsai-8B Q2_0 decode: 2.17 -> 6.92 tok/s (3.2x) Ternary-Bonsai-8B Q2_0 prompt eval: 2.7 -> 8.6 tok/s (3.2x) Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…-Eng#86) * ggml-cpu: x86 AVX512-VNNI repack GEMV/GEMM for Q1_0 and Q2_0 Q2_0 previously had no repack path at all and Q1_0 only a NEON one, so batched CPU mul_mat for both formats fell back to per-row vec_dot on x86. Add: - block_q2_0x4 (4-row, 8-byte-chunk interleave) with generic repack, gemv and gemm implementations - AVX512-VNNI gemv/gemm kernels for the existing q1_0 4x8 layout and the new q2_0 4x8 layout; sum(qy) is computed once per activation sub-block and the horizontal reduction happens once per output tile - repack type selection on AVX512-VNNI CPUs for both formats * ggml-cpu: fix int-overflow-before-widening in Q2_0 repack (CodeQL) Cast to size_t/int64_t before the nrow*nblocks and i*nblocks multiplications so they cannot overflow int before widening. Resolves CodeQL alerts 630/631.
joydolma
changed the title
Collaborator
|
Thanks for the PR, will take a closer look. Is this somethign related to our fork? Or should it be more in the main llama.cpp. Does not seem to be touching our part of code (our main addition is mianly the kernels for 1-bit and ternary models, and few experiments features kv-cache 4-bit and dspark.) |
There was a problem hiding this comment.
Pull request overview
Updates MCP transport selection to support legacy SSE servers while retaining Streamable HTTP defaults.
Changes:
- Detects SSE transport from endpoint URLs.
- Creates the transport specified by the MCP configuration.
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 3 comments.
| File | Description |
|---|---|
tools/ui/src/lib/utils/mcp.ts |
Adds SSE URL detection. |
tools/ui/src/lib/services/mcp.service.ts |
Selects configured SSE or Streamable HTTP transport. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Author
Thanks for taking a look! You're right—this isn't related to Prism's inference/runtime additions. It affects the MCP client in the bundled UI. While testing with Nuke MCP, ComfyUI MCP, and my own MCP Bridge, I found that the configured transport wasn't being respected. Even when SSE was explicitly configured, the client always created a Streamable HTTP transport, which prevented legacy SSE MCP servers from connecting. I've updated the implementation based on the review feedback, added regression tests, and all unit tests are now passing. If you think this is better suited for upstream llama.cpp, I'm happy to submit it there instead. |
Collaborator
|
@joydolma Yes please, upstream is best for this I think, then we can pick it up ourselves from llama.cpp. Curious any nice MCPs did you try that works well with our models? |
Collaborator
|
@joydolma is this upstreamed, if theere is a PR let us know. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
This PR fixes MCP transport selection in the UI by respecting the configured
transport type instead of always defaulting to Streamable HTTP.
This restores compatibility with legacy SSE-based MCP servers while preserving
support for newer Streamable HTTP servers.
Changes
Motivation
Some MCP servers (including existing third-party implementations) still expose
the legacy SSE transport. The current behavior always prefers Streamable HTTP,
which prevents successful connections to those servers even when SSE is
explicitly configured.
This change allows the configured transport to be used while maintaining
compatibility with both transport implementations.
Compatibility
This change has been verified with:
The goal is to preserve compatibility with both transport implementations
without changing the default behavior for Streamable HTTP servers.
Testing
Tested successfully with:
Verified tool discovery and execution with both legacy SSE and Streamable HTTP
servers.