engram: batch-cosine Adapter/Strategy/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

Conversation 0 Commits 3 Files Changed 9

+1581 -5

will.anderson commented 2026-08-15 22:20:39 +00:00

Owner

Copy Link

Copy Source

Why

Directive from Will: stop hand-rolling GPU compute kernels — use a real, proven, permissively-licensed library instead. ggml (MIT, the compute library underneath llama.cpp) is installed on this machine as a standalone Homebrew package (brew info ggml, v0.20.0), independent of llama.cpp itself: ggml.h / ggml-backend.h / ggml-metal.h / ggml-cpu.h / ggml-alloc.h under /opt/homebrew/include, libggml.dylib / libggml-base.dylib under /opt/homebrew/lib.

This is not a rip-and-replace of #114's hand-rolled Metal shader. #114 is real, carefully verified work — its numbers (17.2x, 0.9990 id-recall) are correct and its Metal shader is preserved, unchanged, as one of three selectable strategies behind a new Adapter/Strategy/Factory. This supersedes #114 (which is being closed with a pointer here) rather than sitting alongside it as a second, competing implementation.

ggml is used here strictly as a bounded compute utility — batched cosine-similarity math — analogous to a VBD Accessor calling out to infrastructure. It does not touch the engram's actual reasoning/graph-traversal/activation-spreading core, which stays 100% own-code.

The real ggml API shape (verified against the installed headers + a standalone probe, not assumed)

ggml ships its CPU and Metal implementations as dynamically loaded plugin .so files, not statically linkable symbols:

nm -gU libggml.dylib libggml-base.dylib   # ggml_backend_metal_init NOT exported
nm -gU $(brew --prefix ggml)/libexec/libggml-metal.so   # exported ONLY here

Correct usage (verified end-to-end in a standalone probe against a plain-C dot-product reference — bit-correct within float rounding):

1. ggml_backend_load_all_from_path(dir) — dlopen()s every backend plugin .so in dir and registers its device(s). We point this at $(brew --prefix ggml)/libexec (resolved via the stable /opt/homebrew/opt/ggml/libexec symlink, overridable via EL_GGML_BACKEND_PATH), falling back to ggml_backend_load_all()'s own default search only if that finds nothing — avoids double-registering the same plugins.
2. ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_GPU) → the registered Metal device.
3. ggml_backend_dev_init(dev, NULL) → a live ggml_backend_t.
4. Build a tiny ggml_context (no_alloc=true), declare 2D F32 tensors, ggml_mul_mat(node_matrix[dim,n], query_matrix[dim,nq]) → out[n,nq] where out[j*n+i] = dot(node_i, query_j) — ggml's documented convention (A: [k,n], B: [k,m] transposed internally, result [n,m]) maps exactly onto the row-major packed layout #114's kernel already used. One matmul replaces the whole per-row dot-product loop.
5. ggml_backend_alloc_ctx_tensors(ctx, backend) to allocate device buffers, ggml_backend_tensor_set()/_get() to upload/read back, ggml_backend_graph_compute() to run.

ggml_mul_mat only computes raw dot products — no notion of "cosine" or this codebase's -2.0 sentinel. Per spec: only valid, uniform-dim rows are gathered into the packed matrix sent to the GPU; every excluded row (null/dim-mismatch) gets -2.0 scattered back without ever reaching the GPU. Norms are computed on the CPU host in the same pass that already touches every element to gather — free — using the same 4-way partial-sum accumulation #114's kernel and the CPU oracle both use, so the float32 error profile stays comparable across strategies.

Structure: Adapter + Strategy + Factory

eg_cosine_batch.h                      — PUBLIC adapter. Zero #ifdef at call sites.
                                          Renamed from eg_metal_cosine.h since multiple
                                          backends now live under it — same names
                                          (eg_cosine_batch / _available / _multi) minus
                                          the "_metal" that no longer describes what's
                                          underneath. Contract is byte-identical to
                                          #114's original: false/untouched on ANY
                                          failure, never partial, caller always falls
                                          back to its own scalar loop.
