Skip to content

POWER10/11 MMA acceleration for all quantized formats (silicon-validated) - #100

Open
mavin2009 wants to merge 66 commits into
PrismML-Eng:prismfrom
mavin2009:power-mma
Open

POWER10/11 MMA acceleration for all quantized formats (silicon-validated)#100
mavin2009 wants to merge 66 commits into
PrismML-Eng:prismfrom
mavin2009:power-mma

Conversation

@mavin2009

Copy link
Copy Markdown

What this is

MMA (Matrix-Multiply Assist) acceleration for every quantized format this fork ships on POWER10/POWER11 — the Bonsai formats (Q1_0/Q2_0) plus the K-quants, the IQ grid codebooks, the legacy types, TQ1_0/TQ2_0, IQ4_NL/XS and MXFP4/NVFP4. Today only Q4_0/Q8_0 and the float types ride MMA; everything else falls back to scalar/VSX vec_dot.

Integration is via llamafile_sgemm behind #if defined(__MMA__): zero impact on any other architecture. The series adds ten kernel translation units, a tensor-keyed weight-pack cache (packs once per tensor per model lifetime, parallel first-touch across the op's threads), and small-n dispatch policy that routes each format to whichever path measured faster on silicon.

The branch is 21 commits, one per patch of the numbered series it was developed as; each commit message carries its measured result.

Measured on POWER10

IBM 9105-42A LPAR (4 cores/SMT2), RHEL 9.7, GCC 11.5. Baseline = the same tree built -mcpu=power9 (MMA compiled out), same machine, minutes apart.

Prompt processing: 3.7×–48× across every format tested. Highlights (pp128 @8t): Q2_0 1.7B 8.7→416 t/s (48×), IQ3_XS 1.5B 33→331 t/s (10×), TQ1_0 13→332 t/s (26×), Q4_K 27B 5.9→21.9 t/s (3.7×).

Token generation: parity or better everywhere. Q2_0 4.6× (its fallback is scalar); GER-GEMV over the cached tiles buys TQ1_0 3.1×, IQ2_XXS +42%, IQ3_XXS/S +38%, IQ1_S +23%, IQ1_M +58%; formats whose vec_dot measured better (TQ2_0, IQ2_S/XS, K-quants, IQ4, legacy) keep vec_dot at small n by explicit dispatch policy.

How it was verified

  • End-to-end: temperature-0 greedy token-identity gates (MMA build vs no-MMA build, same tree) across 18 model probes covering every accelerated format family — all PASS. Where bit-identity is impossible by construction, a three-tier gate certifies divergences against the machine's measured cross-codegen envelope (a -mcpu=power8 control build containing none of this code).
  • Decoders: bit-exact against ggml's own dequantize_row_* — every grid/sign table compared exhaustively, ~7.6M random dequantized elements at maxrel = 0.
  • Kernels: 15 standalone suites vs exact float64 references (random data, ragged shapes, multi-slab, n=1), UBSan clean, run under qemu and natively on POWER10.

Full methodology, the validation harness, benchmark data and the complete engineering log (including the experiments that lost and stayed in-tree as documented negative space) are in the companion repo: https://github.com/mavin2009/ppc-mma-kernels — start with docs/VALIDATION-POWER10.md and docs/REVIEW.md.

Caveats, stated plainly

  • Validated on one POWER10 machine. Power11 is untested (no P11-specific paths; it should simply inherit). If you have Power silicon, scripts/validate-on-power.sh in the companion repo runs the whole protocol against a checkout of this branch and emits a paste-ready report — independent results very welcome.
  • MXFP4/NVFP4 are kernel-level + decoder-cross-check verified; they cannot be produced by requantization in this fork, so no end-to-end gate exists for them yet.
  • The pack cache holds decoded int8 weights alongside the mmap'd model (defaults to a 2 GiB cap, PPC_MMA_PACK_CACHE_MB to change, refusals print loudly). Migrating it into the repack buffer-type machinery is designed but deliberately sequenced later; the design note explains why.

Happy to split this into smaller PRs (kernels first, cache second, dispatch policy third) if that's easier to review.

