ggml-cpu: enable Q2_0 VNNI kernel on AVX-VNNI-only CPUs - #76
Open
gondoi wants to merge 40 commits into
Open
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.
The fast path in ggml_vec_dot_q2_0_q8_0 was gated on __AVX512VNNI__ && __AVX512VL__, which is stricter than needed: the kernel body uses only 256-bit registers, so it runs unchanged on CPUs that have AVX-VNNI but no AVX512 (e.g. Intel Alder Lake / Raptor Lake). The only difference is the intrinsic name: _mm256_dpbusd_avx_epi32 instead of _mm256_dpbusd_epi32. Extend the gate with defined(__AVXVNNI__) and select the intrinsic via a local macro. Builds with -DGGML_AVX_VNNI=ON (or -march=native on supporting CPUs) now take the vectorized path. Measured on i7-12650H (Raptor/Alder Lake, no AVX512) with Ternary-Bonsai-27B Q2_g64: ~8x decode speedup vs the scalar fallback, approaching the memory-bandwidth limit. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
khosravipasha
approved these changes
Jul 30, 2026
khosravipasha
left a comment
khosravipasha
left a comment
Collaborator
There was a problem hiding this comment.
somehow missed this one.
Thanks for adding, what is the main speed change, its on certain hardware that it fallsback?
We now have both Q1_0 and Q2_0 in main llama.cpp this could be a good cnadiate to upstream there too, can send PR there and we can later pick it up. Or can merge here as well, need to test it a bit
cc @bri-prism
There was a problem hiding this comment.
Pull request overview
Extends the Q2_0 x86 dot-product fast path to AVX-VNNI-only CPUs.
Changes:
- Broadens the VNNI compile-time gate.
- Selects the appropriate AVX-VNNI or AVX512-VNNI intrinsic.
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
bri-prism
requested changes
Aug 12, 2026
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Thanks for this, and the kernel itself checks out. I verified the dpbusd operand order and saturation bounds, the 2-bit extraction across all byte values, and that the SIMD path is bit-identical to the scalar fallback. Two things need fixing before this can merge, plus some suggestions inline.
- The widened gate loses the AVX2 guarantee the old AVX512 gate carried implicitly. On MSVC we define
__AVXVNNI__outside the /arch chain (ggml-cpu/CMakeLists.txt:297), so an AVX-VNNI-only config either fails to link (hsum_i32_8lives inside the__AVX__/__AVX2__/__AVX512F__region) or, with-DGGML_AVX=ON -DGGML_AVX2=OFF, builds cleanly while emitting AVX2 encodings into an /arch:AVX binary. GCC and Clang are immune since -mavxvnni implies AVX2. Suggest:
#if ((defined(__AVX512VNNI__) && defined(__AVX512VL__)) || defined(__AVXVNNI__)) && (defined(__AVX2__) || defined(__AVX512F__))
- The commit message carries a
Co-Authored-By: Claudetrailer and the PR body a "Generated with Claude Code" footer. AGENTS.md only sanctionsAssisted-by:and lists both of these forms under prohibited usage. Please amend the commit and edit the body. Worth getting right here since this was flagged as an upstreaming candidate and ggml-org enforces this harder than we do.
One correction to the PR body: -DGGML_AVX_VNNI=ON only takes effect with -DGGML_NATIVE=OFF (the flag block sits in the else of the GGML_NATIVE branch), so the enable instructions as written are a silent no-op on most builds. Also the i7-12650H is Alder Lake (not Raptor), the artifact name says g64 where our Q2_0 is g128, and 0.54 tok/s x 8 threads over a 7.6 GB model works out to roughly 33 GB/s, which is about 43% of that part's DDR5 ceiling, so I would drop the "approaching the memory-bandwidth limit" line. The gap is good news: there is still kernel-side headroom here.
Two smaller things on lines GitHub will not let me comment on directly:
- The byte replication (movq load, two 128-bit vpshufb, vinserti128) can be one vpbroadcastq plus a single 256-bit vpshufb with a combined [idxlo|idxhi] index vector. Both lanes hold the same 8 source bytes and vpshufb is in-lane, so the result is identical. Saves 2 uops per sub-block, more on E-cores.
- Nit: the
_mm256_set_m128icall should beMM256_SET_M128I(hi, lo); the file defines the macro for a reason and this is the only raw call among 42 sites.
Not blocking, for follow-up: the per-row dpbusd(ones, qy) recompute is about 20% of the inner-loop work and would vanish if Q2_0 moved to Q8_1 activations (block_q8_1.s), but that is a cross-arch traits change for another PR. And no CI currently executes an AVX-VNNI-only config (the SDE step is commented out), so a SIMD-vs-generic bit-exactness test would be cheap insurance.
| #if (defined(__AVX512VNNI__) && defined(__AVX512VL__)) || defined(__AVXVNNI__) | ||
| // VNNI: unpack 2-bit codes c in {0,1,2,3} (value = c-1), then | ||
| // dot((c-1), qy) = dpbusd(c, qy) - dpbusd(1, qy). | ||
| // The kernel only uses 256-bit registers, so it runs unchanged on |
There was a problem hiding this comment.
This comment is not quite right: the block also uses XMM ops (the idx vectors and the movq load below), and VNNI plus 256-bit registers is not a sufficient condition, which is exactly how the gate above ended up too wide. Suggest: "uses only SSE/AVX2 ops (no 512-bit or AVX512-only instructions); requires AVX2". I would also drop the CPU model list, it is already stale (Meteor Lake, Sierra Forest).
| // The kernel only uses 256-bit registers, so it runs unchanged on | ||
| // AVX-VNNI-only CPUs (e.g. Intel Alder/Raptor Lake, where AVX512 is | ||
| // unavailable); the AVX-VNNI intrinsic differs only in name. | ||
| #if defined(__AVX512VNNI__) && defined(__AVX512VL__) |
There was a problem hiding this comment.
Consider dropping the macro and calling mul_sum_us8_pairs_float (defined ~470 lines up in this file) with the whole block gated on __AVX2__ instead. It is bit-exact for this input range: codes are masked to 0..3, so its maddubs fallback cannot saturate, and the sums stay well inside exact-float territory. That removes the fifth copy of this exact dpbusd dispatch in the tree and extends the speedup to every AVX2 CPU without VNNI (Haswell through Rocket Lake, Zen 1 to 3), including the shipped haswell/skylakex variants, which this PR currently leaves on the scalar loop.
| // unavailable); the AVX-VNNI intrinsic differs only in name. | ||
| #if defined(__AVX512VNNI__) && defined(__AVX512VL__) | ||
| #define GGML_Q2_0_DPBUSD(acc, a, b) _mm256_dpbusd_epi32(acc, a, b) | ||
| #else |
There was a problem hiding this comment.
If the macro stays, make this #elif defined(__AVXVNNI__) with a trailing #else + #error, matching the other dpbusd dispatches in quants.c, repack.cpp, and sgemm.cpp. The bare #else silently emits the AVX-VNNI intrinsic for any future config that reaches it without the feature, and __AVX512VNNI__ without __AVX512VL__ is already constructible today with -DGGML_AVX512_VNNI=ON without -DGGML_AVX512=ON.
| __m256i codes = _mm256_permute4x64_epi64(_mm256_packus_epi16(r0, r1), 0xD8); // 32 codes in order | ||
| const int dp = hsum_i32_8(_mm256_dpbusd_epi32(_mm256_setzero_si256(), codes, qy)); | ||
| const int sy = hsum_i32_8(_mm256_dpbusd_epi32(_mm256_setzero_si256(), ones, qy)); | ||
| const int dp = hsum_i32_8(GGML_Q2_0_DPBUSD(_mm256_setzero_si256(), codes, qy)); |
There was a problem hiding this comment.
The two reductions can be one: keep the dpbusd results as vectors and do hsum_i32_8(_mm256_sub_epi32(dpv, syv)). Bit-identical (lanes are bounded well inside int32) and it halves the reduction work, which is latency-serialized through vmovd. The q1_0 sibling below goes further with a float accumulator and a single hsum per call if you want to match that idiom.
| } | ||
| #undef GGML_Q2_0_DPBUSD | ||
| #else | ||
| for (int i = 0; i < nb; i++) { |
There was a problem hiding this comment.
Pre-existing, but this PR rewrites the exact #if/#else around it: this scalar branch is byte-for-byte identical to ggml_vec_dot_q2_0_q8_0_generic, and the ARM version of this function already delegates instead (arch/arm/quants.c). ggml_vec_dot_q2_0_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc); return; deletes the duplicate and leaves one copy to maintain.
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.
Problem
The Q2_0 ternary fast path in
ggml_vec_dot_q2_0_q8_0(ggml/src/ggml-cpu/arch/x86/quants.c) is gated on#if defined(__AVX512VNNI__) && defined(__AVX512VL__). That gate is stricter than the code requires: CPUs with AVX-VNNI but no AVX512 — notably Intel Alder Lake and Raptor Lake client parts, where AVX512 is fused off — fall through to the scalar loop and leave most of the available throughput on the table.Change
Extend the gate to
(defined(__AVX512VNNI__) && defined(__AVX512VL__)) || defined(__AVXVNNI__)and select the dot-product intrinsic via a local macro:_mm256_dpbusd_epi32(unchanged)_mm256_dpbusd_avx_epi32Why this is safe
__m256i) and 128-bit registers — no 512-bit state, no masking, nothing AVX512-specific.vpdpbusdoperation on ymm registers; the intrinsic name (and VEX vs EVEX encoding) is the only difference. Numeric behavior is bit-identical.Measured results
On an i7-12650H (Alder Lake, AVX-VNNI, no AVX512) running Ternary-Bonsai-27B Q2_g64 (7.6 GB): stock scalar path decodes at 0.54 tok/s; with this patch the vectorized path is enabled, giving ~8x decode speedup, approaching the memory-bandwidth limit.
Enabling it
The build system already supports this:
-DGGML_AVX_VNNI=ON(which defines__AVXVNNI__, seeggml/src/ggml-cpu/CMakeLists.txt) or the default-march=nativeon a supporting CPU enables the path automatically.Compile checks
Both configurations build cleanly (gcc 15.2,
-DGGML_NATIVE=OFF -DGGML_AVX=ON -DGGML_AVX2=ON -DGGML_FMA=ON -DGGML_F16C=ON, with and without-DGGML_AVX_VNNI=ON). Disassembly confirmsvpdpbusdis emitted inggml_vec_dot_q2_0_q8_0for the VNNI build and absent (scalar fallback) in the default build.🤖 Generated with Claude Code