eg_cosine_batch_strategy.h              — INTERNAL vtable (EgCosineBatchStrategy:
                                           name/available/batch/batch_multi). Only the
                                           factory + strategy .c/.m files see this.
eg_cosine_batch.c                       — the FACTORY. Selects one strategy, lazily,
                                           cached after first call.
eg_cosine_batch_strategy_ggml.c         — NEW. ggml + dynamic Metal backend plugin.
eg_cosine_batch_strategy_metal_hand.m   — #114's ORIGINAL hand-rolled Metal bridge,
                                           preserved almost verbatim (functions renamed
                                           static, exposed via the vtable getter).
eg_cosine_batch.metal                   — #114's shader, byte-identical, unchanged.
eg_cosine_batch_strategy_cpu.c          — universal always-false fallback (direct
                                           descendant of #114's eg_metal_cosine_stub.c).

Selection — env var + build-time + runtime capability probe, all three:

• Build-time: Darwin links all three strategies (-DEG_HAVE_STRATEGY_GGML -DEG_HAVE_STRATEGY_METAL_HAND); non-Darwin links only the CPU fallback — no Objective-C compiler, no Metal frameworks, matching #114's original Linux behavior exactly (verified: compiles clean, eg_cosine_batch() returns false unconditionally, -Wall -Wextra silent).
• Runtime probe: each strategy's own available() does the real, cheap-after-first-call check (device present, plugin loaded, pipeline/graph compiles) — never assumed from build-time alone.
• Env var: EL_COSINE_BATCH_STRATEGY=ggml|metal|cpu|auto forces a specific strategy (verified all four paths); EL_METAL_COSINE=0 (back-compat with #114) disables every GPU-backed strategy outright.
• Default (auto, unset): ggml → metal-hand → cpu, first available wins. ggml is what actually runs by default on this machine today — verified via eg_cosine_batch_strategy_name() — which is what makes "stop hand-rolling GPU kernels" real rather than nominal.

Any other genuine batch-numeric call site would route through this same eg_cosine_batch.h — see the sweep below for why none currently qualifies.

Real numbers

Real store snapshot (~/.neuron/backups/pre-cutover-hnsw-20260815-000735/neuron.egm, a file-based backup — never touched live :8742/:7770), 13,415 real embedded nodes, dim=768, 300 real queries, via vindex_bench store:

BRUTE-FORCE (CPU):    1.980 ms/query
BRUTE-GGML:           0.109 ms/query   (18.1x; id-recall 0.9933 vs CPU oracle; same-rank |Δdist| max 6.80e-05, mean 1.43e-05, n=2647)
BRUTE-METAL (hand):   0.108 ms/query   (18.4x; id-recall 0.9997 vs CPU oracle; same-rank |Δdist| max 3.58e-07, mean 7.55e-08, n=2930)

Synthetic scaling sweep (vindex_bench sweep 768 13000,25000,50000, dim=768, 150 queries/step):

N	CPU ms/q	ggml ms/q (speedup, recall)	hand-Metal ms/q (speedup, recall)
13,000	1.927	0.146 (13.2x, 1.0000)	0.179 (10.8x, 1.0000)
25,000	3.724	0.392 (9.5x, 1.0000)	0.297 (12.6x, 1.0000)

(50,000-node step was still building its HNSW index — ~200s+ for the index alone at this size — when the harness's own timeout cut it off; the trend across 13k→25k already shows both GPU strategies holding a stable ~10-18x margin over CPU as N grows, consistent with #114's original 13k→50k finding.)

Honest tradeoff: ggml vs. hand-rolled Metal

Throughput is statistically tied — sometimes ggml edges ahead (13k synth: 13.2x vs 10.8x), sometimes hand-rolled does (real store: 18.4x vs 18.1x; 25k synth: 12.6x vs 9.5x). Neither strategy is a clear throughput winner at these scales.

Precision is NOT tied. ggml's -2.0-sentineled cosine consistently shows ~150-190x larger same-rank distance error than the hand-rolled kernel against the CPU double-precision oracle (e.g. real-store: max 6.80e-05 vs 3.58e-07), and correspondingly slightly lower id-recall (0.9933 vs 0.9997 over 300 queries — a handful of near-tied ranks reorder). Both strategies compute in float32 throughout; the gap is most likely ggml's Metal matmul kernel using a different (more parallel, SIMD-group-reduction) accumulation order than the hand-rolled kernel's explicit sequential 4-way unroll — a real, measured difference, not assumed. In absolute terms 6.8e-05 is still tiny for 768-dim float32 cosine and >99% recall is still strong — but it is a genuine, honest tradeoff, not a wash.

Cold-start cost is also not tied. ggml's Metal library init took ~7.8s on a cold cache (first-ever call in the process — it loads/parses its full embedded kernel library, covering every op/type/config combination) vs. ~20ms warm on a subsequent process. The hand-rolled kernel only ever compiles the 2 tiny kernels it actually uses, at "tens of ms." Both costs are one-time-per-process and amortize over a long-lived daemon, but ggml's is materially (~150-400x) larger.

Per Will's directive, ggml remains the default — the point was "stop hand-rolling GPU kernels," not "prove ggml strictly dominates" — but this tradeoff should be visible to reviewers, not buried.

Directive 3 sweep — what else was checked

Re-verified #114's own investigation by reading the current dev code directly (not re-derived from scratch): the 5 eg_cosine( call sites in el_runtime.c — CPU reference impl, O(K=8) dedup (not O(N²)), the lazy/memoized eg_cosq_at (M8.1 fix — batching this would re-introduce the O(N·D) prescan the codebase deliberately killed), eg_knn_for_node's ~16-candidate HNSW loop (too small for dispatch overhead), and a single pairwise API function — none qualify, confirming #114's read.

New candidates investigated this PR, both rejected with evidence:

1. ingest.el's find_existing_by_content dedup (per a mid-flight ask) — traced its call chain fully: /api/search → engram_retrieve_geometric_json (el_runtime.c ~L12609). Read the function in full: it is purely lexical/tag-based (token seed → skill:/shape:/op: region addressing → bounded graph spread) — its own comment says "no cold-embed seed." eg_similarity() exists in ingest.el but is dead code, never called. Not embedding-shaped at all currently — does not qualify.
2. elp/ faculty code + projector pipeline — swept every file (audio-surface.el, image-surface.el, speech.el, voice-profile.el, voice-ingest.el, speech-ingest.el, elp/projector/*.py). Found exactly two genuine embarrassingly-parallel batch-numeric loops, both in speech.el: voice_f0 (pitch-lag autocorrelation, ~53-213 lag-candidates × ~1500-3000-sample window) and voice_peak_in_band (per-frequency-bin Goertzel DFT, ~28 bins × ≤3000 samples). Same "1-query-vs-N-candidates" shape as batch-cosine — but both run only on sub-second sustained-vowel clips today, too small for GPU dispatch overhead to pay off, same reasoning #114 used to correctly reject HNSW's bounded candidate lists. Everything else in elp/ is sequential-state synthesis, tiny fixed loops, or small-N string/keyword scoring.

Investigated, evidence-backed, deliberately left OUT of this PR (per explicit instruction not to force it in):

3. Embedding generation itself (eg_embed_fetch, el_runtime.c ~L6829) — currently one text → one HTTP POST to Ollama's /api/embeddings per call, looped sequentially in engram_embed_backfill (up to 64 nodes/call). Real feasibility check: llama.cpp's llama-embedding CLI loads the Ollama-managed nomic-embed-text blob directly (~/.ollama/models/blobs/sha256-970aa7... — confirmed to be a raw GGUF file despite the missing extension) via ggml/Metal, and with --embd-normalize -1 (raw output) produces embeddings cosine 0.9999994 identical to Ollama's own HTTP response for the same text — i.e. numerically compatible with the 13k+ already-stored Ollama-produced embeddings. This is real and promising, but substantially bigger than this PR's scope: it needs batched tokenization via llama.cpp's C API, a persistently-loaded model inside the daemon process, and exact truncation/pooling parity with the current code. Reporting as a strong, evidence-backed follow-on, not half-building it here.

Verification performed

• -Wall -Wextra clean on every new file (ggml strategy, factory, CPU strategy, ObjC hand-rolled strategy).
• Non-Darwin build path simulated directly (compiled + linked with zero ggml/Metal deps, ran vindex_bench synth successfully, both GPU strategies correctly reported "not compiled into this build", eg_cosine_batch() returns false unconditionally).
• Public factory API (eg_cosine_batch.h) exercised standalone exactly as a real call site would: default strategy resolves to ggml, -2.0 sentinel verified through the full path for a dim-mismatched node, and all four EL_COSINE_BATCH_STRATEGY values (ggml/metal/cpu/unset) plus the legacy EL_METAL_COSINE=0 gate all resolve to the correct strategy.
• vindex_bench's hand-rolled-Metal path re-verified still works unmodified through the new strategy wrapper (id-recall 0.9997-1.0000 across every run above).

Blocked / not attempted

Nothing structural. ggml's dynamic-plugin backend loading works cleanly from a plain cc-built binary once pointed at the right libexec directory — no build-system blocker.

## Why Directive from Will: stop hand-rolling GPU compute kernels — use a real, proven, permissively-licensed library instead. `ggml` (MIT, the compute library underneath `llama.cpp`) is installed on this machine as a standalone Homebrew package (`brew info ggml`, v0.20.0), independent of `llama.cpp` itself: `ggml.h` / `ggml-backend.h` / `ggml-metal.h` / `ggml-cpu.h` / `ggml-alloc.h` under `/opt/homebrew/include`, `libggml.dylib` / `libggml-base.dylib` under `/opt/homebrew/lib`. This is **not** a rip-and-replace of #114's hand-rolled Metal shader. #114 is real, carefully verified work — its numbers (17.2x, 0.9990 id-recall) are correct and its Metal shader is preserved, unchanged, as one of three selectable strategies behind a new Adapter/Strategy/Factory. This supersedes #114 (which is being closed with a pointer here) rather than sitting alongside it as a second, competing implementation. `ggml` is used here strictly as a bounded **compute utility** — batched cosine-similarity math — analogous to a VBD Accessor calling out to infrastructure. It does not touch the engram's actual reasoning/graph-traversal/activation-spreading core, which stays 100% own-code. ## The real ggml API shape (verified against the installed headers + a standalone probe, not assumed) `ggml` ships its CPU and Metal implementations as **dynamically loaded plugin `.so` files**, not statically linkable symbols: ``` nm -gU libggml.dylib libggml-base.dylib # ggml_backend_metal_init NOT exported nm -gU $(brew --prefix ggml)/libexec/libggml-metal.so # exported ONLY here ``` Correct usage (verified end-to-end in a standalone probe against a plain-C dot-product reference — bit-correct within float rounding): 1. `ggml_backend_load_all_from_path(dir)` — `dlopen()`s every backend plugin `.so` in `dir` and registers its device(s). We point this at `$(brew --prefix ggml)/libexec` (resolved via the stable `/opt/homebrew/opt/ggml/libexec` symlink, overridable via `EL_GGML_BACKEND_PATH`), falling back to `ggml_backend_load_all()`'s own default search only if that finds nothing — avoids double-registering the same plugins. 2. `ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_GPU)` → the registered Metal device. 3. `ggml_backend_dev_init(dev, NULL)` → a live `ggml_backend_t`. 4. Build a tiny `ggml_context` (`no_alloc=true`), declare 2D F32 tensors, `ggml_mul_mat(node_matrix[dim,n], query_matrix[dim,nq])` → `out[n,nq]` where `out[j*n+i] = dot(node_i, query_j)` — ggml's documented convention (`A: [k,n]`, `B: [k,m]` transposed internally, result `[n,m]`) maps exactly onto the row-major packed layout #114's kernel already used. One matmul replaces the whole per-row dot-product loop. 5. `ggml_backend_alloc_ctx_tensors(ctx, backend)` to allocate device buffers, `ggml_backend_tensor_set()`/`_get()` to upload/read back, `ggml_backend_graph_compute()` to run. `ggml_mul_mat` only computes raw dot products — no notion of "cosine" or this codebase's `-2.0` sentinel. Per spec: only valid, uniform-dim rows are gathered into the packed matrix sent to the GPU; every excluded row (null/dim-mismatch) gets `-2.0` scattered back without ever reaching the GPU. Norms are computed on the CPU host in the same pass that already touches every element to gather — free — using the same 4-way partial-sum accumulation #114's kernel and the CPU oracle both use, so the float32 error profile stays comparable across strategies. ## Structure: Adapter + Strategy + Factory ``` eg_cosine_batch.h — PUBLIC adapter. Zero #ifdef at call sites. Renamed from eg_metal_cosine.h since multiple backends now live under it — same names (eg_cosine_batch / _available / _multi) minus the "_metal" that no longer describes what's underneath. Contract is byte-identical to #114's original: false/untouched on ANY failure, never partial, caller always falls back to its own scalar loop. eg_cosine_batch_strategy.h — INTERNAL vtable (EgCosineBatchStrategy: name/available/batch/batch_multi). Only the factory + strategy .c/.m files see this. eg_cosine_batch.c — the FACTORY. Selects one strategy, lazily, cached after first call. eg_cosine_batch_strategy_ggml.c — NEW. ggml + dynamic Metal backend plugin. eg_cosine_batch_strategy_metal_hand.m — #114's ORIGINAL hand-rolled Metal bridge, preserved almost verbatim (functions renamed static, exposed via the vtable getter). eg_cosine_batch.metal — #114's shader, byte-identical, unchanged. eg_cosine_batch_strategy_cpu.c — universal always-false fallback (direct descendant of #114's eg_metal_cosine_stub.c). ``` **Selection** — env var + build-time + runtime capability probe, all three: - **Build-time**: Darwin links all three strategies (`-DEG_HAVE_STRATEGY_GGML -DEG_HAVE_STRATEGY_METAL_HAND`); non-Darwin links only the CPU fallback — no Objective-C compiler, no Metal frameworks, matching #114's original Linux behavior exactly (verified: compiles clean, `eg_cosine_batch()` returns `false` unconditionally, `-Wall -Wextra` silent). - **Runtime probe**: each strategy's own `available()` does the real, cheap-after-first-call check (device present, plugin loaded, pipeline/graph compiles) — never assumed from build-time alone. - **Env var**: `EL_COSINE_BATCH_STRATEGY=ggml|metal|cpu|auto` forces a specific strategy (verified all four paths); `EL_METAL_COSINE=0` (back-compat with #114) disables every GPU-backed strategy outright. - **Default (`auto`, unset)**: `ggml` → `metal-hand` → `cpu`, first available wins. **ggml is what actually runs by default on this machine today** — verified via `eg_cosine_batch_strategy_name()` — which is what makes "stop hand-rolling GPU kernels" real rather than nominal. Any other genuine batch-numeric call site would route through this same `eg_cosine_batch.h` — see the sweep below for why none currently qualifies. ## Real numbers **Real store snapshot** (`~/.neuron/backups/pre-cutover-hnsw-20260815-000735/neuron.egm`, a file-based backup — never touched live `:8742`/`:7770`), 13,415 real embedded nodes, dim=768, 300 real queries, via `vindex_bench store`: ``` BRUTE-FORCE (CPU): 1.980 ms/query BRUTE-GGML: 0.109 ms/query (18.1x; id-recall 0.9933 vs CPU oracle; same-rank |Δdist| max 6.80e-05, mean 1.43e-05, n=2647) BRUTE-METAL (hand): 0.108 ms/query (18.4x; id-recall 0.9997 vs CPU oracle; same-rank |Δdist| max 3.58e-07, mean 7.55e-08, n=2930) ``` **Synthetic scaling sweep** (`vindex_bench sweep 768 13000,25000,50000`, dim=768, 150 queries/step): | N | CPU ms/q | ggml ms/q (speedup, recall) | hand-Metal ms/q (speedup, recall) | |---|---|---|---| | 13,000 | 1.927 | 0.146 (13.2x, 1.0000) | 0.179 (10.8x, 1.0000) | | 25,000 | 3.724 | 0.392 (9.5x, 1.0000) | 0.297 (12.6x, 1.0000) | (50,000-node step was still building its HNSW index — ~200s+ for the index alone at this size — when the harness's own timeout cut it off; the trend across 13k→25k already shows both GPU strategies holding a stable ~10-18x margin over CPU as N grows, consistent with #114's original 13k→50k finding.) ## Honest tradeoff: ggml vs. hand-rolled Metal **Throughput is statistically tied** — sometimes ggml edges ahead (13k synth: 13.2x vs 10.8x), sometimes hand-rolled does (real store: 18.4x vs 18.1x; 25k synth: 12.6x vs 9.5x). Neither strategy is a clear throughput winner at these scales. **Precision is NOT tied.** ggml's `-2.0`-sentineled cosine consistently shows **~150-190x larger** same-rank distance error than the hand-rolled kernel against the CPU double-precision oracle (e.g. real-store: max `6.80e-05` vs `3.58e-07`), and correspondingly slightly lower id-recall (0.9933 vs 0.9997 over 300 queries — a handful of near-tied ranks reorder). Both strategies compute in float32 throughout; the gap is most likely ggml's Metal matmul kernel using a different (more parallel, SIMD-group-reduction) accumulation order than the hand-rolled kernel's explicit sequential 4-way unroll — a real, measured difference, not assumed. In absolute terms `6.8e-05` is still tiny for 768-dim float32 cosine and >99% recall is still strong — but it is a genuine, honest tradeoff, not a wash. **Cold-start cost is also not tied.** ggml's Metal library init took **~7.8s on a cold cache** (first-ever call in the process — it loads/parses its full embedded kernel library, covering every op/type/config combination) vs. `~20ms` warm on a subsequent process. The hand-rolled kernel only ever compiles the 2 tiny kernels it actually uses, at "tens of ms." Both costs are one-time-per-process and amortize over a long-lived daemon, but ggml's is materially (~150-400x) larger. Per Will's directive, `ggml` remains the default — the point was "stop hand-rolling GPU kernels," not "prove ggml strictly dominates" — but this tradeoff should be visible to reviewers, not buried. ## Directive 3 sweep — what else was checked Re-verified #114's own investigation by reading the current `dev` code directly (not re-derived from scratch): the 5 `eg_cosine(` call sites in `el_runtime.c` — CPU reference impl, O(K=8) dedup (not O(N²)), the lazy/memoized `eg_cosq_at` (M8.1 fix — batching this would re-introduce the O(N·D) prescan the codebase deliberately killed), `eg_knn_for_node`'s ~16-candidate HNSW loop (too small for dispatch overhead), and a single pairwise API function — none qualify, confirming #114's read. **New candidates investigated this PR, both rejected with evidence:** 1. **`ingest.el`'s `find_existing_by_content` dedup** (per a mid-flight ask) — traced its call chain fully: `/api/search` → `engram_retrieve_geometric_json` (`el_runtime.c` ~L12609). Read the function in full: it is purely lexical/tag-based (token seed → `skill:`/`shape:`/`op:` region addressing → bounded graph spread) — its own comment says "no cold-embed seed." `eg_similarity()` exists in `ingest.el` but is dead code, never called. **Not embedding-shaped at all currently** — does not qualify. 2. **`elp/` faculty code + projector pipeline** — swept every file (`audio-surface.el`, `image-surface.el`, `speech.el`, `voice-profile.el`, `voice-ingest.el`, `speech-ingest.el`, `elp/projector/*.py`). Found exactly two genuine embarrassingly-parallel batch-numeric loops, both in `speech.el`: `voice_f0` (pitch-lag autocorrelation, ~53-213 lag-candidates × ~1500-3000-sample window) and `voice_peak_in_band` (per-frequency-bin Goertzel DFT, ~28 bins × ≤3000 samples). Same "1-query-vs-N-candidates" shape as batch-cosine — but both run only on sub-second sustained-vowel clips today, too small for GPU dispatch overhead to pay off, same reasoning #114 used to correctly reject HNSW's bounded candidate lists. Everything else in `elp/` is sequential-state synthesis, tiny fixed loops, or small-N string/keyword scoring. **Investigated, evidence-backed, deliberately left OUT of this PR** (per explicit instruction not to force it in): 3. **Embedding generation itself** (`eg_embed_fetch`, `el_runtime.c` ~L6829) — currently one text → one HTTP POST to Ollama's `/api/embeddings` per call, looped sequentially in `engram_embed_backfill` (up to 64 nodes/call). Real feasibility check: `llama.cpp`'s `llama-embedding` CLI loads the Ollama-managed `nomic-embed-text` blob **directly** (`~/.ollama/models/blobs/sha256-970aa7...` — confirmed to be a raw GGUF file despite the missing extension) via ggml/Metal, and with `--embd-normalize -1` (raw output) produces embeddings **cosine 0.9999994 identical** to Ollama's own HTTP response for the same text — i.e. numerically compatible with the 13k+ already-stored Ollama-produced embeddings. This is real and promising, but substantially bigger than this PR's scope: it needs batched tokenization via llama.cpp's C API, a persistently-loaded model inside the daemon process, and exact truncation/pooling parity with the current code. Reporting as a strong, evidence-backed follow-on, not half-building it here. ## Verification performed - `-Wall -Wextra` clean on every new file (ggml strategy, factory, CPU strategy, ObjC hand-rolled strategy). - Non-Darwin build path simulated directly (compiled + linked with zero ggml/Metal deps, ran `vindex_bench synth` successfully, both GPU strategies correctly reported "not compiled into this build", `eg_cosine_batch()` returns `false` unconditionally). - Public factory API (`eg_cosine_batch.h`) exercised standalone exactly as a real call site would: default strategy resolves to `ggml`, `-2.0` sentinel verified through the full path for a dim-mismatched node, and all four `EL_COSINE_BATCH_STRATEGY` values (`ggml`/`metal`/`cpu`/unset) plus the legacy `EL_METAL_COSINE=0` gate all resolve to the correct strategy. - `vindex_bench`'s hand-rolled-Metal path re-verified still works unmodified through the new strategy wrapper (id-recall 0.9997-1.0000 across every run above). ## Blocked / not attempted Nothing structural. `ggml`'s dynamic-plugin backend loading works cleanly from a plain `cc`-built binary once pointed at the right `libexec` directory — no build-system blocker.

will.anderson added 1 commit 2026-08-15 22:20:41 +00:00

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

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

will.anderson referenced this pull request 2026-08-15 22:21:02 +00:00

engram: real Metal batch-cosine kernel for vindex_bench's brute-force oracle #114

will.anderson referenced this issue from a commit 2026-08-15 22:57:17 +00:00

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

will.anderson referenced this pull request 2026-08-15 22:58:03 +00:00

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

will.anderson added 2 commits 2026-08-15 23:01:33 +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.

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 1c9de03fdb

will.anderson merged commit 38a8e32d6c into dev 2026-08-15 23:22:57 +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#116