engram: make the ggml batch-cosine strategy actually compute in fp32 (recall 0.9933 -> 0.9987) #121

Merged

will.anderson merged 1 commits from improve/ggml-cosine-fp32-and-init into feat/engram-ggml-cosine-batch 2026-08-15 23:01:31 +00:00

Conversation 0 Commits 1 Files Changed 1

+235 -26

will.anderson commented 2026-08-15 22:58:03 +00:00

Owner

Copy Link

Copy Source

Follow-up to #116. Base is feat/engram-ggml-cosine-batch (#116's branch), not dev — #116 is still open, so dev does not contain eg_cosine_batch_strategy_ggml.c yet. Targeting #116's branch keeps this diff to just the improvement; retarget to dev if #116 merges first. Nothing in #116's history is rebased, amended or rewritten.

Why

Will, on #116's own numbers (ggml recall 0.9933 vs the hand-rolled kernel's 0.9997): "we can keep ggml, but why do it worse?"

Answer: we don't have to. The precision gap was a real, fixable usage bug in how this file called ggml — not a ggml limitation. The cold-start gap was not a bug at all. Both are settled below with measurements, including the fixes that didn't work.

The precision bug

ggml-metal has two F32xF32 matmul kernels and picks between them purely on ne11 — the number of B rows, which for us is the query-batch size:

ne11	kernel	template
<= 8	kernel_mul_mv_ext_f32_f32_*	<float, float> — genuine F32
> 8	kernel_mul_mm_f32_f32	<half, half4x4, simdgroup_half8x8, half, half2x4, ...> — both operands narrowed to F16

Both operands are narrowed to F16 even though the tensors are GGML_TYPE_F32 on both sides. Verify it yourself, no guessing — the kernel templates are literal strings in the shipped plugin:

$ strings $(brew --prefix ggml)/libexec/libggml-metal.so | grep -E 'host_name\("kernel_mul_m[mv]_f32_f32'
template [[host_name("kernel_mul_mm_f32_f32")]]   kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, ...>;
template [[host_name("kernel_mul_mv_f32_f32")]]   kernel mul_mv_t_t kernel_mul_mv_t_t<float, float>;

#116 issued one ggml_mul_mat with ne11 = nq (300 in the benchmark), landing squarely on the F16 path. ggml confirms the pick at runtime:

ggml_metal_library_compile_pipeline: compiling pipeline: base = 'kernel_mul_mm_f32_f32'   <- before
ggml_metal_library_compile_pipeline: compiling pipeline: base = 'kernel_mul_mv_ext_f32_f32_r1_4'  <- after

The file's own header comment asserted the opposite ("computes in F32 on the Metal backend"). That claim was wrong; it is replaced with the measurement.

Fix

Emit ceil(nq/8) mul_mats over ne11<=8 ggml_view_2d slices of one query tensor, all expanded into one graph and one ggml_backend_graph_compute — so the node matrix is still uploaded and shared exactly once, which is the entire reason batch_multi exists. EL_GGML_MULMAT_CHUNK overrides the 8; setting it >= nq reproduces the old behaviour exactly, which is how the before/after below was measured in a single binary.

8 is not invented — it is ggml-metal's own threshold, found by sweeping ne11 and watching both the error and which pipeline ggml compiles. 9 flips to mul_mm and the error jumps back in the same step:

ne11=6    ( 50 groups)  max_abs=3.460e-08  mean_abs=3.863e-09     8.38 ms
ne11=8    ( 38 groups)  max_abs=3.460e-08  mean_abs=3.863e-09     7.70 ms
ne11=9    ( 34 groups)  max_abs=7.978e-05  mean_abs=1.027e-05     5.65 ms
ne11=16   ( 19 groups)  max_abs=7.978e-05  mean_abs=1.038e-05     3.31 ms

Results — real store, offline snapshot

vindex_bench store <snapshot> 768 300 10 64 against an offline copy of the store (13,415 live embedded nodes, dim=768, 300 real queries, CPU double-accumulated oracle). No live service on :8742/:7770 was touched.

	id-recall vs oracle	same-rank |Δdist| max	mean
