Factory over ggml (supersedes #114) #116

Merged

will.anderson merged 3 commits from feat/engram-ggml-cosine-batch into dev

2026-08-15 23:22:57 +00:00

Author	SHA1	Message	Date
will.anderson	1c9de03fdb	Merge pull request 'engram: make the ggml batch-cosine strategy actually compute in fp32 (recall 0.9933 -> 0.9987)' (#121 ) from improve/ggml-cosine-fp32-and-init into feat/engram-ggml-cosine-batch El SDK CI - dev / build-and-test (pull_request) Failing after 3m43s Details	2026-08-15 23:01:29 +00:00
bigmerge	c008b7228a	engram: make the ggml batch-cosine strategy actually compute in fp32 #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.	2026-08-15 17:57:09 -05:00
bigmerge	b3f410fc91	engram: batch-cosine Adapter/Strategy/Factory over ggml, supersedes hand-rolled PR #114 El SDK CI - dev / build-and-test (pull_request) Failing after 4m29s Details Stop hand-rolling GPU kernels for batch cosine similarity — use ggml (the MIT-licensed compute library underneath llama.cpp, installed standalone via Homebrew) as the preferred backend, without ripping out PR #114's carefully-verified hand-rolled Metal shader. Structure: one stable public adapter (eg_cosine_batch.h, zero #ifdef at call sites) backed by three selectable concrete Strategies behind an internal vtable (eg_cosine_batch_strategy.h) chosen by a Factory (eg_cosine_batch.c): - eg_cosine_batch_strategy_ggml.c — NEW. ggml + dynamically-loaded Metal backend plugin (ggml_backend_load_all_from_path + ggml_mul_mat for the batched dot product), gather/scatter around the -2.0 sentinel contract. - eg_cosine_batch_strategy_metal_hand.m — PR #114's original hand-rolled Metal shader bridge, preserved almost verbatim, now one strategy among several rather than the only option. eg_cosine_batch.metal kept byte-identical to the original. - eg_cosine_batch_strategy_cpu.c — universal always-false fallback (direct descendant of PR #114's eg_metal_cosine_stub.c). Selection: EL_COSINE_BATCH_STRATEGY=ggml\|metal\|cpu\|auto (default: ggml first, then hand-rolled Metal, then CPU — first available wins), plus back-compat EL_METAL_COSINE=0 to disable every GPU-backed strategy. build_vindex_bench.sh compiles all three strategies on Darwin, CPU-fallback-only elsewhere. vindex_bench.c now reports BRUTE-GGML and BRUTE-METAL side by side against the same CPU oracle, on the same dataset, in one run (real numbers vs. real store snapshot in the PR body).	2026-08-15 17:16:41 -05:00

Author

SHA1

Message

Date

will.anderson

1c9de03fdb

Merge pull request 'engram: make the ggml batch-cosine strategy actually compute in fp32 (recall 0.9933 -> 0.9987)' (#121 ) from improve/ggml-cosine-fp32-and-init into feat/engram-ggml-cosine-batch

El SDK CI - dev / build-and-test (pull_request) Failing after 3m43s

Details

2026-08-15 23:01:29 +00:00

bigmerge

c008b7228a

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

#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.

2026-08-15 17:57:09 -05:00

bigmerge

b3f410fc91

engram: batch-cosine Adapter/Strategy/Factory over ggml, supersedes hand-rolled PR #114

El SDK CI - dev / build-and-test (pull_request) Failing after 4m29s

Details

Stop hand-rolling GPU kernels for batch cosine similarity — use ggml (the
MIT-licensed compute library underneath llama.cpp, installed standalone via
Homebrew) as the preferred backend, without ripping out PR #114's
carefully-verified hand-rolled Metal shader.

Structure: one stable public adapter (eg_cosine_batch.h, zero #ifdef at call
sites) backed by three selectable concrete Strategies behind an internal
vtable (eg_cosine_batch_strategy.h) chosen by a Factory (eg_cosine_batch.c):

  - eg_cosine_batch_strategy_ggml.c    — NEW. ggml + dynamically-loaded Metal
                                          backend plugin (ggml_backend_load_all_from_path
                                          + ggml_mul_mat for the batched dot
                                          product), gather/scatter around the
                                          -2.0 sentinel contract.
  - eg_cosine_batch_strategy_metal_hand.m — PR #114's original hand-rolled
                                          Metal shader bridge, preserved
                                          almost verbatim, now one strategy
                                          among several rather than the only
                                          option. eg_cosine_batch.metal kept
                                          byte-identical to the original.
  - eg_cosine_batch_strategy_cpu.c     — universal always-false fallback
                                          (direct descendant of PR #114's
                                          eg_metal_cosine_stub.c).

Selection: EL_COSINE_BATCH_STRATEGY=ggml|metal|cpu|auto (default: ggml first,
then hand-rolled Metal, then CPU — first available wins), plus back-compat
EL_METAL_COSINE=0 to disable every GPU-backed strategy. build_vindex_bench.sh
compiles all three strategies on Darwin, CPU-fallback-only elsewhere.

vindex_bench.c now reports BRUTE-GGML and BRUTE-METAL side by side against
the same CPU oracle, on the same dataset, in one run (real numbers vs. real
store snapshot in the PR body).

2026-08-15 17:16:41 -05:00

engram: batch-cosine Adapter/Strategy/Factory over ggml (supersedes #114) #116

3 Commits