pl752 and others added 30 commits June 7, 2026 22:42
* 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.
thadreber-web and others added 22 commits July 19, 2026 23:50
MMA outer-product kernels for the 1-bit and ternary Bonsai formats,
with a select-and-sum GEMV that reads raw 1/2-bit weights for token
generation. Dispatch via llamafile_sgemm behind #if __MMA__.
K-quant GEMM kernels: raw unsigned codes ride the unsigned GER
operand, offsets handled algebraically (separable rank-one update
from the activation quantizer's existing sums).
Q3_K completes the K-quants; the IQ4 family flips operand
orientation (signed codebook values on the signed side, activations
XOR-flipped) with a per-row 128*codesum correction at repack.
The Q4_1/Q5_1 min term reduces to one multiply-add because Q8_1
blocks already store the scaled activation sum.
TQ1_0, TQ2_0, IQ2_XXS/XS/S, IQ3_XXS/S, IQ1_S/M collapse into shared
32-deep and 16-deep chunk kernels plus per-format scalar decoders;
decode runs at repack time, off the hot path.
Replaces silent skips in the one-shot drivers with GGML_ABORT.
First llamafile_sgemm call for a tensor packs the whole matrix once;
subsequent calls reuse it for the model lifetime. Capacity-bounded
(PPC_MMA_PACK_CACHE_MB).
…shared B-pack

FNV-1a content fingerprint detects model reloads at the same
address; admission-without-eviction makes the cache monotonic;
column-partitioned activation packing removes per-thread duplicate
pack work on cache hits.
Two accumulator sets alternate so the next chunk's GERs issue before
the previous chunk drains. On POWER10 silicon: pp -1..-5%, tg +3-4%
pre-0015; standard remains the default.
Replaces tinyBLAS_Q0_PPC for these types; small-n shapes stay on
vec_dot pending GEMV forwarding.
The 128*W correction is exactly representable and subtracted from
the exact integer accumulator before scaling, matching ggml's
scalar rounding order. Field fix from POWER10 validation.
n too small to feed every thread by columns (worst case n=1) now
row-partitions with the cached pack. Field fix, Q4_K tg32.
The packed path reads int8-expanded weights (1.8-3x native bytes)
and loses generation to vec_dot below one column tile; measured on
POWER10 (tg 4.18 -> 35.7 t/s on a 1.5B IQ2_M after this change
plus the cache fix).
128 fixed slots lost to the ~197 tensors of a 28-layer model -- a
third of the weights re-packed every call. Slot count can no longer
bind; the byte cap is the only bound and crossing it prints what it
costs. Stats API added.
These formats' vec_dot runs at 2-3% of the memory wall on POWER10;
the packed path wins despite the int8 expansion (tg +50%/+98%).
…flag

dcbt TH=8 stream hints measured +13% pp128 on IQ4_XS silicon and
become that kernel's default (PPC_DCBT_LINES restores line touches);
neutral on K-quants. lxvp vector-pair loads measured -1..-22% and
stay behind IQGRID_LXVP as documented negative space.
All threads of the op pack disjoint row-tile slices; the last one
publishes. Cold start 4.4s -> 1.05s on a 1.5B IQ2_M, +6% steady pp
from NUMA-friendly first touch.
Built to test the stream-parallelism hypothesis at n=1; measured
27-44% slower than vec_dot -- n=1 is issue-rate-bound on POWER10.
Kernel and float64 tests remain for wider silicon to re-run.
…asurement

Each xvi8ger4pp retires 4 rows x 4 depth over tiles already decoded
in the pack cache. On POWER10: TQ1_0 3.1x over vec_dot, IQ2_XXS
+42%, IQ1_S +23%, IQ3_XXS/S +38%, IQ1_M +58%; TQ2_0/IQ2_S/IQ2_XS
keep vec_dot, which measured better for them.
@khosravipasha

Copy link
Copy Markdown
Collaborator

nice, this is pretty cool. I missed this somehow.
It might be better to upstream to main llama.cpp speically for Q1_0 everything is already in the upstream.

For Q2_0 we are still in middle a migration but that will also eventually be upstreamed.

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Adds POWER10/11 MMA acceleration for quantized CPU matrix multiplication.

Changes:

  • Adds MMA kernels for K-quants, IQ formats, legacy formats, and Q1/Q2.
  • Introduces cached weight packing and format-specific dispatch.
  • Registers the new implementation with llamafile SGEMM.

Reviewed changes

Copilot reviewed 15 out of 15 changed files in this pull request and generated 2 comments.

Show a summary per file
File Description
ggml/src/ggml-cpu/CMakeLists.txt Registers new POWER sources.
llamafile/sgemm.cpp Dispatches quantized operations to MMA kernels.
llamafile/qbit_ppc_mma.{cpp,h} Implements Q1_0/Q2_0 kernels.
llamafile/q2k_ppc_mma.cpp Implements Q2_K kernels.
llamafile/q3k_ppc_mma.cpp Implements Q3_K kernels.
llamafile/q4k_ppc_mma.cpp Implements Q4_K kernels and GEMV.
llamafile/q5k_ppc_mma.cpp Implements Q5_K kernels.
llamafile/q6k_ppc_mma.cpp Implements Q6_K kernels.
llamafile/iq4_ppc_mma.cpp Implements IQ4/MXFP4 kernels.
llamafile/iq_grid_ppc_mma.cpp Implements grid-codebook and NVFP4 kernels.
llamafile/iq_grids_ppc.h Provides IQ lookup tables.
llamafile/legacy_ppc_mma.cpp Implements legacy quantized kernels.
llamafile/kquants_ppc_mma.h Declares MMA and cache APIs.
llamafile/ppc_pack_cache.cpp Implements the packed-weight cache.
Comments suppressed due to low confidence (3)

ggml/src/ggml-cpu/llamafile/ppc_pack_cache.cpp:98

  • The cache identity omits lda, although llamafile_sgemm explicitly permits any row stride at least k. Calling the kernel for two views with the same base pointer, m, k, and format but different row strides reuses the first packed layout and produces incorrect output. Include the source stride (and immutable tensor/lifetime identity) in the key and all acquire/publish calls.
extern "C" void * ppc_apack_cache_acquire(const void * key, int64_t m, int64_t k,
                                          int variant, size_t bytes, int * fresh) {

ggml/src/ggml-cpu/llamafile/ppc_pack_cache.cpp:164

  • This invalidation function is never called by the model or backend-buffer teardown paths, so admitted packs survive model unload and remain resident until process exit. Repeatedly loading models can consume the 2 GiB process-wide cap with unreachable packs, after which later models permanently fall back to per-call packing. Tie cache entries to the owning buffer/model lifetime rather than relying on an uncalled global clear.
// explicit invalidation for embedders that unload models
extern "C" void ppc_apack_cache_clear(void) {

ggml/src/ggml-cpu/llamafile/ppc_pack_cache.cpp:219

  • An unready hit is assumed to belong to the same ggml operation. If two graph executions use the same tensor concurrently, the second execution also receives fresh=1; its threads duplicate row slices and decrement the first fill's pending, potentially publishing the buffer before every unique slice is packed. Add a per-fill operation identity, or make unrelated executions wait for readiness while only the creating operation joins the fill.
        if (hit) {
            void * b = hit->buf;
            if (!hit->ready) *fresh = 1;   // join the parallel fill
            pthread_mutex_unlock(&g_mu);
            return b;

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment on lines 83 to 98
list(APPEND GGML_CPU_SOURCES
ggml-cpu/llamafile/sgemm.cpp
ggml-cpu/llamafile/qbit_ppc_mma.cpp
ggml-cpu/llamafile/q4k_ppc_mma.cpp
ggml-cpu/llamafile/q5k_ppc_mma.cpp
ggml-cpu/llamafile/q6k_ppc_mma.cpp
ggml-cpu/llamafile/q2k_ppc_mma.cpp
ggml-cpu/llamafile/q3k_ppc_mma.cpp
ggml-cpu/llamafile/iq4_ppc_mma.cpp
ggml-cpu/llamafile/legacy_ppc_mma.cpp
ggml-cpu/llamafile/iq_grid_ppc_mma.cpp
ggml-cpu/llamafile/ppc_pack_cache.cpp
ggml-cpu/llamafile/iq_grids_ppc.h
ggml-cpu/llamafile/kquants_ppc_mma.h
ggml-cpu/llamafile/qbit_ppc_mma.h
ggml-cpu/llamafile/sgemm.h)
Comment on lines +61 to +65
const size_t approx = (size_t)m * (size_t)(k/8);
const size_t tail = approx >= 32 ? approx - 32 : 0;
uint64_t h = 1469598103934665603ull;
for (int i = 0; i < 32; i++) { h ^= b[i]; h *= 1099511628211ull; }
for (int i = 0; i < 32; i++) { h ^= b[tail + i]; h *= 1099511628211ull; }
@mavin2009

Copy link
Copy Markdown
Author

Thanks for taking a look and for the feedback.

I'll go ahead and upstream it if you think that's the best place for this.

@khosravipasha

Copy link
Copy Markdown
Collaborator

@mavin2009 yes please upstream and we can pick up the changes from there.
Main thing to note taht Q1_0 is equivalent in both places, but Q2_0 ours is group size 128 and the official one is group size 64 so make sure to test that for correctness.

You can use the ggufs ending with Q2_0_g64.gguf in our huggingface to test with the official Q2_0 llama.cpp.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

8 participants