ggml old (ne11=300)	0.9933	6.80e-05	1.43e-05
ggml new (ne11<=8)	0.9987	4.77e-07	9.30e-08
hand-rolled Metal (untouched)	0.9997	3.58e-07	7.55e-08

~145x better max error, ~154x better mean — now the same order of magnitude as the hand-rolled kernel instead of 150x off it. These three lines are bit-identical on every repeat run.

The cost, stated rather than buried

The ms/query column in vindex_bench is too noisy on this machine to carry a claim (the untouched hand-rolled strategy swings 0.064–0.131 across runs). So speed was measured separately: median of 15 reps of the whole batch_multi() call — gather + norms + upload + GPU + scatter — after a discarded warm-up, three separate runs:

	median batch (300q)	ms/query
ggml old (unchunked)	13.19 / 13.35 / 14.42 ms	~0.044
ggml new (chunked)	19.92 / 20.08 / 20.23 ms	~0.067
hand-rolled Metal	17.74 / 17.88 / 17.99 ms	~0.060

Correctness costs ~+6.7ms per 300-query batch (~1.5x on this call), and leaves ggml ~12% behind the hand-rolled kernel instead of ~35% ahead of it.

This is not recoverable inside ggml. An fp32 matmul on Metal must re-stream the whole node matrix once per <=8 queries (38 dispatches x ~41MB here), where the F16 mul_mm kernel tiles it in threadgroup memory and reads it far fewer times. ggml's Metal backend ships no fp32 tiled matmul, so on this backend "fast" and "fp32" are genuinely exclusive. The hand-rolled kernel escapes the choice only because it is an fp32 kernel written for this one shape.

Decision for you, not taken here: the factory still prefers ggml first, per your original directive. On this hardware and shape ggml is now measurably behind the hand-rolled kernel on both axes (0.067 vs 0.060 ms/query, 0.9987 vs 0.9997). Flipping the default preference order is a policy call, so it is left alone and flagged rather than changed unilaterally.

Two things that did NOT work

Recorded so nobody retries them.

• ggml_mul_mat_set_prec(t, GGML_PREC_F32) does nothing here. Tried it; error was bit-identical with and without — 1.038e-05 either way. ggml-metal has no F32-accumulating mul_mm kernel to switch to, so the flag has nowhere to go. ne11 is the only lever.

  A: mul_mat ne11=nq (current)  max_rel_err=3.045e+01  mean_rel_err=2.240e-03   8.0 ms
  B: + GGML_PREC_F32            max_rel_err=3.045e+01  mean_rel_err=2.240e-03   5.2 ms
  

• The ACCEL/BLAS device is a mirage. In an isolated compute-only probe it looked excellent — 3.4–4.0ms at mean |Δdot| 1.5e-08, i.e. as fast as the old F16 path and far more accurate. End-to-end on the real store it is dominated on both axes: 0.191 ms/query at 0.9973 recall, vs 0.125–0.142 at 0.9987 for the Metal default. The probe wasn't competing for the same CPU cores the real call path is. It stays reachable via EL_GGML_DEVICE=accel as a no-Metal fallback, labelled measured-and-rejected.

Cold start — not a bug, and not ours

The ~7.8s first-call cost attributed to this file is not this file re-initialising per call. Init was already cached behind g_init_attempted (and always was — the hypothesis that it re-ran ggml_backend_load_all_from_path every call is false; see the original lines 109–111).

It is Apple's Metal shader cache missing on ggml's embedded metallib (~650 kernels in one __ggml_metallib section). That cache is keyed on the library, not our binary, and is shared across processes:

first-ever load  : ggml_metal_library_init: loaded in 7.670 sec
next run, different binary, same ggml : ggml_metal_library_init: loaded in 0.009 sec

Corroborated by the cache itself — ~/…/C/com.apple.metal/32024/libraries3.data grew to 48MB at the moment of the 7.67s run. So: once per machine per ggml version, not once per process, and nothing this file does can avoid it. The hand-rolled strategy escapes it only because its shader is two small kernels instead of six hundred.

