Merge pull request 'promote dev -> stage: transduce unification + HNSW + ggml adapter + reconciliation (2026-08-15)' (#119) from dev into stage

transduce() is now THE single mechanism: one function, no content-type
branch inside it. It never asks whether `source` is prose, JSON, or
raw/opaque bytes (audio, etc.) — it runs one algorithm unconditionally:
split on "\n\n" as a universal boundary-marker check, and if that finds
no boundary, fall back to fixed 4096-char windows. Same node/edge wiring
(root -contains-> chunk, chunk -precedes-> next, "#"-prefixed chunk gets
a heading/section_of link) regardless of what's inside a chunk. Dedup is
the existing find_existing_by_content path via merge_manifold, applied
uniformly. The old transduce_structured JSON dataset/records/feature-node
interpretation is deleted outright, not just unused — a JSON file now
gets chunked and deduped like anything else, with no pre-computed
structure. All five ingest_* entry points still exist unchanged in name
and role; ingest_file/ingest_dir/ingest_url/ingest_llm now call the one
transduce() (ingest_stream builds its own turn-nodes directly and never
called either old function, so it's untouched).

This unlocks raw/opaque content (audio, or anything else with no natural
text/JSON shape) without any DSP, LLM call, or external API: transduce()
chunks it exactly like it chunks anything else. There is zero semantic
understanding of audio (or any payload) claimed or built here — any
meaning is expected to emerge later from Neuron's own existing mechanisms
(embedding, spreading activation, dedup) acting on this real geometry
over time.

Two small C builtins added to el_runtime.c/h (fs_size, fs_read_b64_chunk)
because El strings are NUL-unsafe under strlen-based ops and fs_read()'s
result silently truncates at the first embedded NUL, which is routine in
real binary/audio bytes. ingest_file compares fs_read()'s string length
against a real fs_size() stat() count; on mismatch it rebuilds the
payload as base64-encoded fixed 3072-byte windows read directly off disk
(binary-safe in C, verbatim, no invention), joined with the same "\n\n"
marker transduce()'s boundary scan already looks for. This is a
mechanical fidelity fix, not interpretation of content — transduce()
never learns a fallback happened. Registered both builtins' arity in
codegen.el; did not rebuild the elc compiler binary itself (unrelated,
pre-existing gap: self-hosting elc via el_seed.c fails on this worktree
independent of this change, reproduced with codegen.el reverted) — the
existing elc binary compiles calls to unregistered builtins via its
already-existing arity=-1 passthrough, confirmed by an actual clean
`elc ingest.el` + `cc` build against the modified el_runtime.c.

INGEST_KIND keeps existing only as an acquisition-mechanism selector
(dir/file/url/llm/stream — which RPC to use to fetch bytes), not as a
content-type flag; the redundant "structured" value (an alias for "file"
that hinted the now-deleted JSON branch) is removed. ingest_dir drops its
file-extension filter for the same reason: transduce() takes anything now.

Verification: local manifold construction confirmed correct against a
real captured audio file (will_clean.wav, 304288 bytes, and a 12288-byte
real prefix slice) — exact expected node/edge counts both times
(101 nodes/199 edges full file; 5 nodes/7 edges for the slice, matching
ceil(bytes/3072)+1 nodes and 2n-1 edges), with real, verbatim base64
content confirmed decoding back to the actual WAV header bytes. Compiles
clean via the real elc + the modified el_runtime.c/engram_*.c (built and
booted an actual sandbox engram off this exact source with `nsbx create
--branch`).

NOT verified this session, disclosed rather than papered over: end-to-end
server-confirmed persistence (a real before/after /api/stats delta, and a
fetched node by id) for the audio, prose, and JSON-fixture cases. Every
local nsbx sandbox engram tried tonight (two stock pre-#109 binaries
hitting the known O(N*D) brute-force scan bug, then a fresh #109/HNSW
binary built from current dev) took minutes-to indefinitely long on the
final /api/load-merge write's embedding step and hit the client's 60s
HTTP timeout before responding, even for a 5-node write. This is
confirmed as real (if slow) forward progress, not a hang: the sandbox's
WAL file was observed growing steadily across every attempt. The code's
own pre-existing HONESTY GATE correctly refused to report success in
every case, returning "load-merge failed: ..." with a
"nothing below this manifold was confirmed persisted by the server" note
instead — exactly as designed. This is an environment/infrastructure
limitation, not a defect introduced by this change: the engram server
binary itself is untouched by this commit.

ingest/.gitignore

@@ -0,0 +1 @@
+build/

ingest/src/ingest.el

@@ -1,17 +1,28 @@
 // ingest.el — the native EL AFFERENT INGEST ORGAN
 //
 // The source-polymorphic ingest(source) primitive: point it at a directory,
-// file, url, llm-query, structured-primitive set, or stream; it EXTRACTS the
-// real content faithfully (no invention), TRANSDUCES it into a DISCRETE
-// MANIFOLD (multiple nodes + internal edges — meaning-structure, never a
-// single blob; the conversion from extracted surface content into geometry
-// is automatic and invisible to the caller, the way digestion is invisible
-// to the one who chose to eat — ingest is the conscious act, transduce is
-// the mechanism underneath it, and it is no less real for being unseen),
-// and MERGES that manifold into the engram geometry: shared
-// meanings DEDUP onto existing nodes (search + exact/cosine match), genuinely
-// new meanings add nodes, relations add edges. Every node enters with
-// PROVENANCE + grounding-level + stewardship class from the moment of entry.
+// file, url, llm-query, or stream; it EXTRACTS the real content faithfully
+// (no invention), TRANSDUCES it into a DISCRETE MANIFOLD (multiple nodes +
+// internal edges — meaning-structure, never a single blob; the conversion
+// from extracted surface content into geometry is automatic and invisible
+// to the caller, the way digestion is invisible to the one who chose to
+// eat — ingest is the conscious act, transduce is the mechanism underneath
+// it, and it is no less real for being unseen), and MERGES that manifold
+// into the engram geometry: shared meanings DEDUP onto existing nodes
+// (search + exact/cosine match), genuinely new meanings add nodes,
+// relations add edges. Every node enters with PROVENANCE + grounding-level
+// + stewardship class from the moment of entry.
+//
+// transduce() is THE single mechanism — one function, polymorphic, with no
+// content-type branch inside it. It does not ask whether a payload is
+// prose, structured data, or raw/opaque bytes (audio, or anything else);
+// it runs one boundary-scan-with-fixed-window-fallback chunking algorithm
+// and one dedup mechanism on whatever bytes it's handed, unconditionally.
+// Any deeper structure a payload might have (shared fields, relationships,
+// what a chunk of audio "means") is NOT interpreted here — that's left
+// entirely to the engram's own mechanisms (embedding, spreading activation,
+// dedup) acting on this real geometry over time. This organ claims zero
+// semantic understanding of any payload it transduces.
 //
 // It is a pure HTTP CLIENT of the engram server — it links only el_runtime.c
 // via fs/http/json/string builtins; it never links el_seed.c or the engram
@@ -46,60 +57,6 @@ fn j_q(s: String) -> String {
     return "\"" + j_esc(s) + "\""
 }
 
-// Extract the top-level keys of a JSON object string. A thin, self-contained
-// scanner (FLAGGED: the one non-trivial parser in this organ — everything else
-// is faithful text handling). Tracks string state + brace/bracket depth; a key
-// is a string at object-interior depth 1 immediately followed by ':'.
-fn json_object_keys(obj: String) -> [String] {
-    let keys: [String] = el_list_empty()
-    let n: Int = str_len(obj)
-    let i: Int = 0
-    let depth: Int = 0
-    let in_str: Bool = false
-    let esc: Bool = false
-    let str_start: Int = -1
-    let cur: String = ""
-    let have_key: Bool = false
-    while i < n {
-        let c: String = str_char_at(obj, i)
-        if in_str {
-            if esc {
-                esc = false
-            } else {
-                if str_eq(c, "\\") {
-                    esc = true
-                } else {
-                    if str_eq(c, "\"") {
-                        in_str = false
-                        cur = str_slice(obj, str_start + 1, i)
-                        have_key = true
-                    }
-                }
-            }
-        } else {
-            if str_eq(c, "\"") {
-                in_str = true
-                str_start = i
-            }
-            if str_eq(c, "{") { depth = depth + 1 }
-            if str_eq(c, "}") { depth = depth - 1 }
-            if str_eq(c, "[") { depth = depth + 1 }
-            if str_eq(c, "]") { depth = depth - 1 }
-            if str_eq(c, ":") {
-                if have_key {
-                    if depth == 1 {
-                        keys = el_list_append(keys, cur)
-                    }
-                }
-                have_key = false
-            }
-            if str_eq(c, ",") { have_key = false }
-        }
-        i = i + 1
-    }
-    return keys
-}
-
 // ═══════════════════════════════════════════════════════════════════════════
 // SECTION B — engram HTTP client (provenance-carrying afferent LOAD)
 // ═══════════════════════════════════════════════════════════════════════════
@@ -402,168 +359,125 @@ fn head80(s: String) -> String {
 //             We accumulate into module-level lists carried by the caller.
 // ═══════════════════════════════════════════════════════════════════════════
 
-// PROSE: chunk text into a discrete manifold. Split on blank lines into
-// paragraphs; every non-empty paragraph is its own node (NEVER one blob).
-// Edges: doc-root -contains-> chunk; chunk -precedes-> next chunk;
-// most-recent-heading -section_of-> chunk. Content is verbatim (substring of
-// the source) — pure extraction of ground truth.
-fn transduce_prose(nodes: [String], edges: [String], text: String,
-               prov: String, ground: String, steward: String,
-               root_lid: String, root_title: String) -> [String] {
-    // returns [nodes_json_list_encoded, edges_json_list_encoded] is awkward in
-    // EL; instead we mutate by returning a 2-list. We package results as a
-    // single JSON array string carrying {nodes:[...],edges:[...]} additions.
-    // (Kept simple: caller passes empty lists and receives the packaged pair.)
+// TRANSDUCE — the single mechanism. Takes ANY payload (prose, structured
+// data, raw/opaque bytes — audio, whatever) as one opaque string and turns
+// it into a discrete manifold: nodes + internal edges. There is no
+// content-type branch anywhere in this function. It never asks "is this
+// text," "is this JSON," "is this audio" — it runs ONE algorithm on the
+// bytes it is given, unconditionally:
+//
+//   1. BOUNDARY SCAN — split on "\n\n". This is a property of the bytes
+//      (does a blank-line-style marker occur in them, yes or no), not a
+//      classification of what the content IS. Prose paragraphs split on it
+//      because that's how prose is typically written; that's a fact about
+//      the bytes, not a rule this function knows about prose. Anything else
+//      that happens to contain the same marker splits on it too, and
+//      anything that doesn't, doesn't — same code path either way.
+//   2. FIXED-WINDOW FALLBACK — if step 1 found no boundary (0 or 1 non-empty
+//      piece), the payload is cut into fixed-size windows instead. Same
+//      chunk-per-node, edge-per-adjacency structure as step 1 produces; only
+//      the source of the cut point differs.
+//
+// Every resulting chunk becomes its own node (never one blob), wired with
+// the same edges regardless of what's inside a chunk: root -contains->
+// chunk, chunk -precedes-> next chunk, and — if a chunk happens to start
+// with "#" — most-recent-heading -section_of-> chunk. That "#" check is a
+// structural marker (a fact about a chunk's first byte), not a decision
+// about whether this run is "the text case": chunks from any payload that
+// never happen to start with "#" simply never trigger it.
+//
+// Dedup is the existing, fully generic mechanism (find_existing_by_content,
+// via merge_manifold downstream of merge_packed) applied uniformly to every
+// chunk from every payload — there is no separate "structured" dedup path.
+// Any deeper structure that might exist inside a payload (shared fields,
+// repeated records, relationships) is NOT pre-computed here; that's left to
+// the engram's own mechanisms (embedding, spreading activation, dedup)
+// acting on this geometry over time, which is the whole point of handing it
+// raw bytes instead of a hand-coded interpretation of them.
+//
+// Byte-safety note: `source` must already be a string this function can
+// safely str_split/str_slice. Protecting it from silent truncation (El
+// strings are NUL-unsafe under strlen-based ops; fs_read()'s result
+// truncates at the first embedded NUL, which is routine in real binary
+// bytes) is a MECHANICAL fidelity concern that belongs to whatever produced
+// `source` (see ingest_file's file_source_string below) — not a
+// content-type judgment made in here. transduce() never learns whether a
+// chunk is plain text or a base64-encoded raw-byte window; every chunk is
+// handled identically either way.
+fn transduce(nodes: [String], edges: [String], source: String,
+             prov: String, ground: String, steward: String,
+             root_lid: String, root_title: String) -> [String] {
     let tagbase: String = "prov:" + prov + " ground:" + ground + " steward:" + steward
-    // root node
-    nodes = el_list_append(nodes, mk_node(root_lid, "document: " + root_title,
-        "Concept", "Semantic", "0.6", "0.6", "0.9", tagbase + " kind:document"))
+    nodes = el_list_append(nodes, mk_node(root_lid, "source: " + root_title,
+        "Concept", "Semantic", "0.6", "0.6", "0.9", tagbase + " kind:source"))
 
-    let paras: [String] = str_split(text, "\n\n")
-    let np: Int = el_list_len(paras)
-    let idx: Int = 0
+    // step 1: universal boundary scan
+    let boundary_parts: [String] = str_split(source, "\n\n")
+    let chunks: [String] = el_list_empty()
+    let bp_n: Int = el_list_len(boundary_parts)
+    let bp_i: Int = 0
+    while bp_i < bp_n {
+        let piece: String = str_trim(el_list_get(boundary_parts, bp_i))
+        if !str_eq(piece, "") { chunks = el_list_append(chunks, piece) }
+        bp_i = bp_i + 1
+    }
+
+    // step 2: no boundary found -> fixed-size windows over the whole
+    // payload. 4096 chars/node: low kilobytes — big enough to keep node
+    // count sane on a large unbroken payload, small enough that each node
+    // stays a legible, individually embeddable/dedupable unit rather than
+    // one giant blob.
+    if el_list_len(chunks) <= 1 {
+        chunks = el_list_empty()
+        let total: Int = str_len(source)
+        let win: Int = 4096
+        let off: Int = 0
+        while off < total {
+            let endp: Int = if off + win < total { off + win } else { total }
+            let piece: String = str_slice(source, off, endp)
+            if !str_eq(piece, "") { chunks = el_list_append(chunks, piece) }
+            off = off + win
+        }
+    }
+
+    let nc: Int = el_list_len(chunks)
+    let ci: Int = 0
     let last_chunk: String = ""
     let last_heading: String = ""
-    let ci: Int = 0
-    while idx < np {
-        let raw: String = str_trim(el_list_get(paras, idx))
-        if !str_eq(raw, "") {
-            let lid: String = root_lid + ":c" + int_to_str(ci)
-            let is_heading: Bool = str_starts_with(raw, "#")
-            let kind: String = if is_heading { "kind:heading" } else { "kind:doc-chunk" }
-            nodes = el_list_append(nodes, mk_node(lid, raw,
-                "Knowledge", "Semantic", "0.55", "0.55", "0.9", tagbase + " " + kind))
-            // containment: document root -contains-> chunk
-            edges = el_list_append(edges, mk_edge(root_lid, "contains", lid))
-            // sequence: previous chunk -precedes-> this chunk
-            if !str_eq(last_chunk, "") {
-                edges = el_list_append(edges, mk_edge(last_chunk, "precedes", lid))
-            }
-            // sectioning: most-recent heading -section_of-> this chunk
-            if is_heading {
-                last_heading = lid
-            } else {
-                if !str_eq(last_heading, "") {
-                    edges = el_list_append(edges, mk_edge(last_heading, "section_of", lid))
-                }
-            }
-            last_chunk = lid
-            ci = ci + 1
+    while ci < nc {
+        let raw: String = el_list_get(chunks, ci)
+        let lid: String = root_lid + ":c" + int_to_str(ci)
+        let is_heading: Bool = str_starts_with(raw, "#")
+        let kind: String = if is_heading { "kind:heading" } else { "kind:chunk" }
+        nodes = el_list_append(nodes, mk_node(lid, raw,
+            "Knowledge", "Semantic", "0.55", "0.55", "0.9", tagbase + " " + kind))
+        // containment: root -contains-> chunk
+        edges = el_list_append(edges, mk_edge(root_lid, "contains", lid))
+        // sequence: previous chunk -precedes-> this chunk
+        if !str_eq(last_chunk, "") {
+            edges = el_list_append(edges, mk_edge(last_chunk, "precedes", lid))
         }
-        idx = idx + 1
-    }
-    // package: we return the two lists concatenated via a sentinel; but EL
-    // lists can't nest heterogeneously here, so we instead return nodes and
-    // rely on the caller holding edges by reference is not possible — so we
-    // encode both into one list: [ "N" + nodejson ... , "E" + edgejson ... ].
-    let packed: [String] = el_list_empty()
-    let a: Int = 0
-    let an: Int = el_list_len(nodes)
-    while a < an { packed = el_list_append(packed, "N" + el_list_get(nodes, a)) a = a + 1 }
-    let b: Int = 0
-    let bn: Int = el_list_len(edges)
-    while b < bn { packed = el_list_append(packed, "E" + el_list_get(edges, b)) b = b + 1 }
-    return packed
-}
-
-// STRUCTURED / RAW-GEOMETRY: ingest structured primitives (phonetics/formants,
-// instrument signatures, scene primitives) as GEOMETRY, faithfully. Normalized
-// input shape:
-//   {"dataset":"<name>","primitive_type":"<t>",
-//    "records":[{"key":"<id>","features":{...categorical...},"attributes":{...}}]}
-// Each record -> a primitive node; each categorical feature -> a SHARED feature
-// node (deduped across records: many primitives -> one feature node = real
-// connective geometry, meaning saturates); numeric attributes fold into the
-// primitive's content (unique values, no dedup benefit). This is knowledge
-// represented as geometry, not prose — the path speech/music/image ingest on.
-fn transduce_structured(nodes: [String], edges: [String], js: String,
-                    prov: String, ground: String, steward: String,
-                    root_lid: String) -> [String] {
-    // grounding integrity: the SOURCE may declare its own epistemic grounding
-    // (measured / derived / convention / ...) via a top-level "grounding" field;
-    // honor it faithfully over the ingest-time default. This keeps the per-node
-    // ground: facet consistent with the source's honest self-description.
-    let src_ground: String = json_get_string(js, "grounding")
-    let use_ground: String = if str_eq(src_ground, "") { ground } else { src_ground }
-    let tagbase: String = "prov:" + prov + " ground:" + use_ground + " steward:" + steward
-    let dsname: String = json_get_string(js, "dataset")
-    let ptype: String = json_get_string(js, "primitive_type")
-    // capture the source's own scholarly provenance citation (verbatim) onto
-    // the dataset root — faithful attribution, retrievable, reachable from every
-    // primitive via its -contains- edge back to the root.
-    let src_cite: String = json_get_string(js, "provenance")
-    let root_content: String = "dataset: " + dsname + " (" + ptype + ")"
-    if !str_eq(src_cite, "") { root_content = root_content + " | provenance: " + src_cite }
-    nodes = el_list_append(nodes, mk_node(root_lid, root_content,
-        "Concept", "Semantic", "0.6", "0.6", "0.9", tagbase + " kind:dataset"))
-
-    let recs: String = json_get_raw(js, "records")
-    let nr: Int = json_array_len(recs)
-    let r: Int = 0
-    while r < nr {
-        let rec: String = json_array_get(recs, r)
-        let rkey: String = json_get_string(rec, "key")
-        let attrs: String = json_get_raw(rec, "attributes")
-        // faithful compact serialization of the primitive's numeric signature
-        let attr_str: String = flatten_pairs(attrs)
-        let content: String = ptype + " " + rkey
-        if !str_eq(attr_str, "") { content = content + " | " + attr_str }
-        let plid: String = root_lid + ":" + rkey
-        nodes = el_list_append(nodes, mk_node(plid, content,
-            "Concept", "Semantic", "0.6", "0.6", "0.92",
-            tagbase + " kind:primitive primitive:" + ptype + " key:" + rkey))
-        edges = el_list_append(edges, mk_edge(root_lid, "contains", plid))
-
-        // categorical features -> SHARED (deduped) feature nodes + labelled edges
-        let feats: String = json_get_raw(rec, "features")
-        let fkeys: [String] = json_object_keys(feats)
-        let fk: Int = el_list_len(fkeys)
-        let k: Int = 0
-        while k < fk {
-            let fname: String = el_list_get(fkeys, k)
-            let fval: String = json_get_string(feats, fname)
-            // shared feature node: content is the feature=value pair; identical
-            // pairs across records dedup onto ONE node (the geometry).
-            let flid: String = "feat:" + fname + "=" + fval
-            let fcontent: String = fname + "=" + fval
-            nodes = el_list_append(nodes, mk_node(flid, fcontent,
-                "Concept", "Semantic", "0.5", "0.5", "0.9",
-                tagbase + " kind:feature feature:" + fname))
-            edges = el_list_append(edges, mk_edge(plid, fname, flid))
-            k = k + 1
+        // sectioning: most-recent heading -section_of-> this chunk
+        if is_heading {
+            last_heading = lid
+        } else {
+            if !str_eq(last_heading, "") {
+                edges = el_list_append(edges, mk_edge(last_heading, "section_of", lid))
+            }
         }
-        r = r + 1
+        last_chunk = lid
+        ci = ci + 1
     }
-    let packed: [String] = el_list_empty()
-    let a: Int = 0
-    let an: Int = el_list_len(nodes)
-    while a < an { packed = el_list_append(packed, "N" + el_list_get(nodes, a)) a = a + 1 }
-    let b: Int = 0
-    let bn: Int = el_list_len(edges)
-    while b < bn { packed = el_list_append(packed, "E" + el_list_get(edges, b)) b = b + 1 }
-    return packed
-}
 
-// flatten a flat JSON object of scalar fields into "k=v k=v" (faithful; values
-// verbatim). Used for numeric attribute signatures.
-fn flatten_pairs(obj: String) -> String {
-    if str_eq(obj, "") { return "" }
-    let keys: [String] = json_object_keys(obj)
-    let n: Int = el_list_len(keys)
-    let out: String = ""
-    let i: Int = 0
-    while i < n {
-        let k: String = el_list_get(keys, i)
-        // json_get_raw returns the raw token — works for NUMBERS (bare, e.g.
-        // "270") where json_get_string yields "" for non-string values. Strip
-        // surrounding quotes if the value happens to be a string token.
-        let raw: String = json_get_raw(obj, k)
-        let v: String = str_replace(raw, "\"", "")
-        let sep: String = if i == 0 { "" } else { " " }
-        out = out + sep + k + "=" + v
-        i = i + 1
-    }
-    return out
+    // package both lists into one, "N"/"E"-prefixed (see merge_packed).
+    let packed: [String] = el_list_empty()
+    let pn_i: Int = 0
+    let pn_n: Int = el_list_len(nodes)
+    while pn_i < pn_n { packed = el_list_append(packed, "N" + el_list_get(nodes, pn_i)) pn_i = pn_i + 1 }
+    let pe_i: Int = 0
+    let pe_n: Int = el_list_len(edges)
+    while pe_i < pe_n { packed = el_list_append(packed, "E" + el_list_get(edges, pe_i)) pe_i = pe_i + 1 }
+    return packed
 }
 
 // unpack the "N"/"E"-prefixed packed list back into two lists, then merge
@@ -594,18 +508,7 @@ fn basename(path: String) -> String {
     return el_list_get(parts, n - 1)
 }
 
-fn ends_with_ci(s: String, suf: String) -> Bool {
-    return str_ends_with(str_to_lower(s), suf)
-}
-
-fn is_text_file(path: String) -> Bool {
-    return ends_with_ci(path, ".md") || ends_with_ci(path, ".txt")
-        || ends_with_ci(path, ".markdown") || ends_with_ci(path, ".text")
-}
-
 // default ingestion grounding; overridable per-invocation via INGEST_GROUND.
-// Note: a source's OWN top-level "grounding" field (structured) takes precedence
-// over this — the author's honest self-description wins.
 fn default_ground() -> String {
     let g: String = env("INGEST_GROUND")
     if str_eq(g, "") { return "extracted" }
@@ -618,25 +521,67 @@ fn default_steward() -> String {
     return s
 }
 
-// ingest one file -> report JSON
+// Mechanical fidelity guard — NOT a content-type test. fs_read()'s el_val_t
+// result truncates at the first embedded NUL byte under El's strlen-based
+// string ops (see fs_size's doc comment in runtime/el_runtime.h); comparing
+// its length against fs_size() (a real stat()-based byte count) is a
+// technical fact about whether the string channel captured the file intact
+// — computed the same way for a poem, a JSON file, or a WAV, and saying
+// nothing about what the file IS. When the counts agree, `text` is
+// trustworthy verbatim. When they don't (silent truncation happened),
+// rebuild the payload as base64-encoded fixed-size windows read directly
+// off disk (fs_read_b64_chunk — binary-safe in C), joined with the same
+// "\n\n" boundary marker transduce()'s generic scan already looks for, so
+// transduce() sees one ordinary boundary-delimited payload and runs its one
+// algorithm on it exactly as it would on prose — it never learns that a
+// fidelity problem occurred upstream, let alone why.
+fn file_source_string(path: String, text: String, real_size: Int) -> String {
+    if real_size <= 0 { return text }
+    if str_len(text) == real_size { return text }
+    // 3072 raw bytes -> 4096 base64 chars (3 divides evenly into base64's
+    // 3-byte/4-char ratio); keeps each resulting node's content a clean,
+    // bounded, low-kilobytes unit, same order of magnitude as the fixed
+    // fallback window in transduce() itself.
+    let win: Int = 3072
+    let out: String = ""
+    let off: Int = 0
+    let first: Bool = true
+    while off < real_size {
+        let chunk_b64: String = fs_read_b64_chunk(path, off, win)
+        if str_eq(chunk_b64, "") {
+            off = real_size
+        } else {
+            let sep: String = if first { "" } else { "\n\n" }
+            out = out + sep + chunk_b64
+            first = false
+            off = off + win
+        }
+    }
+    return out
+}
+
+// ingest one file -> report JSON. Uniform for every file regardless of
+// extension or content — transduce() decides nothing about content-type, so
+// neither does this function; it only decides whether the raw bytes made it
+// through the read intact (file_source_string), which is a fidelity
+// question, not a format one.
 fn ingest_file(path: String) -> String {
+    let real_size: Int = fs_size(path)
     let text: String = fs_read(path)
-    if str_eq(text, "") {
+    let source: String = file_source_string(path, text, real_size)
+    if str_eq(source, "") {
         return "{\"error\":\"empty or unreadable\",\"path\":" + j_q(path) + "}"
     }
     let prov: String = "file:" + path
-    if ends_with_ci(path, ".json") {
-        let packed: [String] = transduce_structured(el_list_empty(), el_list_empty(),
-            text, prov, default_ground(), default_steward(), "ds:" + basename(path))
-        return merge_packed(packed)
-    }
-    let packed: [String] = transduce_prose(el_list_empty(), el_list_empty(),
-        text, prov, default_ground(), default_steward(),
+    let packed: [String] = transduce(el_list_empty(), el_list_empty(),
+        source, prov, default_ground(), default_steward(),
         "doc:" + basename(path), basename(path))
     return merge_packed(packed)
 }
 
-// ingest a directory: walk one level, ingest each supported file, aggregate
+// ingest a directory: walk one level, ingest every file found, aggregate.
+// No extension filter — transduce() handles any payload uniformly now, so
+// there is no content-type gate at the directory boundary either.
 fn ingest_dir(path: String) -> String {
     let entries: [String] = fs_list(path)
     let n: Int = el_list_len(entries)
@@ -649,14 +594,12 @@ fn ingest_dir(path: String) -> String {
         let name: String = str_trim(el_list_get(entries, i))
         if !str_eq(name, "") {
             let full: String = path + "/" + name
-            if is_text_file(full) || ends_with_ci(full, ".json") {
-                println("FILE " + full)
-                let rep: String = ingest_file(full)
-                tot_created = tot_created + json_get_int(rep, "nodes_created")
-                tot_deduped = tot_deduped + json_get_int(rep, "nodes_deduped")
-                tot_edges = tot_edges + json_get_int(rep, "edges_added")
-                files = files + 1
-            }
+            println("FILE " + full)
+            let rep: String = ingest_file(full)
+            tot_created = tot_created + json_get_int(rep, "nodes_created")
+            tot_deduped = tot_deduped + json_get_int(rep, "nodes_deduped")
+            tot_edges = tot_edges + json_get_int(rep, "edges_added")
+            files = files + 1
         }
         i = i + 1
     }
@@ -667,11 +610,12 @@ fn ingest_dir(path: String) -> String {
         ",\"edges_accepted\":" + int_to_str(tot_edges) + "}"
 }
 
-// ingest a url: fetch, treat body as prose (faithful extraction of what's there)
+// ingest a url: fetch, hand the body straight to transduce (faithful
+// extraction of what's there — no interpretation of what it is)
 fn ingest_url(url: String) -> String {
     let body: String = http_get(url)
     if str_eq(body, "") { return "{\"error\":\"empty fetch\",\"url\":" + j_q(url) + "}" }
-    let packed: [String] = transduce_prose(el_list_empty(), el_list_empty(),
+    let packed: [String] = transduce(el_list_empty(), el_list_empty(),
         body, "url:" + url, "extracted", "public-web",
         "url:" + url, url)
     return merge_packed(packed)
@@ -686,7 +630,7 @@ fn ingest_llm(query: String) -> String {
     let resp: String = http_post_json("http://127.0.0.1:11434/api/generate", body)
     let answer: String = json_get_string(resp, "response")
     if str_eq(answer, "") { return "{\"error\":\"no model response\"}" }
-    let packed: [String] = transduce_prose(el_list_empty(), el_list_empty(),
+    let packed: [String] = transduce(el_list_empty(), el_list_empty(),
         answer, "llm:" + model + ":" + query, "candidate-provisional", "guide-provisional",
         "llm:" + query, "guide answer: " + query)
     return merge_packed(packed)
@@ -728,6 +672,19 @@ fn ingest_stream(path: String) -> String {
 // SECTION G — ENTRY
 // ═══════════════════════════════════════════════════════════════════════════
 
+// INGEST_KIND selects an ACQUISITION mechanism only — dir/file/url/llm/
+// stream — i.e. which RPC shape to use to go get the bytes (walk a
+// directory, open a file, fetch a URL, query an LLM, read a turn-stream).
+// That is a genuinely unavoidable choice at the process-entry boundary
+// (nothing about the string "/tmp/x" tells you whether it's a file to read
+// or a stream to read line-by-line, or distinguishes an LLM query from a
+// path), so it cannot be dropped the way content-type dispatch was.
+// It is NOT a content-type flag: it says nothing about what's inside the
+// bytes once fetched, and none of the five ingest_* functions it selects
+// among interpret their payload differently by content shape anymore —
+// they all hand off to the single, format-agnostic transduce(). The old
+// "structured" value (a caller-declared alias for "file", used only to hint
+// the now-removed JSON-vs-prose branch) is gone along with that branch.
 let kind: String = env("INGEST_KIND")
 let arg: String = env("INGEST_ARG")
 
@@ -741,20 +698,16 @@ if str_eq(kind, "dir") {
     if str_eq(kind, "file") {
         report = ingest_file(arg)
     } else {
-        if str_eq(kind, "structured") {
-            report = ingest_file(arg)
+        if str_eq(kind, "url") {
+            report = ingest_url(arg)
         } else {
-            if str_eq(kind, "url") {
-                report = ingest_url(arg)
+            if str_eq(kind, "llm") {
+                report = ingest_llm(arg)
             } else {
-                if str_eq(kind, "llm") {
-                    report = ingest_llm(arg)
+                if str_eq(kind, "stream") {
+                    report = ingest_stream(arg)
                 } else {
-                    if str_eq(kind, "stream") {
-                        report = ingest_stream(arg)
-                    } else {
-                        report = "{\"error\":\"unknown INGEST_KIND: " + kind + "\"}"
-                    }
+                    report = "{\"error\":\"unknown INGEST_KIND: " + kind + "\"}"
                 }
             }
         }

lang/el-compiler/src/codegen.el

@@ -2766,6 +2766,8 @@ fn builtin_arity(name: String) -> Int {
     if str_eq(name, "fs_read") { return 1 }
     if str_eq(name, "fs_write") { return 2 }
     if str_eq(name, "fs_list") { return 1 }
+    if str_eq(name, "fs_size") { return 1 }
+    if str_eq(name, "fs_read_b64_chunk") { return 3 }
     // JSON
     if str_eq(name, "json_get") { return 2 }
     if str_eq(name, "json_parse") { return 1 }

lang/runtime/build_vindex_bench.sh

@@ -1,51 +0,0 @@
-#!/bin/bash
-# build_vindex_bench.sh — build the vindex_bench oracle/proof harness, with
-# the real ggml + hand-rolled-Metal batch-cosine strategies on Darwin and a
-# zero-dependency CPU-only stub everywhere else. Mirrors the two-step recipe
-# documented in vindex_bench.c's own header comment; this script exists so
-# that recipe is one command, not a copy-pasted paragraph.
-#
-# Darwin build links FOUR strategy translation units:
-#   eg_cosine_batch.c                    — the Factory (always)
-#   eg_cosine_batch_strategy_cpu.c       — universal fallback (always)
-#   eg_cosine_batch_strategy_ggml.c      — ggml + dynamic Metal backend plugin
-#   eg_cosine_batch_strategy_metal_hand.m — PR #114's original hand-rolled
-#                                            Metal shader, preserved as one
-#                                            selectable strategy
-# plus -DEG_HAVE_STRATEGY_GGML -DEG_HAVE_STRATEGY_METAL_HAND so the Factory
-# (and vindex_bench.c's own direct strategy comparison) knows both exist.
-#
-# ggml is resolved via `brew --prefix ggml` when available (portable across
-# Intel /usr/local and Apple Silicon /opt/homebrew installs), falling back to
-# /opt/homebrew if brew isn't on PATH. Override with GGML_PREFIX=... env var.
-#
-# Usage: ./build_vindex_bench.sh [output_path]
-set -euo pipefail
-cd "$(dirname "$0")"
-OUT="${1:-./vindex_bench}"
-CC="${CC:-cc}"
-
-if [ "$(uname -s)" = "Darwin" ]; then
-  GGML_PREFIX="${GGML_PREFIX:-$(brew --prefix ggml 2>/dev/null || echo /opt/homebrew)}"
-  echo "== Darwin: building with the ggml + hand-rolled-Metal strategies (ggml prefix: $GGML_PREFIX) =="
-
-  "$CC" -O2 -std=c11 -x objective-c \
-    -c eg_cosine_batch_strategy_metal_hand.m -o /tmp/eg_cosine_batch_strategy_metal_hand.o \
-    -framework Metal -framework Foundation
-
-  "$CC" -O2 -std=c11 -DEG_HAVE_STRATEGY_GGML -DEG_HAVE_STRATEGY_METAL_HAND -w \
-    -I"$GGML_PREFIX/include" \
-    vindex_bench.c engram_vindex.c \
-    eg_cosine_batch.c eg_cosine_batch_strategy_cpu.c eg_cosine_batch_strategy_ggml.c \
-    /tmp/eg_cosine_batch_strategy_metal_hand.o \
-    -L"$GGML_PREFIX/lib" -lggml -lggml-base \
-    -Wl,-rpath,"$GGML_PREFIX/lib" \
-    -lm -framework Metal -framework Foundation -o "$OUT"
-else
-  echo "== non-Darwin: building with the CPU-only fallback strategy (no ggml, no Metal) =="
-  "$CC" -O2 -std=c11 -w vindex_bench.c engram_vindex.c \
-    eg_cosine_batch.c eg_cosine_batch_strategy_cpu.c \
-    -lm -o "$OUT"
-fi
-
-echo "built: $OUT"

lang/runtime/eg_cosine_batch.c

@@ -1,121 +0,0 @@
-/* eg_cosine_batch.c — the Factory. Implements the stable public interface
- * declared in eg_cosine_batch.h by selecting ONE concrete
- * EgCosineBatchStrategy (eg_cosine_batch_strategy.h) and dispatching every
- * call to it. This is the ONLY file that branches on EG_HAVE_STRATEGY_*
- * (build-time: which strategy .c/.m files were actually compiled in for
- * this platform) — call sites never see those macros.
- *
- * Selection is lazy (first call) and cached — mirrors the lazy-init caching
- * every individual strategy already does internally, so there is no added
- * per-call cost after the first.
- *
- * Selection mechanism (env var + build-time + runtime capability probe, all
- * three, exactly as directed):
- *   - BUILD-TIME decides which strategies exist to choose from at all: a
- *     Darwin build compiles+links the ggml strategy and the hand-rolled
- *     Metal strategy (EG_HAVE_STRATEGY_GGML / EG_HAVE_STRATEGY_METAL_HAND
- *     both defined); a non-Darwin build compiles neither, matching PR #114's
- *     original Linux behavior exactly (CPU-fallback only, no Objective-C
- *     compiler or Metal frameworks required).
- *   - RUNTIME CAPABILITY PROBE: each candidate strategy's own available()
- *     does the real, cheap-after-first-call check (device present, backend
- *     plugin loaded, pipeline compiled) — never assumed from build-time
- *     alone. A build that HAS the ggml strategy compiled in but is running
- *     on hardware/software where it can't actually initialize (backend
- *     plugin missing, no GPU) correctly falls through to the next candidate.
- *   - ENV VAR gives explicit, debuggable override for either axis:
- *       EL_COSINE_BATCH_STRATEGY = "ggml" | "metal" | "cpu" | unset/"auto"
- *         forces a specific strategy (falling back to cpu if the forced one
- *         isn't actually available), or leaves the default auto-preference
- *         order in place.
- *       EL_METAL_COSINE = 0/n/N/f/F   (back-compat with PR #114's vindex_bench
- *         gate) disables ALL GPU-backed strategies outright, same as before.
- *
- * DEFAULT preference order when nothing is forced: ggml, then hand-rolled
- * Metal, then CPU fallback — first candidate whose available() reports true
- * wins. This is what makes "stop hand-rolling GPU kernels, use ggml" real
- * rather than nominal: ggml is what actually runs by default on this
- * machine today (see the PR body for the measured numbers backing that).
- */
-#include "eg_cosine_batch.h"
-#include "eg_cosine_batch_strategy.h"
-
-#include <stdlib.h>
-#include <string.h>
-
-static bool g_selected = false;
-static const EgCosineBatchStrategy* g_active = NULL;
-
-static bool eg_env_truthy_off(const char* v) {
-    return v && (v[0]=='0' || v[0]=='n' || v[0]=='N' || v[0]=='f' || v[0]=='F');
-}
-
-static const EgCosineBatchStrategy* eg_select_strategy(void) {
-    if (g_selected) return g_active;
-    g_selected = true;
-
-    const EgCosineBatchStrategy* cpu = eg_cosine_batch_strategy_cpu();
-    const char* force = getenv("EL_COSINE_BATCH_STRATEGY");
-    const char* legacy_off = getenv("EL_METAL_COSINE");
-
-    if (eg_env_truthy_off(legacy_off)) { g_active = cpu; return g_active; }
-
-    if (force && strcmp(force, "cpu") == 0) { g_active = cpu; return g_active; }
-
-    if (force && strcmp(force, "ggml") == 0) {
-#ifdef EG_HAVE_STRATEGY_GGML
-        const EgCosineBatchStrategy* s = eg_cosine_batch_strategy_ggml();
-        if (s->available()) { g_active = s; return g_active; }
-#endif
-        g_active = cpu; return g_active;
-    }
-
-    if (force && strcmp(force, "metal") == 0) {
-#ifdef EG_HAVE_STRATEGY_METAL_HAND
-        const EgCosineBatchStrategy* s = eg_cosine_batch_strategy_metal_hand();
-        if (s->available()) { g_active = s; return g_active; }
-#endif
-        g_active = cpu; return g_active;
-    }
-
-    /* auto (unset, or any other value): ggml -> metal-hand -> cpu, first
-     * available wins. */
-#ifdef EG_HAVE_STRATEGY_GGML
-    {
-        const EgCosineBatchStrategy* s = eg_cosine_batch_strategy_ggml();
-        if (s->available()) { g_active = s; return g_active; }
-    }
-#endif
-#ifdef EG_HAVE_STRATEGY_METAL_HAND
-    {
-        const EgCosineBatchStrategy* s = eg_cosine_batch_strategy_metal_hand();
-        if (s->available()) { g_active = s; return g_active; }
-    }
-#endif
-    g_active = cpu;
-    return g_active;
-}
-
-bool eg_cosine_batch_available(void) {
-    return eg_select_strategy()->available();
-}
-
-const char* eg_cosine_batch_strategy_name(void) {
-    return eg_select_strategy()->name;
-}
-
-bool eg_cosine_batch(const float* query, int32_t qdim,
-                      const float* const* node_ptrs,
-                      const int32_t* node_dims,
-                      int32_t n,
-                      double* out_scores) {
-    return eg_select_strategy()->batch(query, qdim, node_ptrs, node_dims, n, out_scores);
-}
-
-bool eg_cosine_batch_multi(const float* queries, int32_t qdim, int32_t nq,
-                            const float* const* node_ptrs,
-                            const int32_t* node_dims,
-                            int32_t n,
-                            double* out_scores) {
-    return eg_select_strategy()->batch_multi(queries, qdim, nq, node_ptrs, node_dims, n, out_scores);
-}

lang/runtime/eg_cosine_batch.h

@@ -1,106 +0,0 @@
-/* eg_cosine_batch.h — stable Adapter interface over batch-cosine-similarity
- * BACKEND STRATEGIES. This header is the ONE thing call sites (el_runtime.c,
- * vindex_bench.c, ...) talk to. Plain C11, safe to #include on every
- * platform — the symbols declared here always exist and always link,
- * regardless of what backend actually runs underneath. Zero #ifdef at call
- * sites: which concrete strategy executes (ggml/Metal, hand-rolled Metal, or
- * the always-false CPU fallback) is resolved once, lazily, inside
- * eg_cosine_batch.c's factory — see eg_cosine_batch_strategy.h for that.
- *
- * This supersedes eg_metal_cosine.h (PR #114's single hand-rolled-Metal-only
- * bridge). The contract is UNCHANGED — same shapes, same sentinel, same
- * never-partial guarantee, same "caller must always be prepared to fall back
- * to its own scalar per-node loop" rule — only the name changed, because the
- * thing behind it is no longer "the Metal bridge," it is "whichever batch-
- * cosine strategy the factory picked." eg_metal_cosine.h's original doc
- * comments (byte-for-byte, this file is the direct descendant) are preserved
- * below since they remain the precise spec any strategy must honor.
- *
- * On ANY failure at ANY step — no compute device, compile/init error, alloc
- * failure, bad args — every function here returns false and writes nothing.
- * Out-params are either fully populated or left completely untouched, never
- * partial. Callers MUST always be prepared to fall back to their own scalar
- * per-node CPU loop unconditionally. These functions must never crash, throw,
- * or hang the calling process — several call sites run inside a long-lived
- * daemon's request-handling hot path.
- */
-#ifndef EG_COSINE_BATCH_H
-#define EG_COSINE_BATCH_H
-
-#include <stdint.h>
-#include <stdbool.h>
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-/* Batched cosine similarity: one query vector against `n` node vectors.
- *
- *   query      — qdim floats, the query embedding. Raw/unnormalized.
- *   qdim       — query dimensionality (e.g. 768 for nomic-embed-text).
- *   node_ptrs  — array of n pointers, node_ptrs[i] pointing at a (possibly
- *                differently-owned, possibly NULL) float vector for node i.
- *                NOT required to be contiguous — every strategy performs the
- *                gather into a packed row-major matrix internally, exactly
- *                mirroring how EngramNode.emb is one malloc per node.
- *   node_dims  — array of n ints, node_dims[i] = that node's real emb_dim
- *                (0 or mismatched vs qdim => that node scores -2.0, matching
- *                eg_cosine's null/dim-mismatch/zero-norm sentinel exactly).
- *   n          — number of nodes.
- *   out_scores — caller-owned array of n doubles; out_scores[i] is filled
- *                with the cosine similarity of node i against query, or
- *                -2.0 for a null/dim-mismatched/zero-norm node — bit-for-bit
- *                the same contract as eg_cosine(node_ptrs[i], query, qdim).
- *
- * Returns true iff a real backend strategy ran and out_scores was fully
- * populated. Returns false (out_scores left untouched) on ANY failure or
- * unavailability — no compute device, compile/init failure, allocation
- * failure, n<=0, qdim<=0, null query/node_ptrs/node_dims/out_scores.
- */
-bool eg_cosine_batch(const float* query, int32_t qdim,
-                      const float* const* node_ptrs,
-                      const int32_t* node_dims,
-                      int32_t n,
-                      double* out_scores);
-
-/* True iff a real (non-CPU-fallback) strategy is available right now (cheap
- * after the first call — cached). Purely informational (e.g. a startup log
- * line or /api/stats field); callers should still treat a false return from
- * eg_cosine_batch()/eg_cosine_batch_multi() itself as the authoritative
- * fallback signal, not this function. */
-bool eg_cosine_batch_available(void);
-
-/* Which concrete strategy is currently selected — "ggml", "metal-hand",
- * or "cpu-fallback". Purely informational/diagnostic, same spirit as
- * eg_cosine_batch_available(). Never NULL. */
-const char* eg_cosine_batch_strategy_name(void);
-
-/* Multi-query batched cosine: nq query vectors against the SAME n node
- * vectors, in one call. A real strategy uploads/prepares the node population
- * ONCE and reuses it for every query, instead of nq separate
- * eg_cosine_batch() calls each paying the full gather+upload cost — PR #114
- * measured this necessary: at N~=13.7k/dim=768, repeating the single-query
- * call per query was slower than the CPU baseline; batching queries together
- * is what makes a GPU-backed path a real win at this shape. Use this
- * whenever multiple queries will run against an unchanged (or
- * rarely-changing) node population; use eg_cosine_batch() for a genuinely
- * one-off comparison.
- *
- *   queries    — nq*qdim floats, row-major (query i at queries+i*qdim).
- *   out_scores — caller-owned nq*n doubles, row-major
- *                (out_scores[i*n+j] = cosine(queries[i], node j)), same
- *                -2.0 sentinel semantics as eg_cosine_batch().
- *
- * Returns true iff a real strategy ran and out_scores was fully populated
- * (all nq*n entries); false (untouched) on any failure/unavailability. */
-bool eg_cosine_batch_multi(const float* queries, int32_t qdim, int32_t nq,
-                            const float* const* node_ptrs,
-                            const int32_t* node_dims,
-                            int32_t n,
-                            double* out_scores);
-
-#ifdef __cplusplus
-}
-#endif
-
-#endif /* EG_COSINE_BATCH_H */

lang/runtime/eg_cosine_batch.metal

@@ -1,156 +0,0 @@
-/* eg_cosine_batch.metal — batched cosine similarity, one query vs N node vectors.
- *
- * GPU-shaped counterpart to eg_cosine() in el_runtime.c: same math, same
- * dim-mismatch/zero-norm sentinel (-2.0), applied to N independent rows in
- * parallel instead of one pair at a time in a CPU loop.
- *
- * Semantics MUST match eg_cosine() exactly:
- *   - inputs are raw, UNNORMALIZED vectors (nomic-embed-text magnitudes are
- *     not 1.0) — this kernel computes the full dot/(|a|*|b|) cosine, not a
- *     plain dot product.
- *   - a node whose declared dim differs from the query dim, or whose norm is
- *     zero, scores exactly -2.0 (below any valid cosine in [-1,1]), so a
- *     caller doing `if (score > threshold)` behaves identically whether the
- *     scalar or the batched path filled the array.
- *
- * Precision: Apple GPUs do not support double in Metal Shading Language —
- * everything here is float32. eg_cosine accumulates in CPU double, but its
- * *inputs* are float32 embeddings, so the achievable precision ceiling is
- * bounded by the input data regardless of accumulator width. To keep the
- * float32 reduction from drifting relative to the double-accumulated CPU
- * result across dim=768 terms, each thread accumulates with 4 independent
- * partial sums (unrolled) rather than one running scalar — the same
- * error-reduction trick already used by the CPU brute-force loop in
- * vindex_bench.c. The measured float-vs-double delta is reported in the PR
- * description; this is not assumed to be "close enough" without measurement.
- */
-#include <metal_stdlib>
-using namespace metal;
-
-/* Per-dispatch invariants. `dim` is the query's dimensionality — the
- * dimensionality every comparable node vector must match. */
-struct EgCosineParams {
-    uint n;      /* number of node rows */
-    uint dim;    /* vector width (both query and node rows are `dim` wide in
-                  * the packed buffer; node_dims[] carries each node's REAL
-                  * embedded dim for the mismatch check) */
-};
-
-/* One thread per node row. node_matrix is n*dim floats, row-major, packed at
- * `dim` stride regardless of a row's real dim (the CPU side zero-pads or
- * skips packing rows that don't match — see eg_cosine_batch_metal in
- * eg_metal_cosine.m for the exact packing contract). node_dims[i] is the
- * node's true emb_dim, used only for the mismatch sentinel — never used to
- * index, since every row is packed at uniform `dim` stride. */
-kernel void eg_cosine_batch_kernel(
-    device const float*  query        [[buffer(0)]],
-    device const float*  node_matrix  [[buffer(1)]],
-    device const int*    node_dims    [[buffer(2)]],
-    constant EgCosineParams& p        [[buffer(3)]],
-    device float*        out_scores   [[buffer(4)]],
-    uint gid [[thread_position_in_grid]])
-{
-    if (gid >= p.n) return;
-
-    if (node_dims[gid] != int(p.dim)) {
-        out_scores[gid] = -2.0f;
-        return;
-    }
-
-    device const float* row = node_matrix + (uint64_t)gid * (uint64_t)p.dim;
-
-    /* 4-way partial accumulation — same shape as vindex_bench.c's brute_topk
-     * unroll, done here for float32 accuracy rather than raw throughput. */
-    float dot0 = 0.0f, dot1 = 0.0f, dot2 = 0.0f, dot3 = 0.0f;
-    float na0  = 0.0f, na1  = 0.0f, na2  = 0.0f, na3  = 0.0f;
-    float nb0  = 0.0f, nb1  = 0.0f, nb2  = 0.0f, nb3  = 0.0f;
-
-    uint d = 0;
-    uint dim4 = p.dim & ~3u;
-    for (; d < dim4; d += 4) {
-        float a0 = row[d],   b0 = query[d];
-        float a1 = row[d+1], b1 = query[d+1];
-        float a2 = row[d+2], b2 = query[d+2];
-        float a3 = row[d+3], b3 = query[d+3];
-        dot0 += a0*b0; dot1 += a1*b1; dot2 += a2*b2; dot3 += a3*b3;
-        na0  += a0*a0; na1  += a1*a1; na2  += a2*a2; na3  += a3*a3;
-        nb0  += b0*b0; nb1  += b1*b1; nb2  += b2*b2; nb3  += b3*b3;
-    }
-    float dot = (dot0 + dot1) + (dot2 + dot3);
-    float na  = (na0  + na1)  + (na2  + na3);
-    float nb  = (nb0  + nb1)  + (nb2  + nb3);
-    for (; d < p.dim; d++) {
-        float a = row[d], b = query[d];
-        dot += a*b; na += a*a; nb += b*b;
-    }
-
-    if (na <= 0.0f || nb <= 0.0f) {
-        out_scores[gid] = -2.0f;
-        return;
-    }
-    out_scores[gid] = dot / sqrt(na * nb);
-}
-
-/* ── multi-query variant ──────────────────────────────────────────────────
- * Same per-pair math as eg_cosine_batch_kernel, but amortizes ONE upload of
- * node_matrix (the expensive part at real store size — 13k*768 floats is
- * ~42MB) across `nq` queries instead of re-uploading it once per query.
- * Measured need: a naive one-query-at-a-time loop calling the single-query
- * kernel nq times was SLOWER than the CPU oracle at N≈13.7k (re-gather +
- * re-upload dominated the actual compute) — this is the fix, not a
- * hypothetical optimization.
- *
- * 2D grid: x = node index [0,n), y = query index [0,nq). out_scores is
- * nq*n, row-major by query (out_scores[qid*n + nid]). */
-struct EgCosineMultiParams { uint n; uint dim; uint nq; };
-
-kernel void eg_cosine_batch_multi_kernel(
-    device const float*  queries      [[buffer(0)]],   /* nq*dim */
-    device const float*  node_matrix  [[buffer(1)]],   /* n*dim */
-    device const int*    node_dims    [[buffer(2)]],   /* n */
-    constant EgCosineMultiParams& p   [[buffer(3)]],
-    device float*        out_scores   [[buffer(4)]],   /* nq*n */
-    uint2 gid [[thread_position_in_grid]])
-{
-    uint nid = gid.x, qid = gid.y;
-    if (nid >= p.n || qid >= p.nq) return;
-
-    uint64_t out_idx = (uint64_t)qid * (uint64_t)p.n + (uint64_t)nid;
-
-    if (node_dims[nid] != int(p.dim)) {
-        out_scores[out_idx] = -2.0f;
-        return;
-    }
-
-    device const float* row   = node_matrix + (uint64_t)nid * (uint64_t)p.dim;
-    device const float* query = queries     + (uint64_t)qid * (uint64_t)p.dim;
-
-    float dot0 = 0.0f, dot1 = 0.0f, dot2 = 0.0f, dot3 = 0.0f;
-    float na0  = 0.0f, na1  = 0.0f, na2  = 0.0f, na3  = 0.0f;
-    float nb0  = 0.0f, nb1  = 0.0f, nb2  = 0.0f, nb3  = 0.0f;
-
-    uint d = 0;
-    uint dim4 = p.dim & ~3u;
-    for (; d < dim4; d += 4) {
-        float a0 = row[d],   b0 = query[d];
-        float a1 = row[d+1], b1 = query[d+1];
-        float a2 = row[d+2], b2 = query[d+2];
-        float a3 = row[d+3], b3 = query[d+3];
-        dot0 += a0*b0; dot1 += a1*b1; dot2 += a2*b2; dot3 += a3*b3;
-        na0  += a0*a0; na1  += a1*a1; na2  += a2*a2; na3  += a3*a3;
-        nb0  += b0*b0; nb1  += b1*b1; nb2  += b2*b2; nb3  += b3*b3;
-    }
-    float dot = (dot0 + dot1) + (dot2 + dot3);
-    float na  = (na0  + na1)  + (na2  + na3);
-    float nb  = (nb0  + nb1)  + (nb2  + nb3);
-    for (; d < p.dim; d++) {
-        float a = row[d], b = query[d];
-        dot += a*b; na += a*a; nb += b*b;
-    }
-
-    if (na <= 0.0f || nb <= 0.0f) {
-        out_scores[out_idx] = -2.0f;
-        return;
-    }
-    out_scores[out_idx] = dot / sqrt(na * nb);
-}

lang/runtime/eg_cosine_batch_strategy.h

@@ -1,82 +0,0 @@
-/* eg_cosine_batch_strategy.h — internal Strategy interface, NOT for call
- * sites (they use eg_cosine_batch.h). Only eg_cosine_batch.c's factory and
- * the concrete strategy implementation files include this.
- *
- * Each concrete strategy exposes exactly one getter returning a pointer to a
- * static, immutable EgCosineBatchStrategy vtable. Which getters actually
- * exist as linkable symbols is a BUILD-TIME concern (decided by
- * build_vindex_bench.sh / the engram daemon's own build, via which .c/.m
- * files get compiled per platform) gated by the EG_HAVE_STRATEGY_* macros
- * below — the factory in eg_cosine_batch.c is the ONLY place that branches
- * on those macros. Call sites never see them; that's the whole point of the
- * Adapter in eg_cosine_batch.h.
- *
- * Three concrete strategies exist:
- *   eg_cosine_batch_strategy_ggml()        — ggml + dynamically-loaded Metal
- *                                             backend plugin. Darwin only in
- *                                             this build; the default
- *                                             preferred strategy wherever
- *                                             available. EG_HAVE_STRATEGY_GGML.
- *   eg_cosine_batch_strategy_metal_hand()  — the original hand-rolled Metal
- *                                             compute shader from PR #114
- *                                             (eg_cosine_batch.metal),
- *                                             preserved verbatim as a
- *                                             selectable fallback strategy,
- *                                             not deleted. Darwin only.
- *                                             EG_HAVE_STRATEGY_METAL_HAND.
- *   eg_cosine_batch_strategy_cpu()         — universal always-false
- *                                             fallback. Always compiled, on
- *                                             every platform; this is what a
- *                                             non-Darwin build links
- *                                             exclusively (matching PR #114's
- *                                             eg_metal_cosine_stub.c), and
- *                                             what any platform falls back
- *                                             to when no real strategy is
- *                                             available at runtime.
- */
-#ifndef EG_COSINE_BATCH_STRATEGY_H
-#define EG_COSINE_BATCH_STRATEGY_H
-
-#include <stdint.h>
-#include <stdbool.h>
-
-#ifdef __cplusplus
-extern "C" {
-#endif
-
-typedef struct EgCosineBatchStrategy {
-    /* Stable, short, lowercase-hyphenated identifier — what
-     * eg_cosine_batch_strategy_name() surfaces. Never NULL. */
-    const char* name;
-
-    /* Cheap after the first call (lazy init, cached internally). Must never
-     * throw/crash/hang — mirrors eg_cosine_batch_available()'s contract. */
-    bool (*available)(void);
-
-    /* Same shape/contract as eg_cosine_batch() in eg_cosine_batch.h. */
-    bool (*batch)(const float* query, int32_t qdim,
-                  const float* const* node_ptrs, const int32_t* node_dims,
-                  int32_t n, double* out_scores);
-
-    /* Same shape/contract as eg_cosine_batch_multi() in eg_cosine_batch.h. */
-    bool (*batch_multi)(const float* queries, int32_t qdim, int32_t nq,
-                         const float* const* node_ptrs, const int32_t* node_dims,
-                         int32_t n, double* out_scores);
-} EgCosineBatchStrategy;
-
-#ifdef EG_HAVE_STRATEGY_GGML
-const EgCosineBatchStrategy* eg_cosine_batch_strategy_ggml(void);
-#endif
-
-#ifdef EG_HAVE_STRATEGY_METAL_HAND
-const EgCosineBatchStrategy* eg_cosine_batch_strategy_metal_hand(void);
-#endif
-
-/* Always declared/linked, on every platform/build. */
-const EgCosineBatchStrategy* eg_cosine_batch_strategy_cpu(void);
-
-#ifdef __cplusplus
-}
-#endif
-
-#endif /* EG_COSINE_BATCH_STRATEGY_H */

lang/runtime/eg_cosine_batch_strategy_cpu.c

@@ -1,45 +0,0 @@
-/* eg_cosine_batch_strategy_cpu.c — plain-C, zero-dependency universal
- * fallback strategy. Always returns false / unavailable. Direct descendant
- * of PR #114's eg_metal_cosine_stub.c, generalized from "the Metal stub" to
- * "the strategy vtable's universal fallback entry" now that multiple real
- * strategies can exist.
- *
- * Always compiled, on every platform. On Darwin builds it is the last-resort
- * strategy the factory falls back to when neither ggml nor the hand-rolled
- * Metal strategy is available at runtime (no device, compile failure, ...).
- * On non-Darwin builds it is the ONLY strategy compiled in at all — no
- * Objective-C, no Metal frameworks, no ggml/Metal backend plugin — so
- * eg_cosine_batch()/eg_cosine_batch_multi() always return false there and
- * every call site's existing CPU fallback runs unconditionally, exactly as
- * before this PR.
- */
-#include "eg_cosine_batch_strategy.h"
-
-static bool cpu_available(void) {
-    return false;
-}
-
-static bool cpu_batch(const float* query, int32_t qdim,
-                       const float* const* node_ptrs, const int32_t* node_dims,
-                       int32_t n, double* out_scores) {
-    (void)query; (void)qdim; (void)node_ptrs; (void)node_dims; (void)n; (void)out_scores;
-    return false;
-}
-
-static bool cpu_batch_multi(const float* queries, int32_t qdim, int32_t nq,
-                             const float* const* node_ptrs, const int32_t* node_dims,
-                             int32_t n, double* out_scores) {
-    (void)queries; (void)qdim; (void)nq; (void)node_ptrs; (void)node_dims; (void)n; (void)out_scores;
-    return false;
-}
-
-static const EgCosineBatchStrategy g_cpu_strategy = {
-    .name = "cpu-fallback",
-    .available = cpu_available,
-    .batch = cpu_batch,
-    .batch_multi = cpu_batch_multi,
-};
-
-const EgCosineBatchStrategy* eg_cosine_batch_strategy_cpu(void) {
-    return &g_cpu_strategy;
-}

lang/runtime/eg_cosine_batch_strategy_ggml.c

@@ -1,515 +0,0 @@
-/* eg_cosine_batch_strategy_ggml.c — the GGML Strategy, and the preferred
- * default whenever it is available (see the factory's selection order in
- * eg_cosine_batch.c).
- *
- * WHY: directive from Will Anderson — stop hand-rolling GPU kernels, use a
- * real, proven, permissively-licensed library instead. ggml (the compute
- * library underneath llama.cpp, MIT licensed) is already installed on this
- * machine as a standalone Homebrew package (`brew info ggml`), independent
- * of llama.cpp itself. This file is a genuinely bounded COMPUTE UTILITY —
- * batch cosine-similarity math — analogous to a VBD Accessor calling out to
- * infrastructure. It is explicitly NOT the engram's reasoning/persistence
- * core; using ggml here does not cross the "own the core" line, because
- * batch cosine math is infrastructure, not the graph traversal / activation
- * spreading / "thinking" that IS the core and stays 100% own-code.
- *
- * ── The real API shape (verified against the installed headers + a
- *    standalone probe program, not assumed from memory of other tensor
- *    libraries) ──────────────────────────────────────────────────────────
- *
- * ggml ships its CPU and Metal implementations as DYNAMICALLY LOADED PLUGIN
- * .so files (confirmed by nm: `ggml_backend_metal_init` is NOT an exported
- * symbol of libggml.dylib/libggml-base.dylib — it exists ONLY inside
- * libggml-metal.so under $(brew --prefix ggml)/libexec/). You cannot link
- * `-lggml-metal`; you must go through ggml's backend REGISTRY:
- *
- *   1. ggml_backend_load_all_from_path(dir) — dlopen()s every backend plugin
- *      .so found in `dir` and registers its device(s). We point this at
- *      $(brew --prefix ggml)/libexec (resolved once, at build+init time; see
- *      eg_ggml_backend_dir() below) rather than relying on
- *      ggml_backend_load_all()'s own default search heuristics, which are
- *      tuned for an installed llama.cpp-style app bundle layout, not an
- *      arbitrary `cc`-built binary invoked from an arbitrary cwd — the exact
- *      same "must not silently fall back to CPU for reasons that have
- *      nothing to do with GPU availability" concern PR #114's hand-rolled
- *      bridge already documented for its own embedded-shader-source choice.
- *   2. ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_GPU) — find the
- *      registered Metal device.
- *   3. ggml_backend_dev_init(dev, NULL) — get a live ggml_backend_t.
- *   4. Build a tiny ggml_context (no_alloc=true; it holds only tensor
- *      metadata, not data), declare 2D F32 tensors, ggml_mul_mat(nodes,
- *      query) — ggml's documented convention: A is [k cols, n rows], B is
- *      [k cols, m rows] (transposed internally), result is [n cols, m rows]
- *      — i.e. mul_mat(node_matrix[dim,n], query_matrix[dim,nq]) yields
- *      out[n,nq] where out[j*n+i] = dot(node_i, query_j). A row-major
- *      (dim,n) node matrix and a row-major (dim,nq) query matrix is EXACTLY
- *      the packed layout the hand-rolled Metal kernel already used — one
- *      matmul replaces the whole per-row dot-product loop.
- *   5. ggml_backend_alloc_ctx_tensors(ctx, backend) to actually allocate
- *      device buffers for those tensors, ggml_backend_tensor_set() to upload,
- *      ggml_backend_graph_compute() to run, ggml_backend_tensor_get() to
- *      read back.
- *
- * This exact sequence was verified end-to-end in a standalone probe (build
- * it yourself: see the PR description) against a plain-C CPU dot product —
- * bit-for-bit correct within float rounding. Real numbers against the
- * el_runtime.c CPU oracle are reported in the PR body via vindex_bench.
- *
- * ggml_mul_mat only computes the raw dot products — it has no notion of
- * "cosine" or of this codebase's -2.0 dim-mismatch/null/zero-norm sentinel.
- * Per the adapter's directive: gather only VALID, uniform-dim rows into the
- * packed matrix sent to the GPU (skipping null/mismatched rows entirely,
- * rather than the hand-rolled kernel's zero-pad-and-sentinel-in-shader
- * approach), then scatter -2.0 back for every row that was excluded — same
- * gather/scatter contract eg_cosine_batch.h documents. Norms (||node||,
- * ||query||) are computed on the CPU host in the same pass that already
- * touches every element to gather/convert — essentially free — using the
- * same 4-way-partial-sum accumulation the hand-rolled kernel and the CPU
- * oracle both use, so the float32 error profile stays comparable across all
- * three strategies. Only the O(n*dim*nq) dot-product matmul — the actual
- * expensive part — is offloaded to the GPU.
- *
- * ── Precision: the ne11<=8 chunking, and why it is not optional ──────────
- *
- * The claim in the first version of this file — "ggml_mul_mat on F32 x F32
- * inputs computes in F32 on the Metal backend" — is WRONG, and the 0.9933
- * id-recall it shipped with (vs the hand-rolled kernel's 0.9997) was the
- * symptom. ggml-metal has two F32xF32 matmul kernels and picks between them
- * purely on ne11 (the number of B rows == our query count):
- *
- *   ne11 <= 8  -> kernel_mul_mv_ext_f32_f32_* / kernel_mul_mv_f32_f32_*
- *                 templated <float, float> — genuine F32 accumulation.
- *   ne11 >  8  -> kernel_mul_mm_f32_f32, which is templated
- *                 <half, half4x4, simdgroup_half8x8, half, half2x4,
- *                  simdgroup_half8x8, ...> — i.e. BOTH operands are narrowed
- *                 to F16 and accumulated in simdgroup_half8x8 tiles, even
- *                 though the tensors are GGML_TYPE_F32 on both sides.
- *
- * (Read 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'
- *  and the runtime pick is visible with GGML_METAL_DEBUG-style logging as
- *  "compiling pipeline: base = 'kernel_mul_mm_f32_f32'".)
- *
- * The previous code issued ONE ggml_mul_mat with ne11 = nq (300 in the
- * benchmark), landing squarely on the F16 mul_mm path. Measured on this
- * machine (M4 Pro), n=13415 x dim=768 x nq=300, against a CPU double-
- * accumulated oracle:
- *
- *   ne11=300 (one mul_mat, the old code) : mean |Δdot| = 1.038e-05
- *   ne11=8   (chunked, this code)        : mean |Δdot| = 3.863e-09
- *
- * — a ~2700x reduction in dot-product error, which is exactly the gap that
- * showed up as 0.9933-vs-0.9997 recall.
- *
- * ggml_mul_mat_set_prec(t, GGML_PREC_F32) does NOT fix this. It was tried:
- * the error was bit-identical with and without it (1.038e-05 either way),
- * because ggml-metal only consults the prec flag on paths that have an F32
- * variant to switch to, and there is no F32-accumulating mul_mm kernel in
- * this build to select. The ONLY lever from outside ggml is ne11.
- *
- * So: instead of one mul_mat with ne11=nq, we emit ceil(nq/8) mul_mats, each
- * over an ne11<=8 ggml_view_2d slice of the same query tensor, all into ONE
- * graph and ONE ggml_backend_graph_compute. The node matrix is still uploaded
- * exactly once and still read by the GPU as one shared operand — the whole
- * point of batch_multi is preserved.
- *
- * The cost is real, and stated rather than buried. Timing the whole
- * batch_multi() call (gather + norms + upload + GPU + scatter) on the real
- * shape, median of 15 reps after a discarded warm-up, three separate runs:
- *
- *   unchunked (old, F16 mm) : 13.19 / 13.35 / 14.42 ms   -> ~0.044 ms/query
- *   chunked   (this code)   : 19.92 / 20.08 / 20.23 ms   -> ~0.067 ms/query
- *   hand-rolled Metal       : 17.74 / 17.88 / 17.99 ms   -> ~0.060 ms/query
- *
- * So correctness here costs about +6.7ms per 300-query batch (~1.5x on this
- * call), and leaves us ~12% behind the hand-rolled kernel instead of ~35%
- * ahead of it. That is not free and should not be sold as free. The reason it
- * cannot be recovered inside ggml: an fp32 matmul on Metal has to 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. Trading precision back for speed is a one-line env change; trading
- * the other way was not available before this commit at all.
- *
- * 8 is not a magic number we invented — it is ggml-metal's own mul_mm
- * threshold, measured by sweeping ne11 and watching both the error and which
- * pipeline ggml compiles (9 flips to mul_mm and the error jumps back to
- * 1.0e-05 in the same step). EL_GGML_MULMAT_CHUNK overrides it: raise it to
- * trade this precision back for throughput, or set it >= nq to reproduce the
- * old single-mul_mat behaviour exactly. If a future ggml moves the threshold,
- * the worst case is that we silently land back on mul_mm — the same accuracy
- * we shipped before, never a correctness break.
- *
- * ── Cold start: what is and is not ours to fix ───────────────────────────
- *
- * The ~7.8s first-call cost reported for the first version of this file is
- * NOT this file re-initialising per call (init is, and always was, cached
- * behind g_init_attempted below). It is Apple's Metal shader cache missing
- * on ggml's embedded metallib — ggml-metal ships ~650 kernels in one
- * __ggml_metallib section, and the first newLibraryWithData of it on a given
- * machine costs seconds ("ggml_metal_library_init: loaded in 7.670 sec")
- * while the driver populates ~/…/C/com.apple.metal/. That cache is keyed on
- * the library, not on our binary, and is shared across processes: the very
- * next run of a DIFFERENT binary linking the same ggml reports
- * "loaded in 0.009 sec". So it is a once-per-machine, per-ggml-version cost,
- * not a per-process one, 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.
- *
- * The residual warm init IS ours to look at, and the answer there is "there
- * was nothing much to win": ggml_backend_load_all_from_path() dlopens every
- * plugin in the directory (three CPU micro-arch variants + BLAS + Metal) when
- * we only ever use Metal, so we now load the single Metal plugin instead —
- * but measured warm that is 44.7-52.4ms against 46.9-58.9ms, i.e. the same
- * number inside noise, because libggml-metal.so's own init dominates. Warm
- * ggml init lands at 44-53ms, against 36-117ms for the hand-rolled strategy's
- * device+pipeline setup. Cold start was never the real defect here; precision
- * was.
- */
-#include "eg_cosine_batch_strategy.h"
-
-#include <ggml.h>
-#include <ggml-backend.h>
-#include <ggml-alloc.h>
-
-#include <stdint.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <math.h>
-
-/* ── lazy, one-time backend init, cached ─────────────────────────────────── */
-static bool g_init_attempted = false;
-static bool g_init_ok = false;
-static ggml_backend_t g_backend = NULL;
-
-/* Where to look for the dynamically-loaded backend plugin .so files.
- * EL_GGML_BACKEND_PATH overrides for non-standard installs; otherwise we try
- * the Homebrew opt-prefix symlink (stable across ggml point-version bumps —
- * $(brew --prefix ggml)/libexec — confirmed to exist and contain
- * libggml-metal.so / libggml-cpu-*.so / libggml-blas.so on this machine),
- * falling back to ggml's own default search (ggml_backend_load_all()) in
- * case a different install layout (e.g. a from-source build with a
- * standard-prefix install) makes that succeed instead. */
-static const char* eg_ggml_backend_dir(void) {
-    const char* s = getenv("EL_GGML_BACKEND_PATH");
-    if (s && *s) return s;
-    return "/opt/homebrew/opt/ggml/libexec";
-}
-
-/* "<dir>/libggml-metal.so" in a static buffer. Only ever called once, from
- * eg_ggml_ensure_init(), before any thread could race it. */
-static const char* eg_ggml_metal_plugin_path(const char* dir) {
-    static char buf[1024];
-    snprintf(buf, sizeof buf, "%s/libggml-metal.so", dir);
-    return buf;
-}
-
-/* Largest ne11 (query-batch rows per ggml_mul_mat) that keeps ggml-metal on
- * its F32 mul_mv kernels instead of the F16-accumulating mul_mm kernel — see
- * the precision discussion in this file's header. EL_GGML_MULMAT_CHUNK
- * overrides; a value <= 0 means "use the default". */
-#define EG_GGML_MULMAT_CHUNK_DEFAULT 8
-
-static int32_t eg_ggml_mulmat_chunk(void) {
-    static bool resolved = false;
-    static int32_t chunk = EG_GGML_MULMAT_CHUNK_DEFAULT;
-    if (!resolved) {
-        resolved = true;
-        const char* s = getenv("EL_GGML_MULMAT_CHUNK");
-        if (s && *s) {
-            long v = strtol(s, NULL, 10);
-            if (v > 0 && v <= INT32_MAX) chunk = (int32_t)v;
-        }
-    }
-    return chunk;
-}
-
-/* Which ggml device this strategy computes on. GPU (Metal) is the default
- * because offloading is the architectural point — the engram's own graph
- * traversal and activation spreading are CPU work, and a "GPU" strategy that
- * quietly saturates the CPU steals from them.
- *
- * ACCEL (ggml's BLAS/Accelerate plugin) is reachable here mainly as a
- * portability fallback and a diagnostic, and it is documented as MEASURED AND
- * REJECTED rather than as a recommendation. In an isolated probe that timed
- * only ggml_backend_graph_compute, BLAS looked excellent — 3.4-4.0ms for the
- * 300-query batch 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 through vindex_bench it does
- * not hold up: 0.191 ms/query at id-recall 0.9973, against 0.125-0.142 ms/query
- * at 0.9987 for the Metal default. It is dominated on BOTH axes, because the
- * isolated probe was not competing with the rest of the batch for the same CPU
- * cores and the real call path is. Kept because a machine with no usable Metal
- * device still wants a working ggml strategy — not because it is faster. */
-static enum ggml_backend_dev_type eg_ggml_device_type(void) {
-    const char* s = getenv("EL_GGML_DEVICE");
-    if (s && *s) {
-        if (strcmp(s, "accel") == 0) return GGML_BACKEND_DEVICE_TYPE_ACCEL;
-        if (strcmp(s, "cpu")   == 0) return GGML_BACKEND_DEVICE_TYPE_CPU;
-    }
-    return GGML_BACKEND_DEVICE_TYPE_GPU;
-}
-
-static bool eg_ggml_ensure_init(void) {
-    if (g_init_attempted) return g_init_ok;
-    g_init_attempted = true;
-
-    const char* dir = eg_ggml_backend_dir();
-    const enum ggml_backend_dev_type want = eg_ggml_device_type();
-
-    /* Metal is the only backend this strategy uses by default, so load just
-     * that one plugin rather than dlopening the whole directory (three CPU
-     * micro-arch variants + BLAS + Metal here).
-     *
-     * Be honest about what this buys: almost nothing in wall time. Measured
-     * warm, three runs each — load-everything 58.9/46.9/55.1ms, Metal-only
-     * 52.4/51.2/44.7ms. The cost is dominated by dlopening and initialising
-     * libggml-metal.so itself, not by the four plugins we skip, so the two
-     * overlap inside noise. It is kept because registering four device types
-     * we will never dispatch to is untidy and makes ggml_backend_dev_by_type
-     * ambiguous, not because it is a speedup — do not cite it as one.
-     *
-     * ggml_backend_load() returns NULL for a missing or unloadable path,
-     * which simply falls through to the broader searches below; it is never
-     * fatal. Any non-default device needs the full directory scan to find
-     * its plugin, so skip the fast path there. */
-    if (want == GGML_BACKEND_DEVICE_TYPE_GPU)
-        ggml_backend_load(eg_ggml_metal_plugin_path(dir));
-
-    ggml_backend_dev_t dev = ggml_backend_dev_by_type(want);
-    if (!dev) {
-        /* Non-standard layout, a ggml built with a differently-named Metal
-         * plugin, or a non-default device: dlopen every plugin in `dir`. */
-        ggml_backend_load_all_from_path(dir);
-        dev = ggml_backend_dev_by_type(want);
-    }
-    if (!dev) {
-        /* Fall back to ggml's own default search heuristics only if the
-         * explicit path above found nothing — avoids double-registering the
-         * same plugins (ggml does not dedupe two different paths that
-         * happen to resolve to the same files, e.g. our stable opt-prefix
-         * symlink vs. its own Cellar-relative guess) in the common case
-         * where the explicit path already worked. */
-        ggml_backend_load_all();
-        dev = ggml_backend_dev_by_type(want);
-    }
-    if (!dev) return false;
-
-    ggml_backend_t backend = ggml_backend_dev_init(dev, NULL);
-    if (!backend) return false;
-
-    g_backend = backend;
-    g_init_ok = true;
-    return true;
-}
-
-static bool ggml_strategy_available(void) {
-    return eg_ggml_ensure_init();
-}
-
-/* ── shared core: gather valid rows + norms, matmul, scatter ────────────── */
-
-/* 4-way partial-sum squared-norm accumulation over `dim` floats — same shape
- * as eg_cosine_batch.metal's per-thread accumulation and vindex_bench.c's
- * CPU brute_topk unroll, kept consistent on purpose so the float32 error
- * profile is comparable across all three strategies. */
-static float eg_norm_sq_f32(const float* v, int32_t dim) {
-    float s0 = 0, s1 = 0, s2 = 0, s3 = 0;
-    int32_t d = 0, dim4 = dim & ~3;
-    for (; d < dim4; d += 4) {
-        s0 += v[d] * v[d]; s1 += v[d+1] * v[d+1];
-        s2 += v[d+2] * v[d+2]; s3 += v[d+3] * v[d+3];
-    }
-    float s = (s0 + s1) + (s2 + s3);
-    for (; d < dim; d++) s += v[d] * v[d];
-    return s;
-}
-
-/* Runs one ggml_mul_mat(node_matrix[dim,n_valid], query_matrix[dim,nq]) and
- * combines it with CPU-computed norms into cosine scores, scattering into
- * out_scores at ORIGINAL (ungathered) indices. out_scores must already be
- * fully sized for n*nq (or n for the single-query case, nq=1) — every entry
- * gets written (valid rows get a real cosine, invalid rows get -2.0), so
- * this never leaves a partial result. Returns false only on a genuine
- * failure (alloc, compute) at which point out_scores is left as whatever a
- * caller-supplied scratch buffer already contained — callers here always
- * pass a fresh buffer they discard on false, matching the adapter contract
- * of "on failure, out_scores is treated as untouched" from the caller's
- * point of view. */
-static bool eg_ggml_run(const float* queries, int32_t qdim, int32_t nq,
-                         const float* const* node_ptrs, const int32_t* node_dims,
-                         int32_t n, double* out_scores) {
-    if (!queries || qdim <= 0 || nq <= 0 || !node_ptrs || !node_dims || n <= 0 || !out_scores)
-        return false;
-    if (!eg_ggml_ensure_init()) return false;
-
-    /* Pass 1 (CPU): gather valid rows (non-NULL ptr, dim == qdim) into a
-     * packed (dim, n_valid) row-major matrix, remembering the original index
-     * of each packed row, and compute each valid row's squared norm in the
-     * same pass. Rows excluded here get -2.0 scattered for every query
-     * below without ever touching the GPU. */
-    int32_t* valid_orig = (int32_t*)malloc((size_t)n * sizeof(int32_t));
-    float* node_norm_sq  = (float*)malloc((size_t)n * sizeof(float)); /* indexed by packed position */
-    float* node_matrix   = NULL;
-    if (!valid_orig || !node_norm_sq) { free(valid_orig); free(node_norm_sq); return false; }
-
-    int32_t n_valid = 0;
-    for (int32_t i = 0; i < n; i++) {
-        if (node_ptrs[i] && node_dims[i] == qdim) n_valid++;
-    }
-
-    if (n_valid > 0) {
-        node_matrix = (float*)malloc((size_t)n_valid * (size_t)qdim * sizeof(float));
-        if (!node_matrix) { free(valid_orig); free(node_norm_sq); return false; }
-        int32_t w = 0;
-        for (int32_t i = 0; i < n; i++) {
-            if (!node_ptrs[i] || node_dims[i] != qdim) continue;
-            memcpy(node_matrix + (size_t)w * qdim, node_ptrs[i], (size_t)qdim * sizeof(float));
-            node_norm_sq[w] = eg_norm_sq_f32(node_ptrs[i], qdim);
-            valid_orig[w] = i;
-            w++;
-        }
-    }
-
-    /* Query norms — nq is typically small (1 or the size of one batch of
-     * comparison queries), so this loop is cheap regardless. */
-    float* q_norm_sq = (float*)malloc((size_t)nq * sizeof(float));
-    if (!q_norm_sq) { free(valid_orig); free(node_norm_sq); free(node_matrix); return false; }
-    for (int32_t j = 0; j < nq; j++) q_norm_sq[j] = eg_norm_sq_f32(queries + (size_t)j * qdim, qdim);
-
-    /* Nothing valid to compare against: every output is -2.0. Still a fully
-     * and correctly populated result — no GPU dispatch was needed to know
-     * that. */
-    if (n_valid == 0) {
-        for (size_t k = 0; k < (size_t)n * (size_t)nq; k++) out_scores[k] = -2.0;
-        free(valid_orig); free(node_norm_sq); free(node_matrix); free(q_norm_sq);
-        return true;
-    }
-
-    /* Pass 2 (GPU via ggml): dot[j*n_valid + i] = dot(node_i, query_j),
-     * computed as ceil(nq/chunk) separate ggml_mul_mat ops over ne11<=chunk
-     * ggml_view_2d slices of ONE query tensor, all expanded into ONE graph
-     * and run by ONE ggml_backend_graph_compute. Chunking is what keeps
-     * ggml-metal on its F32 mul_mv kernels rather than the F16-accumulating
-     * mul_mm kernel (see this file's header); sharing one graph and one
-     * t_nodes tensor is what keeps the node matrix uploaded exactly once,
-     * which is the entire reason batch_multi exists. */
-    const int32_t chunk = eg_ggml_mulmat_chunk();
-    const int32_t ngroups = (nq + chunk - 1) / chunk;
-
-    /* Tensors held by the context: t_nodes, t_query, plus one view and one
-     * mul_mat result per group. The graph holds at most one node per view and
-     * one per mul_mat. Slack on both so a ggml that bookkeeps slightly
-     * differently cannot silently overflow the arena. */
-    const size_t n_tensors  = (size_t)2 * (size_t)ngroups + 8;
-    const size_t graph_size = (size_t)2 * (size_t)ngroups + 16;
-    struct ggml_init_params gp = {
-        .mem_size   = ggml_tensor_overhead() * n_tensors
-                    + ggml_graph_overhead_custom(graph_size, false),
-        .mem_buffer = NULL,
-        .no_alloc   = true,
-    };
-    struct ggml_context* ctx = ggml_init(gp);
-    if (!ctx) { free(valid_orig); free(node_norm_sq); free(node_matrix); free(q_norm_sq); return false; }
-
-    struct ggml_tensor** t_dots = (struct ggml_tensor**)malloc((size_t)ngroups * sizeof(*t_dots));
-    if (!t_dots) { ggml_free(ctx); free(valid_orig); free(node_norm_sq); free(node_matrix); free(q_norm_sq); return false; }
-
-    struct ggml_tensor* t_nodes = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, qdim, n_valid);
-    struct ggml_tensor* t_query = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, qdim, nq);
-    struct ggml_cgraph*  gf     = t_nodes && t_query
-                                ? ggml_new_graph_custom(ctx, graph_size, false) : NULL;
-    if (!gf) { free(t_dots); ggml_free(ctx); free(valid_orig); free(node_norm_sq); free(node_matrix); free(q_norm_sq); return false; }
-
-    bool built = true;
-    for (int32_t g = 0; g < ngroups; g++) {
-        const int32_t start = g * chunk;
-        const int32_t count = (start + chunk <= nq) ? chunk : (nq - start);
-        struct ggml_tensor* t_qv = ggml_view_2d(ctx, t_query, qdim, count,
-                                                t_query->nb[1],
-                                                (size_t)start * t_query->nb[1]);
-        t_dots[g] = t_qv ? ggml_mul_mat(ctx, t_nodes, t_qv) : NULL;
-        if (!t_dots[g]) { built = false; break; }
-        ggml_build_forward_expand(gf, t_dots[g]);
-    }
-    if (!built) { free(t_dots); ggml_free(ctx); free(valid_orig); free(node_norm_sq); free(node_matrix); free(q_norm_sq); return false; }
-
-    struct ggml_backend_buffer* buf = ggml_backend_alloc_ctx_tensors(ctx, g_backend);
-    if (!buf) { free(t_dots); ggml_free(ctx); free(valid_orig); free(node_norm_sq); free(node_matrix); free(q_norm_sq); return false; }
-
-    ggml_backend_tensor_set(t_nodes, node_matrix, 0, (size_t)n_valid * qdim * sizeof(float));
-    ggml_backend_tensor_set(t_query, queries, 0, (size_t)nq * qdim * sizeof(float));
-    free(node_matrix); /* uploaded; the packed CPU copy is no longer needed */
-
-    enum ggml_status st = ggml_backend_graph_compute(g_backend, gf);
-    if (st != GGML_STATUS_SUCCESS) {
-        free(t_dots); ggml_backend_buffer_free(buf); ggml_free(ctx);
-        free(valid_orig); free(node_norm_sq); free(q_norm_sq);
-        return false;
-    }
-
-    /* Each group's result is [n_valid, count] contiguous, so reading group g
-     * into dot + start*n_valid reconstructs exactly the same flat
-     * dot[j*n_valid + w] layout a single ne11=nq mul_mat would have produced —
-     * Pass 3 below is unchanged by the chunking. */
-    float* dot = (float*)malloc((size_t)n_valid * (size_t)nq * sizeof(float));
-    if (!dot) { free(t_dots); ggml_backend_buffer_free(buf); ggml_free(ctx); free(valid_orig); free(node_norm_sq); free(q_norm_sq); return false; }
-    for (int32_t g = 0; g < ngroups; g++) {
-        const int32_t start = g * chunk;
-        const int32_t count = (start + chunk <= nq) ? chunk : (nq - start);
-        ggml_backend_tensor_get(t_dots[g], dot + (size_t)start * n_valid, 0,
-                                (size_t)count * (size_t)n_valid * sizeof(float));
-    }
-    free(t_dots);
-
-    /* Pass 3 (CPU): combine dot/(||a||*||b||) per (query,node) pair, scatter
-     * into out_scores at ORIGINAL node indices; every excluded row gets
-     * -2.0 for every query. out_scores is fully populated either way. */
-    for (int32_t j = 0; j < nq; j++) {
-        double* orow = out_scores + (size_t)j * n;
-        for (int32_t i = 0; i < n; i++) orow[i] = -2.0; /* default: excluded */
-        for (int32_t w = 0; w < n_valid; w++) {
-            float na = node_norm_sq[w], nb = q_norm_sq[j];
-            int32_t oi = valid_orig[w];
-            if (na <= 0.0f || nb <= 0.0f) { orow[oi] = -2.0; continue; }
-            float d = dot[(size_t)j * n_valid + w];
-            orow[oi] = (double)(d / sqrtf(na * nb));
-        }
-    }
-
-    free(dot);
-    ggml_backend_buffer_free(buf);
-    ggml_free(ctx);
-    free(valid_orig); free(node_norm_sq); free(q_norm_sq);
-    return true;
-}
-
-static bool ggml_strategy_batch(const float* query, int32_t qdim,
-                                 const float* const* node_ptrs, const int32_t* node_dims,
-                                 int32_t n, double* out_scores) {
-    if (!query || qdim <= 0 || !node_ptrs || !node_dims || n <= 0 || !out_scores) return false;
-    /* out_scores here is n doubles (nq=1); eg_ggml_run writes n*nq = n of
-     * them, laid out identically to the single-query contract. */
-    return eg_ggml_run(query, qdim, 1, node_ptrs, node_dims, n, out_scores);
-}
-
-static bool ggml_strategy_batch_multi(const float* queries, int32_t qdim, int32_t nq,
-                                       const float* const* node_ptrs, const int32_t* node_dims,
-                                       int32_t n, double* out_scores) {
-    return eg_ggml_run(queries, qdim, nq, node_ptrs, node_dims, n, out_scores);
-}
-
-static const EgCosineBatchStrategy g_ggml_strategy = {
-    .name = "ggml",
-    .available = ggml_strategy_available,
-    .batch = ggml_strategy_batch,
-    .batch_multi = ggml_strategy_batch_multi,
-};
-
-const EgCosineBatchStrategy* eg_cosine_batch_strategy_ggml(void) {
-    return &g_ggml_strategy;
-}

lang/runtime/eg_cosine_batch_strategy_metal_hand.m

@@ -1,358 +0,0 @@
-/* eg_cosine_batch_strategy_metal_hand.m — the HAND-ROLLED-METAL Strategy.
- *
- * This is PR #114's original Objective-C bridge (formerly eg_metal_cosine.m)
- * exposing the hand-written Metal compute shader (eg_cosine_batch.metal) as
- * one concrete EgCosineBatchStrategy. It is preserved here almost verbatim —
- * real, carefully verified work, not discarded — now living behind the
- * Adapter/Strategy/Factory restructuring (see eg_cosine_batch.h and
- * eg_cosine_batch_strategy.h) alongside the new ggml-Metal strategy
- * (eg_cosine_batch_strategy_ggml.c) and the universal CPU fallback
- * (eg_cosine_batch_strategy_cpu.c). The factory in eg_cosine_batch.c prefers
- * ggml by default when both are available; this strategy remains selectable
- * via EL_COSINE_BATCH_STRATEGY=metal, and is what the factory falls back to
- * if ggml's backend plugin fails to load/init for any reason.
- *
- * Apple-only (Metal has no other platform). This file is excluded from the
- * build entirely on non-Darwin — see build_vindex_bench.sh, which only
- * compiles/links this file and defines EG_HAVE_STRATEGY_METAL_HAND when
- * `uname` is Darwin. On Linux the factory never sees this strategy at all —
- * callers must always be prepared for the "no real strategy available"
- * fallback via the CPU strategy, which is also exactly what happens here on
- * Apple hardware with no usable GPU.
- *
- * Design (unchanged from PR #114):
- *   - Device/queue/pipeline are created lazily, once, and cached in static
- *     globals — every call after the first only allocates buffers + submits.
- *   - The Metal shader source is embedded as a C string literal (kMetalSrc
- *     below) rather than loaded from a file at runtime or shipped as a
- *     precompiled .metallib. Chosen over newLibraryWithFile: /a .metallib
- *     because the engram binary can be invoked from an arbitrary working
- *     directory (launchd job, nsbx sandbox, CI) and a file-path shader would
- *     be one relocation away from silently falling back to CPU for reasons
- *     that have nothing to do with Metal availability. Embedding costs one
- *     runtime shader compile (~tens of ms) on first use, amortized over the
- *     process lifetime, in exchange for a genuinely self-contained binary.
- *     Source of truth for review/tooling is eg_cosine_batch.metal — this
- *     string MUST be kept byte-identical to that file (a comment marks both
- *     ends of the copy).
- *   - Buffers use MTLResourceStorageModeShared: on Apple Silicon's unified
- *     memory, CPU and GPU read the same physical pages, so filling a buffer
- *     is a plain memcpy and there is no separate "upload" step.
- *   - ANY failure at ANY step (no device, pipeline compile error, buffer
- *     allocation failure, bad args) returns false and leaves out_scores
- *     untouched. This function is called from the request-handling hot path
- *     of a long-lived daemon — it must never throw, crash, or hang it.
- */
-#import <Foundation/Foundation.h>
-#import <Metal/Metal.h>
-#include "eg_cosine_batch_strategy.h"
-#include <string.h>
-#include <stdlib.h>
-
-/* ── BEGIN embedded shader source (keep in sync with eg_cosine_batch.metal) ── */
-static const char* kEgCosineBatchMetalSrc =
-"#include <metal_stdlib>\n"
-"using namespace metal;\n"
-"struct EgCosineParams { uint n; uint dim; };\n"
-"kernel void eg_cosine_batch_kernel(\n"
-"    device const float*  query        [[buffer(0)]],\n"
-"    device const float*  node_matrix  [[buffer(1)]],\n"
-"    device const int*    node_dims    [[buffer(2)]],\n"
-"    constant EgCosineParams& p        [[buffer(3)]],\n"
-"    device float*        out_scores   [[buffer(4)]],\n"
-"    uint gid [[thread_position_in_grid]])\n"
-"{\n"
-"    if (gid >= p.n) return;\n"
-"    if (node_dims[gid] != int(p.dim)) { out_scores[gid] = -2.0f; return; }\n"
-"    device const float* row = node_matrix + (uint64_t)gid * (uint64_t)p.dim;\n"
-"    float dot0 = 0.0f, dot1 = 0.0f, dot2 = 0.0f, dot3 = 0.0f;\n"
-"    float na0  = 0.0f, na1  = 0.0f, na2  = 0.0f, na3  = 0.0f;\n"
-"    float nb0  = 0.0f, nb1  = 0.0f, nb2  = 0.0f, nb3  = 0.0f;\n"
-"    uint d = 0;\n"
-"    uint dim4 = p.dim & ~3u;\n"
-"    for (; d < dim4; d += 4) {\n"
-"        float a0 = row[d],   b0 = query[d];\n"
-"        float a1 = row[d+1], b1 = query[d+1];\n"
-"        float a2 = row[d+2], b2 = query[d+2];\n"
-"        float a3 = row[d+3], b3 = query[d+3];\n"
-"        dot0 += a0*b0; dot1 += a1*b1; dot2 += a2*b2; dot3 += a3*b3;\n"
-"        na0  += a0*a0; na1  += a1*a1; na2  += a2*a2; na3  += a3*a3;\n"
-"        nb0  += b0*b0; nb1  += b1*b1; nb2  += b2*b2; nb3  += b3*b3;\n"
-"    }\n"
-"    float dot = (dot0 + dot1) + (dot2 + dot3);\n"
-"    float na  = (na0  + na1)  + (na2  + na3);\n"
-"    float nb  = (nb0  + nb1)  + (nb2  + nb3);\n"
-"    for (; d < p.dim; d++) {\n"
-"        float a = row[d], b = query[d];\n"
-"        dot += a*b; na += a*a; nb += b*b;\n"
-"    }\n"
-"    if (na <= 0.0f || nb <= 0.0f) { out_scores[gid] = -2.0f; return; }\n"
-"    out_scores[gid] = dot / sqrt(na * nb);\n"
-"}\n"
-"struct EgCosineMultiParams { uint n; uint dim; uint nq; };\n"
-"kernel void eg_cosine_batch_multi_kernel(\n"
-"    device const float*  queries      [[buffer(0)]],\n"
-"    device const float*  node_matrix  [[buffer(1)]],\n"
-"    device const int*    node_dims    [[buffer(2)]],\n"
-"    constant EgCosineMultiParams& p   [[buffer(3)]],\n"
-"    device float*        out_scores   [[buffer(4)]],\n"
-"    uint2 gid [[thread_position_in_grid]])\n"
-"{\n"
-"    uint nid = gid.x, qid = gid.y;\n"
-"    if (nid >= p.n || qid >= p.nq) return;\n"
-"    uint64_t out_idx = (uint64_t)qid * (uint64_t)p.n + (uint64_t)nid;\n"
-"    if (node_dims[nid] != int(p.dim)) { out_scores[out_idx] = -2.0f; return; }\n"
-"    device const float* row   = node_matrix + (uint64_t)nid * (uint64_t)p.dim;\n"
-"    device const float* query = queries     + (uint64_t)qid * (uint64_t)p.dim;\n"
-"    float dot0 = 0.0f, dot1 = 0.0f, dot2 = 0.0f, dot3 = 0.0f;\n"
-"    float na0  = 0.0f, na1  = 0.0f, na2  = 0.0f, na3  = 0.0f;\n"
-"    float nb0  = 0.0f, nb1  = 0.0f, nb2  = 0.0f, nb3  = 0.0f;\n"
-"    uint d = 0;\n"
-"    uint dim4 = p.dim & ~3u;\n"
-"    for (; d < dim4; d += 4) {\n"
-"        float a0 = row[d],   b0 = query[d];\n"
-"        float a1 = row[d+1], b1 = query[d+1];\n"
-"        float a2 = row[d+2], b2 = query[d+2];\n"
-"        float a3 = row[d+3], b3 = query[d+3];\n"
-"        dot0 += a0*b0; dot1 += a1*b1; dot2 += a2*b2; dot3 += a3*b3;\n"
-"        na0  += a0*a0; na1  += a1*a1; na2  += a2*a2; na3  += a3*a3;\n"
-"        nb0  += b0*b0; nb1  += b1*b1; nb2  += b2*b2; nb3  += b3*b3;\n"
-"    }\n"
-"    float dot = (dot0 + dot1) + (dot2 + dot3);\n"
-"    float na  = (na0  + na1)  + (na2  + na3);\n"
-"    float nb  = (nb0  + nb1)  + (nb2  + nb3);\n"
-"    for (; d < p.dim; d++) {\n"
-"        float a = row[d], b = query[d];\n"
-"        dot += a*b; na += a*a; nb += b*b;\n"
-"    }\n"
-"    if (na <= 0.0f || nb <= 0.0f) { out_scores[out_idx] = -2.0f; return; }\n"
-"    out_scores[out_idx] = dot / sqrt(na * nb);\n"
-"}\n";
-/* ── END embedded shader source ── */
-
-typedef struct EgCosineParamsC { uint32_t n; uint32_t dim; } EgCosineParamsC;
-typedef struct EgCosineMultiParamsC { uint32_t n; uint32_t dim; uint32_t nq; } EgCosineMultiParamsC;
-
-static id<MTLDevice>              g_device   = nil;
-static id<MTLCommandQueue>        g_queue    = nil;
-static id<MTLComputePipelineState> g_pipeline = nil;       /* single-query kernel */
-static id<MTLComputePipelineState> g_pipeline_multi = nil; /* multi-query kernel */
-static bool g_init_attempted = false;
-static bool g_init_ok        = false;
-
-/* Lazy, one-time setup. Never throws — every Metal call here is the
- * "returns nil/NSError on failure" flavor, not an exception-throwing one. */
-static bool eg_metal_ensure_init(void) {
-    if (g_init_attempted) return g_init_ok;
-    g_init_attempted = true;
-
-    @autoreleasepool {
-        id<MTLDevice> dev = MTLCreateSystemDefaultDevice();
-        if (!dev) return false;
-
-        id<MTLCommandQueue> q = [dev newCommandQueue];
-        if (!q) return false;
-
-        NSError* err = nil;
-        NSString* src = [NSString stringWithUTF8String:kEgCosineBatchMetalSrc];
-        MTLCompileOptions* opts = [MTLCompileOptions new];
-        id<MTLLibrary> lib = [dev newLibraryWithSource:src options:opts error:&err];
-        if (!lib) return false;
-
-        id<MTLFunction> fn = [lib newFunctionWithName:@"eg_cosine_batch_kernel"];
-        if (!fn) return false;
-        id<MTLComputePipelineState> pipe = [dev newComputePipelineStateWithFunction:fn error:&err];
-        if (!pipe) return false;
-
-        id<MTLFunction> fnMulti = [lib newFunctionWithName:@"eg_cosine_batch_multi_kernel"];
-        if (!fnMulti) return false;
-        id<MTLComputePipelineState> pipeMulti = [dev newComputePipelineStateWithFunction:fnMulti error:&err];
-        if (!pipeMulti) return false;
-
-        g_device = dev;
-        g_queue = q;
-        g_pipeline = pipe;
-        g_pipeline_multi = pipeMulti;
-        g_init_ok = true;
-        return true;
-    }
-}
-
-static bool mh_available(void) {
-    return eg_metal_ensure_init();
-}
-
-static bool mh_batch(const float* query, int32_t qdim,
-                      const float* const* node_ptrs,
-                      const int32_t* node_dims,
-                      int32_t n,
-                      double* out_scores) {
-    if (!query || qdim <= 0 || !node_ptrs || !node_dims || n <= 0 || !out_scores) return false;
-    if (!eg_metal_ensure_init()) return false;
-
-    @autoreleasepool {
-        const size_t dim = (size_t)qdim;
-        const size_t nu  = (size_t)n;
-
-        /* Gather into a packed row-major matrix — EngramNode.emb is one
-         * malloc per node, not a contiguous array, so this copy is
-         * unavoidable regardless of backend. Rows whose real dim doesn't
-         * match qdim are zero-filled (harmless: the kernel sentinels them
-         * via node_dims before ever reading the row). */
-        float* matrix = (float*)calloc(nu * dim, sizeof(float));
-        int32_t* dims_i32 = (int32_t*)malloc(nu * sizeof(int32_t));
-        if (!matrix || !dims_i32) { free(matrix); free(dims_i32); return false; }
-
-        for (size_t i = 0; i < nu; i++) {
-            dims_i32[i] = node_dims[i];
-            if (node_ptrs[i] && node_dims[i] == qdim) {
-                memcpy(matrix + i * dim, node_ptrs[i], dim * sizeof(float));
-            }
-            /* else: leave zero-filled; node_dims[i] != qdim (or missing)
-             * makes the kernel sentinel it to -2.0 without reading the row. */
-        }
-
-        id<MTLBuffer> bufQuery = [g_device newBufferWithBytes:query
-                                                        length:dim * sizeof(float)
-                                                       options:MTLResourceStorageModeShared];
-        id<MTLBuffer> bufMatrix = [g_device newBufferWithBytes:matrix
-                                                         length:nu * dim * sizeof(float)
-                                                        options:MTLResourceStorageModeShared];
-        id<MTLBuffer> bufDims = [g_device newBufferWithBytes:dims_i32
-                                                       length:nu * sizeof(int32_t)
-                                                      options:MTLResourceStorageModeShared];
-        EgCosineParamsC params = { (uint32_t)nu, (uint32_t)dim };
-        id<MTLBuffer> bufParams = [g_device newBufferWithBytes:&params
-                                                         length:sizeof(params)
-                                                        options:MTLResourceStorageModeShared];
-        id<MTLBuffer> bufOut = [g_device newBufferWithLength:nu * sizeof(float)
-                                                       options:MTLResourceStorageModeShared];
-
-        free(matrix); free(dims_i32);
-
-        if (!bufQuery || !bufMatrix || !bufDims || !bufParams || !bufOut) return false;
-
-        id<MTLCommandBuffer> cmd = [g_queue commandBuffer];
-        if (!cmd) return false;
-        id<MTLComputeCommandEncoder> enc = [cmd computeCommandEncoder];
-        if (!enc) return false;
-
-        [enc setComputePipelineState:g_pipeline];
-        [enc setBuffer:bufQuery  offset:0 atIndex:0];
-        [enc setBuffer:bufMatrix offset:0 atIndex:1];
-        [enc setBuffer:bufDims   offset:0 atIndex:2];
-        [enc setBuffer:bufParams offset:0 atIndex:3];
-        [enc setBuffer:bufOut    offset:0 atIndex:4];
-
-        NSUInteger tgSize = g_pipeline.maxTotalThreadsPerThreadgroup;
-        if (tgSize > 256) tgSize = 256;
-        if (tgSize < 1) tgSize = 1;
-        MTLSize gridSize = MTLSizeMake(nu, 1, 1);
-        MTLSize threadgroupSize = MTLSizeMake(tgSize, 1, 1);
-        [enc dispatchThreads:gridSize threadsPerThreadgroup:threadgroupSize];
-        [enc endEncoding];
-
-        [cmd commit];
-        [cmd waitUntilCompleted];
-
-        if (cmd.status != MTLCommandBufferStatusCompleted) return false;
-
-        const float* results = (const float*)bufOut.contents;
-        if (!results) return false;
-        for (size_t i = 0; i < nu; i++) out_scores[i] = (double)results[i];
-        return true;
-    }
-}
-
-static bool mh_batch_multi(const float* queries, int32_t qdim, int32_t nq,
-                            const float* const* node_ptrs,
-                            const int32_t* node_dims,
-                            int32_t n,
-                            double* out_scores) {
-    if (!queries || qdim <= 0 || nq <= 0 || !node_ptrs || !node_dims || n <= 0 || !out_scores) return false;
-    if (!eg_metal_ensure_init()) return false;
-
-    @autoreleasepool {
-        const size_t dim = (size_t)qdim;
-        const size_t nu  = (size_t)n;
-        const size_t nqu = (size_t)nq;
-
-        float* matrix = (float*)calloc(nu * dim, sizeof(float));
-        int32_t* dims_i32 = (int32_t*)malloc(nu * sizeof(int32_t));
-        if (!matrix || !dims_i32) { free(matrix); free(dims_i32); return false; }
-
-        for (size_t i = 0; i < nu; i++) {
-            dims_i32[i] = node_dims[i];
-            if (node_ptrs[i] && node_dims[i] == qdim) {
-                memcpy(matrix + i * dim, node_ptrs[i], dim * sizeof(float));
-            }
-        }
-
-        /* This is the ONE upload of node_matrix for the whole nq-query batch —
-         * the fix for the measured re-upload-per-query slowdown. */
-        id<MTLBuffer> bufMatrix = [g_device newBufferWithBytes:matrix
-                                                         length:nu * dim * sizeof(float)
-                                                        options:MTLResourceStorageModeShared];
-        id<MTLBuffer> bufDims = [g_device newBufferWithBytes:dims_i32
-                                                       length:nu * sizeof(int32_t)
-                                                      options:MTLResourceStorageModeShared];
-        id<MTLBuffer> bufQueries = [g_device newBufferWithBytes:queries
-                                                          length:nqu * dim * sizeof(float)
-                                                         options:MTLResourceStorageModeShared];
-        EgCosineMultiParamsC params = { (uint32_t)nu, (uint32_t)dim, (uint32_t)nqu };
-        id<MTLBuffer> bufParams = [g_device newBufferWithBytes:&params
-                                                         length:sizeof(params)
-                                                        options:MTLResourceStorageModeShared];
-        id<MTLBuffer> bufOut = [g_device newBufferWithLength:nqu * nu * sizeof(float)
-                                                       options:MTLResourceStorageModeShared];
-
-        free(matrix); free(dims_i32);
-
-        if (!bufMatrix || !bufDims || !bufQueries || !bufParams || !bufOut) return false;
-
-        id<MTLCommandBuffer> cmd = [g_queue commandBuffer];
-        if (!cmd) return false;
-        id<MTLComputeCommandEncoder> enc = [cmd computeCommandEncoder];
-        if (!enc) return false;
-
-        [enc setComputePipelineState:g_pipeline_multi];
-        [enc setBuffer:bufQueries offset:0 atIndex:0];
-        [enc setBuffer:bufMatrix  offset:0 atIndex:1];
-        [enc setBuffer:bufDims    offset:0 atIndex:2];
-        [enc setBuffer:bufParams  offset:0 atIndex:3];
-        [enc setBuffer:bufOut     offset:0 atIndex:4];
-
-        /* 2D dispatch: x over nodes, y over queries. Threadgroup width picked
-         * from the pipeline's own limit, height fixed at 1 — nq is typically
-         * small (tens to low hundreds) relative to n (thousands+), so tiling
-         * the wide axis (n) is what matters for occupancy. */
-        NSUInteger tgWidth = g_pipeline_multi.maxTotalThreadsPerThreadgroup;
-        if (tgWidth > 256) tgWidth = 256;
-        if (tgWidth < 1) tgWidth = 1;
-        MTLSize gridSize = MTLSizeMake(nu, nqu, 1);
-        MTLSize threadgroupSize = MTLSizeMake(tgWidth, 1, 1);
-        [enc dispatchThreads:gridSize threadsPerThreadgroup:threadgroupSize];
-        [enc endEncoding];
-
-        [cmd commit];
-        [cmd waitUntilCompleted];
-
-        if (cmd.status != MTLCommandBufferStatusCompleted) return false;
-
-        const float* results = (const float*)bufOut.contents;
-        if (!results) return false;
-        for (size_t i = 0; i < nqu * nu; i++) out_scores[i] = (double)results[i];
-        return true;
-    }
-}
-
-static const EgCosineBatchStrategy g_metal_hand_strategy = {
-    .name = "metal-hand",
-    .available = mh_available,
-    .batch = mh_batch,
-    .batch_multi = mh_batch_multi,
-};
-
-const EgCosineBatchStrategy* eg_cosine_batch_strategy_metal_hand(void) {
-    return &g_metal_hand_strategy;
-}

lang/runtime/el_runtime.c

@@ -2226,6 +2226,22 @@ el_val_t fs_exists(el_val_t pathv) {
     return (el_val_t)(stat(path, &st) == 0 ? 1 : 0);
 }
 
+/* fs_size — real on-disk byte count of a file, via stat() (not strlen).
+ * Needed alongside fs_read_b64_chunk() below: fs_read()'s el_val_t string
+ * result is NUL-terminated and length-unsafe for arbitrary binary content
+ * (str_len/str_slice fall back to strlen when the buffer isn't a tagged
+ * binary value — see el_input_len), so any caller that needs to walk a
+ * binary file in fixed-size windows (e.g. transduce()'s raw/opaque chunking
+ * of audio bytes) must learn the true length here instead of from the
+ * decoded string. Returns -1 if the path doesn't exist or isn't stat-able. */
+el_val_t fs_size(el_val_t pathv) {
+    const char* path = EL_CSTR(pathv);
+    if (!path || !*path) return -1;
+    struct stat st;
+    if (stat(path, &st) != 0) return -1;
+    return (el_val_t)st.st_size;
+}
+
 /* fs_mkdir — create directory at path with mode 0755, mkdir -p semantics.
  * Returns 1 if path exists or was created (incl. all parents); 0 on failure.
  * Walks the path component-by-component so missing intermediate dirs are
@@ -6847,10 +6863,16 @@ static float* eg_embed_fetch(const char* text, int32_t* out_dim) {
         else esc[w++] = (char)c;
     }
     esc[w] = '\0';
-    size_t blen = w + strlen(eg_embed_model()) + 64;
+    size_t blen = w + strlen(eg_embed_model()) + 96;
     char* body = malloc(blen);
     if (!body) { free(esc); return NULL; }
-    snprintf(body, blen, "{\"model\":\"%s\",\"prompt\":\"%s\"}",
+    /* keep_alive:-1 pins the embed model resident in Ollama indefinitely
+     * (2026-08-15, PR #105 port). Without it the tiny embed model is evicted
+     * whenever a larger generation model loads (unified-memory pressure), so
+     * the NEXT search pays a cold model reload — measured cold reload up to
+     * ~2.2s vs ~0.02-0.05s warm, well inside ENGRAM_EMBED_TIMEOUT_MS but a
+     * real tax on every activate() call that lands cold. Pinning removes it. */
+    snprintf(body, blen, "{\"model\":\"%s\",\"keep_alive\":-1,\"prompt\":\"%s\"}",
              eg_embed_model(), esc);
     free(esc);
     struct curl_slist* h = curl_slist_append(NULL, "Content-Type: application/json");
@@ -9508,6 +9530,36 @@ static inline double eg_cosq_at(EngramStore* g, double* cosq, unsigned char* cos
     return cosq[i];
 }
 
+/* ── Beam cap for engram_activate spreading activation (2026-08-15, PR #105
+ * port) ──────────────────────────────────────────────────────────────────
+ * #105 measured the OLD (pre-adjacency-index, pre-qgate, pre-fan-effect)
+ * BFS reaching multi-second/crash territory at depth 2-3 from unbounded
+ * hub-node fan-out. That specific failure mode is already substantially
+ * mitigated here by mechanisms #105's branch predates: the adjacency index
+ * (O(degree) not O(E) per hop), the query-aware qgate (prunes semantically
+ * irrelevant branches), the ACT-R fan-effect correction (dampens popular-
+ * hub over-connectivity), and the 0.02 firing threshold. A beam cap is still
+ * a genuine additional, orthogonal bound: it caps WORST-CASE per-hop
+ * expansion width regardless of how many targets happen to pass the soft
+ * gates above, so it is kept as defense in depth rather than dropped as
+ * redundant.
+ *
+ * Bounds the number of frontier nodes EXPANDED per hop-level (see the
+ * level-batching in the BFS below). Every reached node still gets its
+ * best_bg[]/reached[] recorded and appears in the returned/promoted set —
+ * the cap bounds only how far ASSOCIATIVE SPREAD continues past a level,
+ * never the direct seed matches or the reported result set. Tunable via
+ * ENGRAM_ACTIVATE_BEAM (default 128, matching #105); set very high (e.g.
+ * the node count) to recover the pre-cap unbounded-per-level behaviour. */
+static int64_t engram_activate_beam(void) {
+    static int64_t v = -1;
+    if (v >= 0) return v;
+    const char* s = getenv("ENGRAM_ACTIVATE_BEAM");
+    int64_t d = 128;
+    if (s && *s) { char* e = NULL; long t = strtol(s, &e, 10); if (e != s && t > 0) d = (int64_t)t; }
+    v = d; return v;
+}
+
 el_val_t engram_activate(el_val_t query, el_val_t depth) {
     EngramStore* g = engram_get();
     const char* q = EL_CSTR(query);
@@ -9552,25 +9604,43 @@ el_val_t engram_activate(el_val_t query, el_val_t depth) {
             backfilled++;
         }
     }
-    /* Query embedding, cached single-slot: the curiosity loop re-issues the
-     * same 4 rotating phrases, so consecutive identical queries skip the
-     * HTTP round-trip entirely. */
-    static char*   _eg_qcache_text = NULL;
-    static float*  _eg_qcache_emb  = NULL;
-    static int32_t _eg_qcache_dim  = 0;
+    /* Query embedding cache (2026-08-15, PR #105 port: single-slot -> direct-
+     * mapped multi-slot). The single-slot cache below this comment's history
+     * only remembered the LAST query, so "the curiosity loop re-issues the
+     * same 4 rotating phrases" only hit when two CONSECUTIVE calls used the
+     * SAME phrase — any rotation among >1 phrase evicted the slot before it
+     * could be reused. #105 measured this as a real cost (a repeated query
+     * costing a full Ollama round-trip whenever a different phrase intervened)
+     * and fixed it with a direct-mapped, FNV-1a-keyed cache sized for the
+     * rotation. Ported here on TOP of the existing cosq/e_eff semantic layer
+     * (this cache only ever supplies q_emb/q_dim into that unchanged
+     * pipeline) rather than replacing it — see the M8/#105 reconciliation
+     * note above eg_cosq_at. ENGRAM_QCACHE_SIZE must be a power of two (mask
+     * indexing below). Full strcmp on lookup rejects hash collisions; each
+     * slot owns its `text`/`vec` and is freed on eviction, matching the old
+     * single-slot free/replace contract — q_emb below still points at cache-
+     * owned memory the caller must NOT free, just as before. */
+#define ENGRAM_QCACHE_SIZE 1024
+    typedef struct { char* text; uint64_t hash; float* vec; int32_t dim; } EgQCacheEntry;
+    static EgQCacheEntry _eg_qcache[ENGRAM_QCACHE_SIZE];
     float*  q_emb = NULL;
     int32_t q_dim = 0;
-    if (_eg_qcache_text && strcmp(_eg_qcache_text, q) == 0) {
-        q_emb = _eg_qcache_emb; q_dim = _eg_qcache_dim;
-    } else {
-        int32_t d = 0;
-        float* v = eg_embed_fetch(q, &d);
-        if (v) {
-            free(_eg_qcache_text); free(_eg_qcache_emb);
-            _eg_qcache_text = strdup(q);
-            _eg_qcache_emb  = v;
-            _eg_qcache_dim  = d;
-            q_emb = v; q_dim = d;
+    {
+        uint64_t qh = engram_id_hash(q);
+        EgQCacheEntry* slot = &_eg_qcache[qh & (ENGRAM_QCACHE_SIZE - 1)];
+        if (slot->vec && slot->hash == qh && slot->text && strcmp(slot->text, q) == 0) {
+            q_emb = slot->vec; q_dim = slot->dim;
+        } else {
+            int32_t d = 0;
+            float* v = eg_embed_fetch(q, &d);
+            if (v) {
+                free(slot->text); free(slot->vec);   /* evict prior occupant */
+                slot->text = strdup(q);
+                slot->hash = qh;
+                slot->vec  = v;
+                slot->dim  = d;
+                q_emb = v; q_dim = d;
+            }
         }
     }
     /* ── Context centroid fold-in (2026-07-29) ──────────────────────────
@@ -9997,8 +10067,45 @@ el_val_t engram_activate(el_val_t query, el_val_t depth) {
     const double FAN_DREF = (g->adj_connected > 0)
         ? (2.0 * (double)g->edge_count / (double)g->adj_connected) : 0.0;
     _eg_act_fan_dref = FAN_DREF;
+    const int64_t activate_beam = engram_activate_beam();
     while (fhead < ftail) {
-        Frontier f = fr[fhead++];
+      /* Level-batch (2026-08-15, PR #105 port): entries sharing .hops are
+       * always contiguous — hop k+1 entries are appended only while
+       * processing hop k, strictly after the current ftail, so they form one
+       * block right after hop k's block (see engram_activate_beam's comment
+       * for why this holds even with the improve-and-re-enqueue behavior
+       * below). Find this level's extent, then beam-select which of it
+       * EXPANDS; every entry in the level still gets recorded via
+       * reached[]/best_bg[] regardless (that happened when it was enqueued,
+       * one level up) — the cap bounds propagation width only. */
+      int64_t level_hops  = fr[fhead].hops;
+      int64_t level_start = fhead;
+      int64_t level_end   = fhead;
+      while (level_end < ftail && fr[level_end].hops == level_hops) level_end++;
+      int64_t level_n = level_end - level_start;
+      unsigned char* expand = NULL;
+      if (level_n > activate_beam) {
+          expand = calloc((size_t)level_n, 1);
+          if (expand) {
+              /* Partial selection: mark the top-`activate_beam` entries by
+               * .act. O(beam*level_n) — beam is the small tunable. */
+              for (int64_t bsel = 0; bsel < activate_beam; bsel++) {
+                  int64_t best = -1;
+                  for (int64_t k = 0; k < level_n; k++) {
+                      if (expand[k]) continue;
+                      if (best < 0 || fr[level_start+k].act > fr[level_start+best].act)
+                          best = k;
+                  }
+                  if (best < 0) break;
+                  expand[best] = 1;
+              }
+          }
+          /* OOM on the selection map: expand stays NULL -> this level runs
+           * unbounded, same as if beam were disabled. Never silently wrong. */
+      }
+      for (int64_t lk = level_start; lk < level_end; lk++) {
+        if (expand && !expand[lk - level_start]) continue;
+        Frontier f = fr[lk];
         if (f.hops >= max_depth) continue;
         int64_t cur = f.idx;
         int64_t new_hops = f.hops + 1;
@@ -10130,6 +10237,9 @@ el_val_t engram_activate(el_val_t query, el_val_t depth) {
                 }
             }
         }
+      }
+      free(expand);
+      fhead = level_end;
     }
     /* Persist layer-1 background_activation to node store. */
     for (int64_t i = 0; i < g->node_count; i++) {
@@ -16375,6 +16485,41 @@ el_val_t el_base64_encode_n(const unsigned char* data, size_t len, int url_safe)
     return el_wrap_str(out);
 }
 
+/* fs_read_b64_chunk — binary-safe windowed file read: read up to `length`
+ * raw bytes starting at byte `offset` from `path` and return them base64-
+ * encoded (RFC 4648, standard alphabet, padded). The raw bytes are read into
+ * a local C buffer and base64-encoded directly here — they never pass
+ * through an el_val_t string as raw bytes, so embedded NUL bytes (common in
+ * real PCM audio) never hit a strlen()-based code path. This mirrors the
+ * existing llm_vision() image-attachment path (read file -> base64 in C ->
+ * hand back a plain-ASCII string) and the http_*_to_file() rationale above:
+ * bypass the string wrapper entirely for the part that must stay binary.
+ *
+ * Returns "" if the path can't be opened, offset is negative or past EOF,
+ * or length <= 0. A short final chunk (less than `length` bytes remaining)
+ * is returned truncated to what's actually on disk — never padded/invented. */
+el_val_t fs_read_b64_chunk(el_val_t pathv, el_val_t offsetv, el_val_t lengthv) {
+    const char* path = EL_CSTR(pathv);
+    int64_t offset = (int64_t)offsetv;
+    int64_t length = (int64_t)lengthv;
+    if (!path || !*path || offset < 0 || length <= 0) return el_wrap_str(el_strdup(""));
+    FILE* f = fopen(path, "rb");
+    if (!f) return el_wrap_str(el_strdup(""));
+    fseek(f, 0, SEEK_END);
+    long sz = ftell(f);
+    if (sz < 0 || offset >= sz) { fclose(f); return el_wrap_str(el_strdup("")); }
+    fseek(f, (long)offset, SEEK_SET);
+    long remain = sz - (long)offset;
+    size_t want = (size_t)(((int64_t)remain < length) ? remain : length);
+    unsigned char* buf = malloc(want > 0 ? want : 1);
+    if (!buf) { fclose(f); return el_wrap_str(el_strdup("")); }
+    size_t got = fread(buf, 1, want, f);
+    fclose(f);
+    el_val_t out = el_base64_encode_n(buf, got, /*url_safe=*/0);
+    free(buf);
+    return out;
+}
+
 /* Decode either alphabet — accepts both '+/' and '-_' transparently, and
  * tolerates missing padding (which JWTs typically omit). Whitespace is
  * skipped for robustness. Invalid characters cause the decode to stop and

lang/runtime/el_runtime.h

@@ -235,6 +235,20 @@ el_val_t  fs_list(el_val_t path);
 el_val_t  fs_exists(el_val_t path);
 el_val_t  fs_mkdir(el_val_t path);   /* mkdir -p, mode 0755                */
 
+/* Real on-disk byte count via stat() — not strlen(). Use this (not
+ * str_len(fs_read(path))) when a file may contain binary content, since
+ * fs_read()'s result truncates at the first embedded NUL under strlen-based
+ * string ops. Returns -1 if the path doesn't exist. */
+el_val_t  fs_size(el_val_t path);
+
+/* Binary-safe windowed read: read up to `length` bytes starting at byte
+ * `offset` from `path` and return them base64-encoded. Bytes are read and
+ * encoded in C without ever passing through an el_val_t string as raw
+ * bytes, so embedded NULs (routine in PCM audio) can't truncate the result.
+ * Returns "" on any failure or when offset is past EOF; a final short
+ * window returns only the bytes that actually exist on disk. */
+el_val_t  fs_read_b64_chunk(el_val_t path, el_val_t offset, el_val_t length);
+
 /* Length-explicit binary write. `length` is an Int (el_val_t holding the
  * byte count). The caller knows the length from context — typically because
  * `bytes` came from base64_decode (which produces a magic-tagged binary

lang/runtime/vindex_bench.c

@@ -7,29 +7,11 @@
  *
  * Read-only: never opens a socket, never writes the store. Safe on an nsbx clone.
  *
- * Also runs the brute-force oracle a second (and third) way, through the
- * batch-cosine Strategies behind eg_cosine_batch_strategy.h — the ggml
- * strategy and the hand-rolled-Metal strategy (Apple/Metal only; see
- * eg_cosine_batch.h/eg_cosine_batch_strategy.h) — and reports each one's
- * latency + a correctness check against the CPU oracle side-by-side with the
- * existing CPU-vs-HNSW numbers. This harness deliberately reaches past the
- * single-selection Factory (eg_cosine_batch.c) to instantiate every
- * compiled-in strategy directly, so it can compare all of them against the
- * SAME dataset in one run — that is the harness's whole job; a real call
- * site (el_runtime.c) never does this, it only ever calls the plain
- * eg_cosine_batch()/eg_cosine_batch_multi() adapter functions.
- * EL_METAL_COSINE=0 forces CPU-only (skips every strategy comparison).
- *
- * Build (macOS, ggml + hand-rolled Metal): see build_vindex_bench.sh.
- * Build (Linux / no Metal): omit every eg_cosine_batch_strategy_*.{c,m} file
- *   except eg_cosine_batch_strategy_cpu.c — this file never references
- *   ggml/Metal directly except through the plain-C strategy header, guarded
- *   by the same EG_HAVE_STRATEGY_* build macros the Factory itself uses.
+ * Build:  cc -O2 -std=c11 vindex_bench.c engram_vindex.c -lm -o vindex_bench
  * Usage:  vindex_bench store <neuron.egm> <dim> [nqueries] [k] [ef_csv]
  *         vindex_bench synth <N> [dim] [clusters] [nqueries] [k] [ef_csv]
  */
 #include "engram_vindex.h"
-#include "eg_cosine_batch_strategy.h"
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
@@ -92,106 +74,6 @@ static double recall_at_k(const int* gt, const uint64_t* ann, int nann, int k){
     return (double)hit / (double)k;
 }
 
-/* EL_METAL_COSINE: 0/off/false disables EVERY strategy comparison outright
- * (falls back to brute_topk() only), matching el_runtime.c's own gate for
- * the same env var (back-compat name kept from PR #114; it now gates all
- * GPU-backed strategies, not just the hand-rolled Metal one). Unset or any
- * other value = try every compiled-in strategy, report each that's
- * available, skip (without failing the run) any that isn't. */
-static bool g_strategy_env_checked = false;
-static bool g_strategy_disabled_by_env = false;
-static void eg_strategy_check_env_once(void){
-    if (g_strategy_env_checked) return;
-    g_strategy_env_checked = true;
-    const char* v = getenv("EL_METAL_COSINE");
-    if (v && (v[0]=='0' || v[0]=='n' || v[0]=='N' || v[0]=='f' || v[0]=='F'))
-        g_strategy_disabled_by_env = true;
-}
-
-/* Batched sibling of brute_topk, generalized over ANY EgCosineBatchStrategy:
- * computes top-k for ALL nq queries in ONE strategy->batch_multi() call,
- * uploading/preparing the node population exactly once instead of once per
- * query. out_ids/out_d are nq*k, row-major (query i's results at
- * out_ids+i*k / out_d+i*k). Returns false (nothing written) on any
- * failure/unavailability; caller treats that as "skip this strategy in the
- * report", never as a hard error. */
-static bool batch_topk_strategy(const EgCosineBatchStrategy* strat,
-                                const float* data, int n, int dim,
-                                const float* queries, int nq,
-                                int k, int* out_ids, float* out_d){
-    if (!strat || !strat->available()) return false;
-
-    const float** row_ptrs = malloc((size_t)n * sizeof(float*));
-    int32_t* dims = malloc((size_t)n * sizeof(int32_t));
-    double* scores = malloc((size_t)nq * (size_t)n * sizeof(double));
-    if (!row_ptrs || !dims || !scores) { free(row_ptrs); free(dims); free(scores); return false; }
-
-    for (int i = 0; i < n; i++) { row_ptrs[i] = data + (size_t)i * dim; dims[i] = dim; }
-
-    bool ok = strat->batch_multi(queries, dim, nq, row_ptrs, dims, n, scores);
-    free(row_ptrs); free(dims);
-    if (!ok) { free(scores); return false; }
-
-    for (int qi = 0; qi < nq; qi++) {
-        int* ids = out_ids + (size_t)qi * k;
-        float* ds = out_d + (size_t)qi * k;
-        const double* srow = scores + (size_t)qi * n;
-        for (int i = 0; i < k; i++) { ids[i] = -1; ds[i] = 3.0f; }
-        for (int i = 0; i < n; i++) {
-            float d = 1.0f - (float)srow[i];   /* same distance convention as brute_topk */
-            if (d >= ds[k-1]) continue;
-            int p = k - 1;
-            while (p > 0 && ds[p-1] > d) { ds[p] = ds[p-1]; ids[p] = ids[p-1]; p--; }
-            ds[p] = d; ids[p] = i;
-        }
-    }
-    free(scores);
-    return true;
-}
-
-/* Runs batch_topk_strategy for one named strategy over ALL nq queries, diffs
- * against the CPU ground truth (gt/gd, both nq*k), and prints a report line
- * in the same shape PR #114 established for BRUTE-METAL — id-recall over
- * every query plus the actual max/mean same-rank distance delta across
- * every (query,rank) pair that was compared, never fabricated or assumed. */
-static void report_strategy_vs_oracle(const char* label, const EgCosineBatchStrategy* strat,
-                                      const float* data, int n, int dim,
-                                      const float* qv, int nq, int k,
-                                      const int* gt, const float* gd, double brute_ms){
-    if (g_strategy_disabled_by_env) { printf("%-13s: disabled via EL_METAL_COSINE\n", label); return; }
-    if (!strat || !strat->available()) { printf("%-13s: not available on this build/host — skipped\n", label); return; }
-
-    int* gtm = malloc((size_t)nq*k*sizeof(int));
-    float* gdm = malloc((size_t)nq*k*sizeof(float));
-    double tm0 = now_s();
-    bool ok = batch_topk_strategy(strat, data, n, dim, qv, nq, k, gtm, gdm);
-    double strat_ms = (now_s()-tm0)*1000.0/nq;
-    if (ok) {
-        double rec_sum = 0; double max_ddiff = 0; double sum_ddiff = 0; int compared = 0;
-        for (int i=0;i<nq;i++) {
-            const int* ids_gt = gt+(size_t)i*k;
-            const float* d_gt = gd+(size_t)i*k;
-            const int* ids_m = gtm+(size_t)i*k;
-            const float* d_m = gdm+(size_t)i*k;
-            uint64_t idset[512]; int m = (k<512)?k:512;
-            for (int j=0;j<m;j++) idset[j] = (uint64_t)ids_m[j];
-            rec_sum += recall_at_k(ids_gt, idset, m, k);
-            for (int j=0;j<k;j++) {
-                if (ids_gt[j] == ids_m[j]) {
-                    double diff = fabs((double)d_gt[j]-(double)d_m[j]);
-                    if (diff>max_ddiff) max_ddiff=diff;
-                    sum_ddiff += diff; compared++;
-                }
-            }
-        }
-        printf("%-13s: %8.3f ms/query   (%.1fx vs CPU brute; id-recall %.4f vs CPU oracle over %d queries; same-rank |Δdist|: max %.2e, mean %.2e over %d compared)\n",
-               label, strat_ms, brute_ms/strat_ms, rec_sum/nq, nq, max_ddiff, compared?sum_ddiff/compared:0.0, compared);
-    } else {
-        printf("%-13s: batch call failed mid-run — skipped\n", label);
-    }
-    free(gtm); free(gdm);
-}
-
 /* Parse "64,128,256" into an int array; returns count. */
 static int parse_csv(const char* s, int* out, int maxo){
     int n=0; if(!s||!*s) return 0;
@@ -260,37 +142,13 @@ static void run_bench(const char* label, float* data, int n, int dim,
         l2norm(dst, dim);
     }
 
-    /* ground truth: brute-force top-k for every query (also the oracle latency).
-     * gd is nq*k (one real slot per query, not a shared scratch buffer) so the
-     * strategy comparisons below can diff against every query's actual
-     * distances, not just whichever query happened to run last. */
+    /* ground truth: brute-force top-k for every query (also the oracle latency). */
     int* gt = malloc((size_t)nq*k*sizeof(int));
-    float* gd = malloc((size_t)nq*k*sizeof(float));
+    float* gd = malloc((size_t)k*sizeof(float));
     double tb0 = now_s();
-    for (int i=0;i<nq;i++) brute_topk(data, n, dim, qv+(size_t)i*dim, k, gt+(size_t)i*k, gd+(size_t)i*k);
+    for (int i=0;i<nq;i++) brute_topk(data, n, dim, qv+(size_t)i*dim, k, gt+(size_t)i*k, gd);
     double brute_ms = (now_s()-tb0)*1000.0/nq;
-    printf("BRUTE-FORCE  : %8.3f ms/query   (oracle; O(N*D), CPU)\n", brute_ms);
-
-    /* GPU-backed oracles: SAME nq queries, SAME top-k contract, via each
-     * compiled-in Strategy's batch_multi() (uploads/prepares the node
-     * population once, not once per query). Run only for strategies that
-     * are actually available (checked internally) — never fabricated, never
-     * assumed. Verified against the CPU ground truth computed above:
-     * id-recall across ALL nq queries, plus the actual max/mean distance
-     * delta across every (query,rank) pair that was compared. */
-    eg_strategy_check_env_once();
-#ifdef EG_HAVE_STRATEGY_GGML
-    report_strategy_vs_oracle("BRUTE-GGML", eg_cosine_batch_strategy_ggml(),
-                              data, n, dim, qv, nq, k, gt, gd, brute_ms);
-#else
-    printf("BRUTE-GGML   : strategy not compiled into this build\n");
-#endif
-#ifdef EG_HAVE_STRATEGY_METAL_HAND
-    report_strategy_vs_oracle("BRUTE-METAL", eg_cosine_batch_strategy_metal_hand(),
-                              data, n, dim, qv, nq, k, gt, gd, brute_ms);
-#else
-    printf("BRUTE-METAL  : strategy not compiled into this build\n");
-#endif
+    printf("BRUTE-FORCE  : %8.3f ms/query   (oracle; O(N*D))\n", brute_ms);
 
     /* HNSW at each ef. */
     uint64_t* aid = malloc((size_t)k*sizeof(uint64_t));