#116 shipped the ggml strategy at 0.9933 id-recall against the CPU oracle
while the hand-rolled Metal kernel it replaced scored 0.9997 — a ~150x worse
error margin. That was not an inherent property of ggml. It was a usage bug in
this file, and this commit fixes it.
ggml-metal has two F32xF32 matmul kernels and picks between them purely on
ne11, the number of B rows, which for us is the query-batch size:
ne11 <= 8 -> kernel_mul_mv_ext_f32_f32_* / kernel_mul_mv_f32_f32_*,
templated <float, float> — genuine F32.
ne11 > 8 -> kernel_mul_mm_f32_f32, templated
<half, half4x4, simdgroup_half8x8, half, half2x4, ...> —
BOTH operands narrowed to F16, despite F32 tensors on both
sides.
The old code issued one ggml_mul_mat with ne11 = nq (300 in the benchmark),
landing squarely on the F16 path. The file's own header comment asserted the
opposite ("computes in F32 on the Metal backend"); that claim was wrong and is
replaced with the measurement.
Fix: emit ceil(nq/8) mul_mats over ne11<=8 ggml_view_2d slices of one query
tensor, all expanded into ONE graph and one ggml_backend_graph_compute, so the
node matrix is still uploaded and shared exactly once. EL_GGML_MULMAT_CHUNK
overrides the 8; setting it >= nq reproduces the old behaviour exactly, which
is also how the before/after below was measured in a single binary.
Measured, real store snapshot, 13415 live embedded nodes, dim=768, 300 real
queries, vs the CPU double-accumulated oracle (vindex_bench, offline copy of
the store — no live service touched):
id-recall same-rank |Δdist| max mean
old (ne11=300) 0.9933 6.80e-05 1.43e-05
new (ne11<=8) 0.9987 4.77e-07 9.30e-08
hand-rolled 0.9997 3.58e-07 7.55e-08
~145x better max error, ~154x better mean — now the same order of magnitude as
the hand-rolled kernel rather than 150x off it.
The cost is real and is documented rather than buried. Median of 15 reps of
the whole batch_multi() call, three runs: 13.2-14.4ms unchunked, 19.9-20.2ms
chunked, 17.7-18.0ms hand-rolled. Correctness costs ~+6.7ms per 300-query
batch and leaves ggml ~12% behind the hand-rolled kernel instead of ~35%
ahead. It cannot be recovered inside ggml: an fp32 matmul on Metal must
re-stream the node matrix once per <=8 queries, and ggml's Metal backend ships
no fp32 TILED matmul, so "fast" and "fp32" are genuinely exclusive there.
Two things that did NOT work, recorded so nobody retries them:
- ggml_mul_mat_set_prec(t, GGML_PREC_F32) does nothing here. Error was
bit-identical with and without it (1.038e-05 either way) — ggml-metal has
no F32-accumulating mul_mm kernel to switch to. ne11 is the only lever.
- The ACCEL/BLAS device looked excellent in an isolated compute-only probe
(3.4-4.0ms, mean |Δdot| 1.5e-08) but is dominated on BOTH axes end-to-end
(0.191 ms/query at 0.9973 recall vs 0.125-0.142 at 0.9987), because the
probe was not competing for the same CPU cores the real call path is. It
stays reachable via EL_GGML_DEVICE as a no-Metal fallback, labelled as
measured-and-rejected, not as a recommendation.
Also corrected: the ~7.8s "cold start" blamed on this file is not this file
re-initialising per call — init was already cached. It is Apple's shader cache
missing on ggml's embedded metallib (~650 kernels), keyed on the library and
shared across processes: the first load on a machine reports
"loaded in 7.670 sec", the next run of a *different* binary reports 0.009 sec.
Once per machine per ggml version, not once per process, and not ours to fix.
Warm ggml init is 44-53ms vs 36-117ms for the hand-rolled strategy.
Loading only libggml-metal.so instead of every plugin in the directory is kept
for tidiness, and explicitly documented as NOT a speedup: 44.7-52.4ms against
46.9-58.9ms, the same number inside noise.
The -2.0 sentinel contract is unchanged and re-verified at batch sizes that
straddle the chunk boundary (1,7,8,9,16,17,33), plus NULL rows, dim
mismatches, zero-norm rows, and an all-invalid population. Notably the old
ne11=300 path fails that same check at a 2e-6 cosine tolerance with 2299
mismatches, which is an independent confirmation of the defect.
bigmerge
c008b7228a