Warm ggml init is 44–53ms, against 36–117ms for the hand-rolled strategy's device+pipeline setup — comparable.

Loading only libggml-metal.so instead of every plugin in the directory is kept for tidiness and explicitly documented as not a speedup: 44.7/51.2/52.4ms against 46.9/55.1/58.9ms, the same number inside noise.

Contract

The -2.0 sentinel contract is unchanged and re-verified at batch sizes straddling the chunk boundary (1, 7, 8, 9, 16, 17, 33), plus NULL rows, dim mismatches, zero-norm rows, and an all-invalid population — all pass, under the default and under EL_GGML_MULMAT_CHUNK=1/300 and EL_GGML_DEVICE=accel/cpu.

Independent confirmation of the defect: the old ne11=300 path fails that same check at a 2e-6 cosine tolerance with 2299 mismatches, while the new default passes with 0.

Scope

One file: lang/runtime/eg_cosine_batch_strategy_ggml.c. eg_cosine_batch.metal and eg_cosine_batch_strategy_metal_hand.m are untouched — read only as reference.

Follow-up to #116. **Base is `feat/engram-ggml-cosine-batch` (#116's branch), not `dev`** — #116 is still open, so `dev` does not contain `eg_cosine_batch_strategy_ggml.c` yet. Targeting #116's branch keeps this diff to just the improvement; retarget to `dev` if #116 merges first. Nothing in #116's history is rebased, amended or rewritten. ## Why Will, on #116's own numbers (ggml recall 0.9933 vs the hand-rolled kernel's 0.9997): *"we can keep ggml, but why do it worse?"* Answer: we don't have to. The precision gap was a real, fixable usage bug in how this file called ggml — not a ggml limitation. The cold-start gap was not a bug at all. Both are settled below with measurements, including the fixes that **didn't** work. ## The precision bug `ggml-metal` has two F32xF32 matmul kernels and picks between them purely on `ne11` — the number of B rows, which for us is the query-batch size: | `ne11` | kernel | template | |---|---|---| | `<= 8` | `kernel_mul_mv_ext_f32_f32_*` | `<float, float>` — genuine F32 | | `> 8` | `kernel_mul_mm_f32_f32` | `<half, half4x4, simdgroup_half8x8, half, half2x4, ...>` — **both operands narrowed to F16** | Both operands are narrowed to F16 *even though the tensors are `GGML_TYPE_F32` on both sides*. Verify it yourself, no guessing — the kernel templates are literal strings in the shipped plugin: ``` $ strings $(brew --prefix ggml)/libexec/libggml-metal.so | grep -E 'host_name\("kernel_mul_m[mv]_f32_f32' template [[host_name("kernel_mul_mm_f32_f32")]] kernel mul_mm_t kernel_mul_mm<half, half4x4, simdgroup_half8x8, half, half2x4, simdgroup_half8x8, ...>; template [[host_name("kernel_mul_mv_f32_f32")]] kernel mul_mv_t_t kernel_mul_mv_t_t<float, float>; ``` `#116` issued **one** `ggml_mul_mat` with `ne11 = nq` (300 in the benchmark), landing squarely on the F16 path. ggml confirms the pick at runtime: ``` ggml_metal_library_compile_pipeline: compiling pipeline: base = 'kernel_mul_mm_f32_f32' <- before ggml_metal_library_compile_pipeline: compiling pipeline: base = 'kernel_mul_mv_ext_f32_f32_r1_4' <- after ``` The file's own header comment asserted the opposite ("computes in F32 on the Metal backend"). That claim was wrong; it is replaced with the measurement. ### Fix Emit `ceil(nq/8)` `mul_mat`s over `ne11<=8` `ggml_view_2d` slices of one query tensor, all expanded into **one** graph and **one** `ggml_backend_graph_compute` — so the node matrix is still uploaded and shared exactly once, which is the entire reason `batch_multi` exists. `EL_GGML_MULMAT_CHUNK` overrides the 8; setting it `>= nq` reproduces the old behaviour exactly, which is how the before/after below was measured **in a single binary**. `8` is not invented — it is ggml-metal's own threshold, found by sweeping `ne11` and watching both the error and which pipeline ggml compiles. `9` flips to `mul_mm` and the error jumps back in the same step: ``` ne11=6 ( 50 groups) max_abs=3.460e-08 mean_abs=3.863e-09 8.38 ms ne11=8 ( 38 groups) max_abs=3.460e-08 mean_abs=3.863e-09 7.70 ms ne11=9 ( 34 groups) max_abs=7.978e-05 mean_abs=1.027e-05 5.65 ms ne11=16 ( 19 groups) max_abs=7.978e-05 mean_abs=1.038e-05 3.31 ms ``` ## Results — real store, offline snapshot `vindex_bench store <snapshot> 768 300 10 64` against an **offline copy** of the store (13,415 live embedded nodes, dim=768, 300 real queries, CPU double-accumulated oracle). No live service on `:8742`/`:7770` was touched. | | id-recall vs oracle | same-rank \|Δdist\| max | mean | |---|---|---|---| | ggml **old** (`ne11=300`) | 0.9933 | 6.80e-05 | 1.43e-05 | | ggml **new** (`ne11<=8`) | **0.9987** | **4.77e-07** | **9.30e-08** | | hand-rolled Metal (untouched) | 0.9997 | 3.58e-07 | 7.55e-08 | **~145x better max error, ~154x better mean** — now the same order of magnitude as the hand-rolled kernel instead of 150x off it. These three lines are bit-identical on every repeat run. ## The cost, stated rather than buried The `ms/query` column in `vindex_bench` is too noisy on this machine to carry a claim (the *untouched* hand-rolled strategy swings 0.064–0.131 across runs). So speed was measured separately: median of 15 reps of the whole `batch_multi()` call — gather + norms + upload + GPU + scatter — after a discarded warm-up, three separate runs: | | median batch (300q) | ms/query | |---|---|---| | ggml old (unchunked) | 13.19 / 13.35 / 14.42 ms | ~0.044 | | ggml new (chunked) | 19.92 / 20.08 / 20.23 ms | ~0.067 | | hand-rolled Metal | 17.74 / 17.88 / 17.99 ms | ~0.060 | Correctness costs **~+6.7ms per 300-query batch (~1.5x on this call)**, and leaves ggml ~12% *behind* the hand-rolled kernel instead of ~35% ahead of it. **This is not recoverable inside ggml.** An fp32 matmul on Metal must re-stream the whole node matrix once per `<=8` queries (38 dispatches x ~41MB here), where the F16 `mul_mm` kernel tiles it in threadgroup memory and reads it far fewer times. ggml's Metal backend ships **no fp32 tiled matmul**, so on this backend "fast" and "fp32" are genuinely exclusive. The hand-rolled kernel escapes the choice only because it is an fp32 kernel written for this one shape. **Decision for you, not taken here:** the factory still prefers `ggml` first, per your original directive. On this hardware and shape ggml is now measurably behind the hand-rolled kernel on *both* axes (0.067 vs 0.060 ms/query, 0.9987 vs 0.9997). Flipping the default preference order is a policy call, so it is left alone and flagged rather than changed unilaterally. ## Two things that did NOT work Recorded so nobody retries them. - **`ggml_mul_mat_set_prec(t, GGML_PREC_F32)` does nothing here.** Tried it; error was *bit-identical* with and without — `1.038e-05` either way. ggml-metal has no F32-accumulating `mul_mm` kernel to switch to, so the flag has nowhere to go. `ne11` is the only lever. ``` A: mul_mat ne11=nq (current) max_rel_err=3.045e+01 mean_rel_err=2.240e-03 8.0 ms B: + GGML_PREC_F32 max_rel_err=3.045e+01 mean_rel_err=2.240e-03 5.2 ms ``` - **The ACCEL/BLAS device is a mirage.** In an isolated compute-only probe it looked excellent — 3.4–4.0ms at mean `|Δdot|` 1.5e-08, i.e. as fast as the old F16 path and far more accurate. End-to-end on the real store it is dominated on **both** axes: 0.191 ms/query at 0.9973 recall, vs 0.125–0.142 at 0.9987 for the Metal default. The probe wasn't competing for the same CPU cores the real call path is. It stays reachable via `EL_GGML_DEVICE=accel` as a no-Metal fallback, labelled measured-and-rejected. ## Cold start — not a bug, and not ours The ~7.8s first-call cost attributed to this file is **not** this file re-initialising per call. Init was already cached behind `g_init_attempted` (and always was — the hypothesis that it re-ran `ggml_backend_load_all_from_path` every call is false; see the original lines 109–111). It is Apple's Metal shader cache missing on ggml's embedded metallib (~650 kernels in one `__ggml_metallib` section). That cache is keyed on **the library, not our binary, and is shared across processes**: ``` first-ever load : ggml_metal_library_init: loaded in 7.670 sec next run, different binary, same ggml : ggml_metal_library_init: loaded in 0.009 sec ``` Corroborated by the cache itself — `~/…/C/com.apple.metal/32024/libraries3.data` grew to 48MB at the moment of the 7.67s run. So: **once per machine per ggml version, not once per process**, and nothing this file does can avoid it. The hand-rolled strategy escapes it only because its shader is two small kernels instead of six hundred. Warm ggml init is 44–53ms, against 36–117ms for the hand-rolled strategy's device+pipeline setup — comparable. Loading only `libggml-metal.so` instead of every plugin in the directory is kept for tidiness and **explicitly documented as not a speedup**: 44.7/51.2/52.4ms against 46.9/55.1/58.9ms, the same number inside noise. ## Contract The `-2.0` sentinel contract is unchanged and re-verified at batch sizes straddling the chunk boundary (1, 7, 8, 9, 16, 17, 33), plus NULL rows, dim mismatches, zero-norm rows, and an all-invalid population — all pass, under the default and under `EL_GGML_MULMAT_CHUNK=1/300` and `EL_GGML_DEVICE=accel/cpu`. Independent confirmation of the defect: the **old** `ne11=300` path *fails* that same check at a 2e-6 cosine tolerance with **2299 mismatches**, while the new default passes with 0. ## Scope One file: `lang/runtime/eg_cosine_batch_strategy_ggml.c`. `eg_cosine_batch.metal` and `eg_cosine_batch_strategy_metal_hand.m` are untouched — read only as reference.

