Repository navigation
metal: q4_0 KV cache on the Mac engine (--kv q4_0, opt-in; the non-q4 path byte-identical when off) - #2533
Closed
afanty2021 wants to merge 26 commits into
Closed
metal: q4_0 KV cache on the Mac engine (--kv q4_0, opt-in; the non-q4 path byte-identical when off)#2533afanty2021 wants to merge 26 commits into
afanty2021 wants to merge 26 commits into
Conversation
…M0-p)
metal/kernels/kv_q4.metal ports src/kernels/cuda/kv_q4.cu's four __global__ kernels
bit for bit: fwht256 (one simdgroup a row, 8 values a lane, orthonormal 1/16 scale),
kv_append_q4_kernel (single step), kv_append_q4_batch_kernel (the prompt path with the
host and stage pools) and kv_gather_q4_kernel, over the q4_group/q4_store device
functions with the deterministic tie rule (a tie in |x| resolves to the larger value in
every lane). The numerics are pinned: the fp16 scale is the software
round-to-nearest-even of f16_bits.hpp's f16_from_f32 (MSL's half() is not used), and
the codeword is int(fma(x, id, 8.5f)) - the contraction the CUDA build's --fmad=true
gives - with int() truncating toward zero as __float2int_rz does.
Every kernel has two binding forms over one body: the CUDA-shaped standalone kernels
(each argument its own buffer, what the parity test drives through mtl_rt) and the *_e
twins that read one argument struct from buffer 0 - the engine dispatch of
strata/gpu/metal.hpp.
metal/tests/kv_q4_parity.cpp checks bit-level parity (whole 18-byte blocks / whole
buffers) against a CPU reference that replicates the device arithmetic exactly: the
same register butterfly, the same tie rule, the same fma. Cases: fwht256; the
single-step append with tie blocks in the rotated domain (x = {v, -v, ...}, one +v
among -v, all equal) and a permuted page table; a 3-step append; the batch kernel
with its host and stage pools; the composite raw input -> device fwht -> device
quantize chain; and the gather's fp16 scratch bits. All PASS on the Apple M4 Pro.
…enerated metal/src/engine/kv_q4_host.cpp implements fwht256_cuda, kv_append_q4_step/steps, kv_append_q4 (the prompt batch path, host and stage pools) and kv_gather_q4_step by name through the device runtime (strata/gpu/metal.hpp's pipeline + dispatch): the source of metal/kernels/kv_q4.metal is compiled at run time, the *_e kernels read their argument struct (the mirrored layouts are pinned by static_asserts) from buffer 0 on the caller's stream, and the page table is passed through unchanged - the engine hands the identity one, Mac having no paging. stubs.cpp regenerated with tools/gen_stubs.py after the port (a link of strata without it as the symbol log): the five kv_q4 stubs are gone. The regeneration also drops iq_dequant_f16 / iq_dequant_gu_f16 / prefill::moe_combine, no longer undefined at this source state (their headers' signatures moved under stubs_144.cpp's overloads), and refreshes the parsed header list - the old file predated those. Bit-level through the engine path too (a one-off probe driving the entry points over gpu::malloc_device memory): fwht256_cuda, kv_append_q4_step, kv_gather_q4_step and kv_append_q4 with the stage pool all match the CPU reference exactly.
…her table, a page-crossing batch Four findings of the final pass: - kv_q4_host.cpp now pins EVERY field's offset of the four mirrored MSL argument structs with static_assert(offsetof(...)), not just the total sizes: a same-size field reorder would previously pass the size asserts and dispatch garbage silently. The engine-path probe still matches bit-level, so the pinned offsets are the layouts the kernels read. - kv_q4.metal's KvQ4Pools struct was dead (no kernel or *_e twin takes it; the twins carry flat host_k/host_v fields) - deleted. - the parity test's 3-step case now runs a PERMUTED page table, and the gather reference computes its row through the table, so the kernel's page-table lookup indirection is actually exercised (the identity hid it). - the batch case starts at pos0 = page_size - 8 with T = 17, so one dispatch crosses a page boundary (cells 56..72). kv_q4_parity: ALL PASS, bit-level, all 17 checks (Apple M4 Pro); the engine-path probe (fwht256_cuda, kv_append_q4_step, kv_gather_q4_step, kv_append_q4 with the stage pool) all PASS through the real gpu::metal dispatch.
…rt, a fast-math note on the divisions, the kernels-dir env override on the engine path - kv_q4.metal: static_assert(sizeof(kv_q4_blk) == 18); the header's numerics list now says a division like mval / -8 is approximate under STRATA_METAL_MATH=fast (the test runtime honors the env, the engine runtime is pinned Safe), so bit-level parity holds in Safe math only; the local kv_q4_blk states its true reason (MSL cannot include C++ headers). - kv_q4_host.cpp: the kernels directory reads STRATA_METAL_KERNELS with STRATA_METAL_KERNELS_DEFAULT as the fallback, as mtl_rt's Context already did - the engine path no longer hardcodes the default.
…ed composite table, a failure locator, the reference factored into kv_q4_ref.h - fixtures: an all-zero group (d = -0.0 -> scale bits 0x8000, id = 0 -> q = 8, qs bytes 0x88 - asserted off the device pool) and the amplitude classes 1e-3 / 40 / 1 of src/kernels/kv_steps_parity.cpp:80 across the 3-step append; - case 5's composite chain goes through a permuted page table like cases 1-4; - cmp_blocks prints the first differing block index and byte offset; - the device-rule reference (fwht, q4 block, append cell) moves to a tests-local kv_q4_ref.h shared with the new engine-path test; the local block_q4_0's comment states its real reason (kv_q4_ref.h keeps one shared include; kv_q4.hpp would drag in the whole qsa.hpp contract) instead of the stale include-path claim.
…oints over strata_gpu_metal kv_q4_engine.cpp drives fwht256_cuda, kv_append_q4_step + kv_gather_q4_step (round trip), a folded planes=1 kv_append_q4_steps and the kv_append_q4 batch with host AND stage pools - the *_e twins' marshalling (one args struct at buffer 0, gpuAddress pointers), the host's grids and the planes folding, bit-level against the shared kv_q4_ref.h. Hand-wired in CMake next to runtime_hello (excluded from the tests glob): it compiles kv_q4_host.cpp in (no ggml) and links strata_gpu_metal + mtl_rt, with STRATA_METAL_KERNELS_DEFAULT pointed at the kernels directory.
…actored out The prelude section (QK4, kv_q4_blk, the software fp16 conversions, q4_group, q4_store) is now banner-marked as the single source of the q4 numerics, with fwht_step the shared butterfly unit: fwht256_lane routes its lane butterflies through it, and qsa.metal's coming q4 mode (function constant 25) will route both its lane and its cross-simdgroup butterflies through the same arithmetic, so the standalone and fused rotations cannot drift. Pure refactor - the lane loop's expression is unchanged - and kv_q4_parity plus the engine-path kv_q4_engine output identical bytes before and after (checked). No separate kv_q4_shared.metalh to #include: metal/src/engine/kv_q4_host.cpp compiles this file ALONE as a bare source string, where no include path exists, so the file stays textually self-contained and sharing goes through metal/'s concatenation convention instead (compile lists prepend this file).
QSA_KV_Q4 (fc 25, free: 20/22/23/24 in use, 7/8 in qsa_prompt's context), off
by default - the Q4_0 KV with the Walsh-Hadamard rotation of kv_q4.hpp fused
into the decode kernels, the PC semantics of layer.cpp's --kv q4_0 path:
qsa_prep the k rows (after norm/RoPE), the v rows (raw - the PC normalizes
only K) and the q attention rows (after norm/RoPE, r < QH) pass
through the orthonormal 256-point FWHT, then pack into the q4
pools [cell][kv head][144] (576 B/cell, the fp16 scale inside each
18-byte block, ksc/vsc unused). The FWHT is the new fwht_row256:
the row/thread layout is NOT restructured (one value per thread,
8 simdgroups x 32 lanes, so a 32-value block is exactly one
simdgroup and q4_group maps onto simd reductions); the low 5 index
bits butterfly within each simdgroup through the shared fwht_step
of kv_q4.metal's prelude, the high 3 across the simdgroups through
threadgroup staging - tg_sum256's part[8] pattern, one slot per
(simdgroup, lane) - with a barrier between every read phase and
write phase. Stage order and arithmetic are kv_q4.metal's
fwht256 (the CUDA register order), so the fused rotation is
bit-identical to the standalone one. The indexer chain (qidx,
tail, dead, pooled) and qfull's gate half never rotate.
qsa_attn reads the q4 pools with every block dequantized to fp32 as
(nibble - 8) * f32_from_f16(d) - the PC native fused path's
arithmetic (qsa_decode_attn.cu's load8_q4), a declared
rounding-path choice stated in the file header; the score and
accumulate math is unchanged.
qsa_merge the merged output row rotates back (the orthonormal FWHT is its
own inverse) after o = o/L, BEFORE the sigmoid(gate) multiply -
the elementwise gate does not commute with the rotation, and the
PC gates b.attn only after fwht256_inplace_cuda. The kernel gains
sg/lane attributes for the staging; msum is free until its own
later use.
Every branch is gated so an fc-off specialization compiles to today's code:
the full current metal test suite is unchanged-green, and qsa_parity's printed
output is byte-identical before and after (checked by diff). qsa.metal now
references kv_q4.metal's prelude, so the compile lists that name it (the
tests' four; the engine window's two are Task 3's) prepend kv_q4.metal, the
same concatenation convention qsa2.metal and qsa_prompt.metal already use.
…V pool bit-level
Drives the whole q4 decode window through mtl_rt Constants (fc 25 on, fc 24
explicitly off): a context of P cells (default 37) built by the fc-25 qsa_prep
plus a 3-token window of prep, score, top-k, attention, merge, against the CPU
reference of kv_q4_ref.h (the device-order FWHT and q4 packing). Checked:
* the V pool's 18-byte blocks BIT-LEVEL against q4_block_ref(fwht256_ref
(vcur)) - the v rows skip norm/RoPE, so the kernel's input bits are known
exactly and the fused cross-simdgroup FWHT + pack chain is verified end to
end. 0 of 640 / 8064 / 32896 blocks differ at P = 37 / 501 / 2053. One
cell's v rows are all zero: their blocks pin the documented all-zero bits
(d = -0.0 -> 0x8000, every qs byte 0x88);
* the K pool DEQUANTIZED against the double-precision norm + RoPE + FWHT
reference at a quantization tolerance (measured 8.2-8.8e-2, limit 0.12:
q4's code step is max|x|/8, the int8 check's 6e-3 is the same derivation
at max|x|/127) - K passes through the fp32 norm/RoPE on the GPU, so its
block bytes are not end-to-end bit-comparable, V above is;
* qcur (the rotated q attention rows) at 1e-5, qidx and the pooled/dead
indexer keys UNROTATED at 1e-5 (4.8e-8 to 1.2e-7 measured), the scores at
1e-5, the ids exactly - identity below and the active selection above the
2051 budget (P = 2051 / 2053);
* attention + merge from the GPU's own qcur, ids and pool bytes: the
reference dequantizes both pools, softmaxes in double, rotates the
256-wide output back with fwht256_ref BEFORE the sigmoid(gate) multiply -
1.9e-7 to 2.7e-7 measured; the wrong placement (gate first, rotate after)
is off by 6.2e-1, printed for contrast, so the check discriminates the
placement by six orders of magnitude.
… wording fixes Review fix round (controller-ruled scope: the compile-list lines of metal/src/engine/* only - no pools, keys or fc wiring; that stays Task 3). qsa.metal's file-scope static_assert references kv_q4.metal's prelude symbols, so EVERY compile of qsa.metal needs the prepend - the tests had it from 4fc0833, but the engine binary compiled qsa.metal bare at THREE sites: window.cpp's decode list ({"qsa.metal", "qsa2.metal"}) and prompt list ({"qsa.metal", "qsa_prompt.metal"}), and prompt_chain.inc's chunk list ({"qsa.metal", "qsa_prompt.metal"}, included by window.cpp:2601). All three now prepend kv_q4.metal, exactly as the tests do; the lists are otherwise unchanged. Two wording minors from the review: * kv_q4.metal's fwht_step comment no longer claims 'every butterfly' of fwht256_lane routes through it - fwht256_lane's high-3-bit register butterflies stay inline (one-thread index arithmetic, nothing to share); it is fwht_row256 that routes all of its stages through fwht_step; * qsa_q4_parity.cpp's header now states what is asserted directly (qcur rotated, qidx unrotated) and what is covered only transitively (pooled / dead, through the score check's reference reading the GPU's rows and the tail's +1e9 bias assertion). Covering tests, Apple M4 Pro: qsa_q4_parity ALL PASS at P = 37 and 2053 (the V pool still 0 blocks differ bitwise, attention+merge 2.7e-7, the wrong gate placement off by 6.2e-1), kv_q4_parity / qsa_parity 3000 ALL PASS, ctest (runtime_hello, kv_q4_engine) 2/2, qsa_prompt_parity ALL PASS - its list is the engines' prompt-list verbatim. The strata binary starts and exits cleanly on --help (0) and bad args (2, usage); its Window::init MSL compile needs a loaded model, and the identical lists are what the green tests above compile.
…lace (pools, qsa keys/constants, the k8v4 startup check, the prompt-path choice); kv_format_test pins the stale-key trap behaviorally
…no scale pools), fc 25 in both the qsa pipeline keys and the Constants (P() looks the key up first), the MTP drafter's pools in the same format, k8v4 refused at startup, q4 prompts kept on the decode window's slow path (one log line, the arena planned without prompt-GEMM slots); the snapshot ruling noted at checkpoint_save
… the chain is int8-only and only the prompt_run_rows guard keeps q4 off it (its attn/merge dispatch under plain keys; prep alone is format-aware); report wording fixed (test library fidelity, stale line citations)
…e two bare prompt compile sites (a dropped prepend is the static_assert's loud MSL error), entry guards on Window::prompt/prompt_chunk under q4 (int8-only chains, M2 removes them), qsa_q4_parity's default run covers the >2051 radix-selection regime, kv_format_test compiles and dispatches the mrope+q4 pair
…d function constant 25 - the pools decode into both right inputs for both cooperative layouts (per-element for !k4, one 4-run a lane for k4: its four k sit in one block, four consecutive bytes on one nibble side, the fp32 product rounded to half like the int8 arm), the scale caches unread under the constant so the host binds a placeholder (the two vsc preloads guarded, folded away with fc off), and qsa_flash's epilogue rotates the accumulated row back through sp (dead staging, fwht_row256's stage order) before the /l and the gate - qsa_merge's placement, which the elementwise gate does not commute with; qsa_prompt_parity's q4 cases (kv_q4_ref-built pools, a fresh RNG stream) check mm and flash against the fc-25 attn+merge chain at 3e-3, discriminate all three epilogue placements, and compile the k4-layout pair, dispatched only on k4 hardware
…double-precision references built from the raw ik fixture (the score stage's reference read the GPU's own pooled/dead rows, so a specialization that polluted them passed it with the same wrong bytes on both sides), and a fifth pass over the M-RoPE + q4 pair (fc 24 and 25 together, a vision prompt under --kv q4_0) with qsa_parity.cpp's random (t, h, w) table - the identity table kv_format_test dispatches cannot tell 'reads mtab' from 'ignores it': the K pool dequantized against the per-pair table positions, qcur/qidx likewise, the V pool bit-level again, and the plain-position reference asserted past the K tolerance (measured 4.8e-01/5.1e-01 against a 0.12 pass); both discriminators counterfactual-checked: pooled un-roped and mtab-ignored each fail exactly the new checks
… out (M2-b) - the five qsa dispatch sites carry fc 25 in BOTH the pipeline key and the Constants (Window::P() looks the key up first: a key without it would hand back the cached int8 specialization over q4 pools): qsa_attn_mm keyed qsa_attn_mm#k4#q4 with fc 7 beside fc 25, qsa_flash qsa_flash#q4 with fc 7 bound unconditionally (its MSL default is true when left out, so a !k4 site must set it false, not drop it - the fc-off Constants stay byte-identical to today's b(7, p.k4) either way), plain qsa_attn/qsa_merge through the decode precedent's qsa_q4_* - so the entry guards on Window::prompt/prompt_chunk, prompt_run_rows' fallback log and prompt_fast_path's fold in stream_arena's pg_on are gone (the arena's min_slots 576 under q4 again, STRATA_METAL_PROMPT=window untouched, prompt_fast_path now true in every format as the one switch point); the chains' ksc/vsc slots need no change: they bind the window members that already hold the 16-byte no_scale placeholder under q4, the same ones decode's fc-25 attn/merge bind, dead-but-harmless before only because the guards kept q4 off the chains; with q4 off every key, Constants set and binding order is today's int8 one, byte for byte - kv_format_test pins the five-key reduction and extends the stale-key counterfactual to the prompt library (every constant set an engine site now binds compiled as its own specialization), and prompt_chain gains --kv q4_0 to run the harness over q4 pools through the same kv_format.h source (engine-level q4 prompt execution deferred to the authorized measurement round)
…-hardware constant set too ({fc 7 true, fc 25} under the same qsa_flash#q4 key as the !k4 set: one process probes one layout, so the engine's cache never holds both under it - compiled beside the cache as its own specialization, the file's every-constant-set-a-site-binds claim now whole, check 36 of 36 on this !k4 host that would never probe it), qsa_flash's int8 vs preload comment states the truth (written but never read - the V arm reads vsc directly; the array and its write stay, removal deferred out of this branch), and qsa_prompt_parity's three-placement epilogue discriminator notes it runs at pos0 0 only (the placement error is the same arithmetic at any pos0; the 30000 flash case stays covered by the rel-L2 reference check)
…uality's run/prompt gain an opt-in --kv that replaces FIXED's fp16 pin only when passed (engine_args() byte-identical without it, diffed over the three gpu_experts modes; compare/pcompare untouched), and bench/results/2026-10-15-mac-kv-q4/ lands the round on the strata-q2_0.json template, same day, only --kv differing: teacher-forced quality over 4,995 positions (q4_0 top-1 94.45%, mean KL 0.034, p99 0.447, ppl 4.517 vs 4.442 = +1.7%, worst segment the agent transcript at +7.3%), with the framing fact measured outright: on Metal fp16 and int8 are one non-q4 path (kv_format.h, 'the window's pools are the int8 ones either way') and the two runs' LOGPOS TSVs are byte-identical, so the fp16 reference IS the shared path; the batched prompt path under q4 (M2-b) answers 20 full-vocab prompts at top-1 100%, KL mean 0.036 / median 0.004 / max 0.195; paired speed has q4_0 at -1.2% decode (20.84 vs 21.09 tok/s) and -1.1% 2200-id prefill (213.2 vs 215.5), both at the ~1% edge the same-code none-vs-int8 pair bounds as run-to-run noise; KV bytes arithmetic (cells x 12 QSA layers x 1056/576 B) meets the log's before-the-arena delta exactly (0.189 vs 0.19 GB at 32768+128 cells, 0.378 vs 0.38 at 65536+128) and the saving lands as +241/+494 resident expert slots in the arena; needles 8/9 under q4_0 (the one miss a 32K@90% injection-suspicion refusal, not a garbled retrieval; every hit the exact codeword) with int8 controls 3/3 - all scripts (speed.sh, kvmem.sh, needle.py over the committed set.ids.json filler) reproduce every number, and the README states the not-testeds (M5 k4 numeric arm on this !k4 host, MTP+q4 end to end, PC consistency, k8v4)
…path pair regenerated exactly as README 'Reproducing' records it (the round's only 2 engine loads: prompt --kv fp16 and --kv q4_0 into a fresh temp dir, the original first.* logits having lived in /tmp and died there) and the 20 per-prompt rows committed as data/prompt-path/per-prompt.tsv so the claim is checkable offline (segment, prompt tokens, top-1 agree, full-vocab KL; raw logits still not committed at ~20 MB/run); the regenerated aggregates are identical to the committed round - top-1 100.00%, KL mean/median/max 0.03565/0.00434/0.19546, the streaming expert set did not drift on these prompts - so no number in the README moved, and pcompare.txt's reference line now points at the committed table instead of the dead /tmp/t6 path; the README also gains the four disclosures the review asked for, one thought each: the needle question ends with 'The access code is' (a completion cue all arms share, so the q4-vs-int8 comparison is internally fair but the 8/9 is under an easier phrasing than a neutral QA form), the single miss has no 90%-depth int8 control (controls ran at 50%, so a shared-path refusal there cannot be excluded) and its reply was cut mid-sentence at the 64-token budget (the injection-suspicion reasoning is explicit, a final refusal sentence is not), '2200-id prompt' now reads 2199 tokens prefilled with the last id seeding generation wherever 2200 stood alone as prefilled tokens (the engine's own line: 'prefill 2199 tokens in 1 chunks'), and the kvmem RAM note closes the brief's watchdog item with the truth that the Metal logs carry no watchdog line and the RAM planning lines are what was copied; tools/quant_quality.py untouched (ast.parse clean)
…3-b) - MAC.md gains the KV-cache memory bullet (1,056 vs 576 B per cell over the 12 QSA layers, one cell per token of context plus 128; measured on the M4 Pro, 48 GB: 0.19 GB less at 32K and 0.38 GB at 64K, which the streaming planner spends as +241/+494 resident expert slots), the Options line (--kv q4_0, opt-in; without it setup writes no --kv at all and the engine's default path is unchanged) and a Quality bullet with the 2026-10-15 round's numbers digit for digit (teacher-forced top-1 94.45%, mean KL 0.034 / p99 0.4467, ppl 4.517 against 4.442 = +1.7%, worst segment the agent transcript +7.3%; the prompt fast path top-1 100% over 20 full-vocab prompts, KL 0.03565/0.00434/0.19546; decode 20.84 against 21.09 tok/s = -1.2% and 2200-id prefill 213.15 against 215.54 = -1.1%, both at the ~1% edge the same-code fp16-vs-int8 pair bounds as noise; needles 8/9, the one miss the 32K@90% injection refusal, int8 controls 3/3) - the bullet states the one-path fact outright (on Metal --kv fp16 and --kv int8 are the same int8 pools; the log prints KV int8 for both and the round's two TSVs came out byte-identical), so no fp16-vs-int8 difference is invented, and it names the not-testeds (the M5 k4-layout kernels, MTP over q4 pools end to end, other Macs) with the bench directory as the entry point; setup_mac's engine_args gains kv=None and returns exactly the old list without it (return args + [--kv, kv] if kv else args), argparse offers choices=[q4_0] only so k8v4 stays the engine's own startup refusal, and tools/test_setup_mac.py pins it with one new test (assertNotIn --kv on the default, q4[:-2] == base, the streaming arm, the source-level choices line) - the file passes 20/20 with no GPU and no downloads, and engine_args/mac_config were diffed against HEAD over a 48-cell matrix (stream x ram x context x shard shape, 0 diffs) plus three full generated configs: byte-identical without the flag, append-only with it
…t stops promising engine logs the repo's root logs/ ignore keeps out of the PR (git check-ignore names .gitignore:14; the 2026-09-27 PC round committed only a README, so no-logs is the repo convention, and nothing was force-added): the speed/kvmem/needle section headers and the Files entry now point at the committed distilled run lines (data/speed-runs.txt, data/kvmem-runs.txt, needle-results.tsv, the data/*/*.stderr.txt summaries) and state that the full engine logs stay out per the ignore and regenerate from the committed scripts; the time-to-first-token delta is now +1.1% computed from the raw ms (10714.3 vs 10600.5, data/speed-runs.txt) with the basis written next to the table, whose one-decimal seconds (10.71 vs 10.60) would give +1.0%; and the prompt-path KL median column says upper-middle (pcompare's kls[n//2] of the sorted 20, so a recomputer averaging the middle pair finds 0.00388, not a mismatch) - every measured value in the tables is untouched, the only number that moved is the derived TTFT percentage
…ut q4_0 is now refused at Window::init (the env generate.cpp parses for the PC's rotated int8/fp16 K/V; this backend builds UNrotated pools while the shared snapshot tag adds +16 for the rotation, so a saved session would lie about its bytes - refused like k8v4, with the env parsed here because there is no shared getter; next to --kv q4_0 the rotation is intrinsic and the tag never adds it, so the env there stays a harmless no-op), qsa_prompt_parity gains a device-free equivalence pin of the two q4 decode arms (the k4 4-run's byte/side mapping vs the !k4 per-element formula, 512k values - this !k4 host cannot dispatch the k4 layout, so the arms are pinned to one arithmetic and the !k4 arm's numeric parity anchors the pair; the comparison is on the extracted nibble, not the fp32 product, because a random fp16 d can be NaN and NaN != NaN), qsa_prompt.metal states the cross-platform rounding note the decode header has (the PC's prompt kernels apply the q4 scale in fp32; this file rounds the operand to half as a matmul2d right input, the same half-rounding the int8 arm always had), and the bench README + MAC.md carry the missing not-tested: long-context quality (every segment is 1,000 tokens; the PC round measured the q4_0 cost growing with context, so +1.7% perplexity is a short-context number) and the load count now includes the regeneration pair (28)
…prompt delta is computed from the raw ms like every other row (806.4 vs 805.3: -0.1%, not the -0.2% the one-decimal tok/s display gives; the basis sentence now covers the whole table, with all four rows' raw-ms figures), the bench README states the round was measured before the upstream merge b056ff8 and what was re-checked after it (the metal tests and the int8/q4_0 determinism runs, byte-identical; the numbers were not re-measured), MAC.md counts three measured Macs now, carries a q4_0 Limits line and points out that setup's default config keeps the MTP draft layer on a resident Mac, so ./setup.sh --kv q4_0 there runs the not-measured MTP-over-q4 combination, and labels the KL median upper-middle-of-20; setup warns instead of silently ignoring --kv q4_0 when a config is already installed (the early-start path names ./setup.sh --setup as the way to change it; the default path is untouched), and the 48-cell engine_args matrix behind the PR's gate 3 - now including the repacked -metal.gguf shard arm - is a test (append-only over stream x ram x context x both shard shapes); tools/test_setup_mac 22/22. The review's independent metal-cpp republish finding did not reproduce: the pin's own text contains the hash segment the reviewer called new, four same-day downloads and the M1-era zip all hash to the pin, and the origin's Last-Modified is Jan 30 - no change made
…used next to --kv int8 only, and allowed next to fp16 where it is a no-op on both engines (the PC gates the rotation on the int8 flag - layer.cpp: kv_rot = kv_q4 || (kv_int8 && env), and generate.cpp sets kv_int8 only for --kv int8, its own notice printing for int8 alone - so with fp16 nothing rotates and the snapshot tag adds nothing: the old refusal blocked a combination byte-identical to the env unset, stricter than the PC's silent no-op it mirrors; kv_format_check gains the int8 flag through the existing qsa_kv_int8 getter, no shared file touched, the message now says 'the PC's rotated int8 K/V', and kv_format_test pins the allow, the narrowed refusal and the env-off case); kv_format.h includes cstdlib for its std::getenv; and the post-merge determinism claim carries its artifact now - data/post-merge-checks.txt records the four g.sh runs (pre-merge f1/f4 on 521b221-era builds, post-merge f5/f6 on b056ff8), the two output token lines with their md5s (int8 4d77bea6..., q4_0 9e398b5e...), and the whole-log diff statement (only timing, throughput, memory-planning and version-string lines differ), and the bench README's provenance note points at it; MAC.md's Limits line narrows to match; kv_format_test ALL PASS
Owner
|
Thanks for the pull request. We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to review properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected. To reopen it takes about two minutes:
The rules are in CONTRIBUTING.md. |
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.
What it changes
The PC engine's q4_0 K/V format (PR #21's FWHT-256 Hadamard rotation) brought to the Metal backend, opt-in
through the engine's existing
--kv q4_0flag:metal/kernels/kv_q4.metal(new): the FWHT-256 rotation, quantize and dequantize kernels, and the poollayout (576 B per cell per QSA layer against the non-q4 path's 1,056).
metal/kernels/qsa.metal,qsa_prompt.metal: the decode and prompt chains made format-aware — the fiveQSA dispatch sites carry function constant 25 in both the pipeline key and the Constants, so a q4 run gets
its own specializations instead of silently reusing the cached int8 ones over q4 pools.
metal/src/engine/:kv_format.h(new, the single source for pool sizes / keys / the format choice),kv_q4_host.cpp(new, the host side),window.cppandprompt_chain.incwired: q4 pools at init, the MTPdrafter swaps pools in the same format,
--kv k8v4refused at startup.fallback that read q4 prompts through decode windows is gone,
STRATA_METAL_PROMPT=windowstill forces it.metal/setup_mac.py: opt-in--kv q4_0that lands in the generated engine config'sargs(
tools/test_setup_mac.pyextended).tools/quant_quality.py: opt-in--kvonrun/prompt(replaces the harness's fixed fp16 pin only whenpassed).
docs/MAC.md: the option, the quality/speed numbers, the KV-memory numbers, the not-testeds.bench/results/2026-10-15-mac-kv-q4/(new): the whole Mac measurement round — scripts, raw TSVs, distilledrun lines, README.
What it leaves alone
src/orinclude/is touched, so no CUDA/HIP/SYCL build is affected (the Metal engine'shost code is the PC's, compiled for macOS as before).
byte for byte (pinned by tests, below). Defaults everywhere: setup writes no
--kvunless asked, theengine's default is the same non-q4 path as the last release.
refused by the shared code, as before).
Opt-in, and how the default path was checked byte-identical
Three independent gates:
metal/tests/kv_format_test.cpppins the five-key reduction (everyconstant set an engine site now binds is compiled as its own specialization) and a stale-key
counterfactual;
qsa_q4_parity.cppandqsa_prompt_parity.cppassert the fc-off outputs againstdouble-precision references — with q4 off every key, Constants set and binding order is the current int8
one, byte for byte.
output byte-identical to the pre-branch baseline run (bd1f315), zero stub lines — the engine's int8 path
does not move.
tools/quant_quality.py'sengine_args()output was diffed before vsafter the opt-in
--kvchange over the three gpu_experts modes: identical without the flag. For setup,metal/setup_mac.py'sengine_args()was compared against HEAD over a 48-cell matrix(stream x ram x context x shard shapes; 0 diffs) plus the full generated config on three of those cells:
identical without
--kv, and with it the pair--kv q4_0is appended at the end and nothing else moves(pinned by
tools/test_setup_mac.py).Measured (Apple M4 Pro, 20 cores, 48 GB unified memory, macOS 27.0.1; Qwen3.8-Flash-Next Q2_0, experts
streamed from the SSD; engine built from this branch; every number same-day 2026-10-15, nothing extrapolated)
The framing fact for the reference column: on Metal,
--kv fp16and--kv int8are one path — thewindow's pools are the int8 ones either way (
kv_format.h; the fp16-flagged run's own log printsKV int8), and the two teacher-forced TSVs of the round are byte-identical(md5
4e942a7375715ac014956fbdbf6110f6). So the reference below is that shared non-q4 path, q4_0 is theonly smaller K/V format on Mac, and no fp16-vs-int8 delta exists on this backend.
fp16-vs-int8 pair (one path on Metal) measured −1.3% decode, bounding noise near 1%. Single runs per cell.
to repeat it (a behavior miss, not a garbled retrieval; all 8 retrievals were the exact codeword).
Disclosures, as recorded in the bench README: the question ends with "The access code is" — a completion cue
all arms share, so the comparison is internally fair but the 8/9 is under an easier phrasing than a neutral
QA form; the miss has no int8 control at 90% depth (controls ran at 50%), so a shared-path refusal there
cannot be excluded; and its reply was cut mid-sentence at the 64-token budget — the injection-suspicion
reasoning is explicit, a final refusal sentence is not.
Disclosures (from the final segment review)
plans 576 stand-by slots for the prompt-GEMM chain (64 without it). Escape hatches:
STRATA_METAL_PROMPT_GEMM=0(prompts off the matrix-unit GEMM chain; the arena is then planned without thestand-by slots) and
STRATA_METAL_PROMPT=window(prompts through decode windows — the slow, exactfallback).
coop_k4=false), so the prompt fast path behindevery prompt-path and prefill number above is the matrix-unit GEMM chain over the v1 layout — the flash chain is off
by default on !k4 hardware and never dispatched here (same hardware fact as the k4 arm below).
chunk (the matrix units), so its numerics are a different floating-point order from the decode window's
(fp32 fused) — outputs can differ on near-ties, the same class as any reordering. Measured under q4_0 on
the 20-prompt set: argmax never flipped (top-1 100%), full-vocab KL as in the table.
this M4 Pro is !k4 hardware and never dispatches them (see Not tested).
--kv k8v4: the only statement is the explicit startup refusal — "metal window: --kv k8v4 is notimplemented on Metal (use fp16, int8 or q4_0)". No partial implementation exists.
STRATA_KV_ROT=1(the PC's rotated int8 K/V env): refused at startup next to--kv int8, like k8v4 — therethe PC rotates the pools and the shared conversation-snapshot tag adds +16 for the rotation, so a Metal session
would save a tag claiming a rotation its pools do not have (the env is parsed in kv_format_check as
generate.cpp does; the int8 flag comes through the existing qsa_kv_int8 getter). Next to fp16 (and no
--kv)the env is a no-op on both engines — the rotation is gated on the int8 flag (layer.cpp:
kv_rot = kv_q4 || (kv_int8 && env), and generate.cpp sets kv_int8 only for--kv int8), so nothing rotates and the tag addsnothing; allowed, matching the PC's silent no-op. Next to
--kv q4_0the rotation is intrinsic and the tagnever adds it, so the env there is a harmless no-op too.
Not tested
The M5 k4-layout q4 numeric arm (fc-7 right-input layouts): compiled but never dispatched on this !k4 host.
MTP + q4_0 end to end: every measured run used the template's
--spec 2 --suffix-draft 0; the drafter'ssame-format pool swap is structurally validated by tests, not exercised end to end.
PC consistency: no PC run this round; the 2026-09-27 PC table (
bench/results/2026-09-27-kv-q4/) is adifferent backend and a
--kv-residentstreaming configuration, so its numbers are not comparable and arenot cited.
Speed repeats (one run per cell; the same-code pair bounds noise near 1%) and needle breadth (one codeword,
one filler corpus, one run per case, no distractors, int8 controls at 50% depth only).
Other Macs: every number above is from one M4 Pro, 48 GB.
Long-context quality: every quality segment is 1,000 tokens. The PC round measured the q4_0 perplexity
cost growing with context (+8% at 1K to +12% at 8K on a document), so the +1.7% here is a short-context
number and must not be read as the format's cost at 32K.
Known follow-ups, none in this PR: the flash q4 epilogue is barrier-dense (a perf item only k4 hardware would
feel, and this !k4 host never dispatched it), the flash
vsdead array is still to be removed,prompt_chain's pre-existing argv bounds pattern and kv_format_test's pre-existing non-q4 flash constant sets
predate the branch and stay as found, and the Task-1-era nits stand (the
pipe()per-call cache lookup, theidentity-page gather case, the offsetof asserts).
Reproducing
bench/results/2026-10-15-mac-kv-q4/README.mdis the entry point: the exact commands for theteacher-forced round (
tools/quant_quality.py run/compare), the prompt-path round (prompt/pcompare), thespeed pairs (
speed.sh), the KV/arena numbers (kvmem.sh) and the needle matrix (needle.py), all with thecommitted raw TSVs and distilled run lines beside them (
data/speed-runs.txt,data/kvmem-runs.txt,needle-results.tsv, the stderr summaries); the full engine logs stay out per the repo'slogs/ignore andregenerate from the committed scripts. The in-tree test gates:
ninja -C build-metalthen the binaries inmetal/tests/(kv_q4_parity,kv_q4_engine,kv_format_test,qsa_q4_parity,qsa_prompt_parity).Setup's default-identity is
python3 -m unittest tools.test_setup_mac(no GPU, no downloads).note:
data/prompt-path/pcompare.txt's reference line was edited by hand to point at the committedper-prompt.tsv (the regeneration commands are in the README); the aggregates were never hand-touched.
Prior art searched (issues and PRs)
Re-checked read-only at packaging time (
gh search issues/gh pr viewagainst the upstream repo, terms: metal kv,kv q4, k8v4, mac backend, apple silicon, metal PRs): no issue or PR covers KV quantization on the Mac/Metal
backend — this PR opens that ground, and the PC lineage below is the prior art.
--kv k8v4: hybrid KV cache — INT8 K + Hadamard-rotated Q4_0 V (816 B/cell) #120 (k8v4, closed) and KV: optional Hadamard rotation for int8 K/V (STRATA_KV_ROT=1), with the drafter and batched verify made rotation-consistent #293 (int8rotation, closed) are the format's siblings; PR qsa_prompt_attn: Q4_0 KV on tensor cores (mode 4) - long prompts with --kv q4_0 ~19% faster #452/Verify window: PDL on sm_90+, graph branches, batched Q4_0 KV append and PLE post-ops (bit-identical) #904 are the PC prompt/append follow-ups.
--kv q4_0faults on a PC (HIP) when the KV itself streams (--kv-resident) with MTPdrafts. Different code path — the Metal engine has no
--kv-residentKV streaming (a Mac streams experts,not KV) — noted here so nobody reads the PC issue as covering this backend. The k8v4 PC streaming bugs
(k8v4: deterministic token repetition on sequential generation (RTX 5090, 0.1.40) — byte-identical across two quants #1135, --kv k8v4 produces degenerate repetition when --kv-resident streams (0.1.40) #1188, k8v4 + --kv-resident: the batched KV append (#783) skips the host mirror for hybrid — parking and KV streaming serve stale bytes (cross-conversation contamination on 0.1.40) #1169, all closed) are the same story:
--kv-resident-path-specific, not applicable here.moe_mm_parity's default tile does not exist in the v1 kernel set. Pre-existing onmain, reproduced byte-identically on a clean tree, untouched by this branch (that test, its kernels and the
failure are outside this diff; the branch's test sweeps treat it as the known exception).