will.anderson added 1 commit 2026-08-15 22:58:04 +00:00

engram: make the ggml batch-cosine strategy actually compute in fp32 ... c008b7228a

#116 shipped the ggml strategy at 0.9933 id-recall against the CPU oracle
while the hand-rolled Metal kernel it replaced scored 0.9997 — a ~150x worse
error margin. That was not an inherent property of ggml. It was a usage bug in
this file, and this commit fixes it.

ggml-metal has two F32xF32 matmul kernels and picks between them purely on
ne11, the number of B rows, which for us is the query-batch size:

  ne11 <= 8  -> kernel_mul_mv_ext_f32_f32_* / kernel_mul_mv_f32_f32_*,
                templated <float, float> — genuine F32.
  ne11 >  8  -> kernel_mul_mm_f32_f32, templated
                <half, half4x4, simdgroup_half8x8, half, half2x4, ...> —
                BOTH operands narrowed to F16, despite F32 tensors on both
                sides.

The old code issued one ggml_mul_mat with ne11 = nq (300 in the benchmark),
landing squarely on the F16 path. The file's own header comment asserted the
opposite ("computes in F32 on the Metal backend"); that claim was wrong and is
replaced with the measurement.

Fix: emit ceil(nq/8) mul_mats over ne11<=8 ggml_view_2d slices of one query
tensor, all expanded into ONE graph and one ggml_backend_graph_compute, so the
node matrix is still uploaded and shared exactly once. EL_GGML_MULMAT_CHUNK
overrides the 8; setting it >= nq reproduces the old behaviour exactly, which
is also how the before/after below was measured in a single binary.

Measured, real store snapshot, 13415 live embedded nodes, dim=768, 300 real
queries, vs the CPU double-accumulated oracle (vindex_bench, offline copy of
the store — no live service touched):

  id-recall   same-rank |Δdist| max   mean
  old (ne11=300)   0.9933   6.80e-05   1.43e-05
  new (ne11<=8)    0.9987   4.77e-07   9.30e-08
  hand-rolled      0.9997   3.58e-07   7.55e-08

~145x better max error, ~154x better mean — now the same order of magnitude as
the hand-rolled kernel rather than 150x off it.

The cost is real and is documented rather than buried. Median of 15 reps of
the whole batch_multi() call, three runs: 13.2-14.4ms unchunked, 19.9-20.2ms
chunked, 17.7-18.0ms hand-rolled. Correctness costs ~+6.7ms per 300-query
batch and leaves ggml ~12% behind the hand-rolled kernel instead of ~35%
ahead. It cannot be recovered inside ggml: an fp32 matmul on Metal must
re-stream the node matrix once per <=8 queries, and ggml's Metal backend ships
no fp32 TILED matmul, so "fast" and "fp32" are genuinely exclusive there.

Two things that did NOT work, recorded so nobody retries them:

  - ggml_mul_mat_set_prec(t, GGML_PREC_F32) does nothing here. Error was
    bit-identical with and without it (1.038e-05 either way) — ggml-metal has
    no F32-accumulating mul_mm kernel to switch to. ne11 is the only lever.
  - The ACCEL/BLAS device looked excellent in an isolated compute-only probe
    (3.4-4.0ms, mean |Δdot| 1.5e-08) but is dominated on BOTH axes end-to-end
    (0.191 ms/query at 0.9973 recall vs 0.125-0.142 at 0.9987), because the
    probe was not competing for the same CPU cores the real call path is. It
    stays reachable via EL_GGML_DEVICE as a no-Metal fallback, labelled as
    measured-and-rejected, not as a recommendation.

Also corrected: the ~7.8s "cold start" blamed on this file is not this file
re-initialising per call — init was already cached. It is Apple's shader cache
missing on ggml's embedded metallib (~650 kernels), keyed on the library and
shared across processes: the first load on a machine reports
"loaded in 7.670 sec", the next run of a *different* binary reports 0.009 sec.
Once per machine per ggml version, not once per process, and not ours to fix.
Warm ggml init is 44-53ms vs 36-117ms for the hand-rolled strategy.

Loading only libggml-metal.so instead of every plugin in the directory is kept
for tidiness, and explicitly documented as NOT a speedup: 44.7-52.4ms against
46.9-58.9ms, the same number inside noise.

The -2.0 sentinel contract is unchanged and re-verified at batch sizes that
straddle the chunk boundary (1,7,8,9,16,17,33), plus NULL rows, dim
mismatches, zero-norm rows, and an all-invalid population. Notably the old
ne11=300 path fails that same check at a 2e-6 cosine tolerance with 2299
mismatches, which is an independent confirmation of the defect.

will.anderson merged commit 1c9de03fdb into feat/engram-ggml-cosine-batch 2026-08-15 23:01:31 +00:00

Sign in to join this conversation.

Reviewers

No Reviewers

Labels

Clear labels

BETA-CRITICAL

On the beta critical path — must resolve before ship

HELD

Blocked pending an external decision (counsel / dual sign-off)

ORTHOGONAL

Parallel workstream (testing/hardening) — safe to assign off critical path

POST-BETA

Important but not blocking the beta ship

No labels

Milestone

No items

No Milestone

Projects

Clear projects

No projects

Assignees

Clear assignees

andre.almeida (Andre Almeida) tim.lingo (Tim Lingo) will.anderson (Will Anderson)

No Assignees

1 Participants

Notifications

Subscribe

Due Date

No due date set.

Dependencies

No dependencies set.

Reference: neuron-technologies/el#121