diff --git a/docs/PERF_NOTES.md b/docs/PERF_NOTES.md
index 3872941..82d8c53 100644
--- a/docs/PERF_NOTES.md
+++ b/docs/PERF_NOTES.md
@@ -1063,3 +1063,71 @@ If a later product path chases GPU for real U-NSGA-III, it looks like this:
 6. **Non-goals for v0 of that path.** Rewriting all of `Run` on GPU. The Array host-rebuild (Phase C, discarded as noise). HIP on the RX 5700 XT (parked: Win10 + gfx1010 unsupported, HSA reported no adapters; bendlang#979 stays on the shelf). CUDA only if someone asks later. First host remains the Mac mini M4 / Metal.
 
 Not started. The open checkboxes live under Bend 0.3 in [ROADMAP.md](ROADMAP.md). Needs Jason's go before code.
+
+## Experimental hand-written Metal associate kernel (2026-11-10)
+
+Isolated experiment under [`experiments/metal-kernels/`](../experiments/metal-kernels/). **No product `Run` change.** No `!` on the default algorithm. No IGD claim (the shipping IGD path is untouched). Raw logs and the clean table: [`experiments/metal-kernels/results/TABLES.md`](../experiments/metal-kernels/results/TABLES.md).
+
+### Chosen hot spot
+
+Fresh warm-native `ab/profile_bend_run.py` on the shared Linux box (8 cores, `bend 2.0.40`, `--threads 1`, seed=1). Host load was high (~6.3–6.7), so `nds` moved between runs; `associate` and `niche` were stable.
+
+**DTLZ2 M=3, p=12, pop=92, gens=150** (oracle):
+
+| run | run_s | nds | niche | associate | normalize | offspring |
+|-----|------:|----:|------:|----------:|----------:|----------:|
+| warm 1 | 2.544 | 566 ms (23%) | 680 ms (27%) | 448 ms (18%) | 379 ms (15%) | 290 ms (11%) |
+| warm 2 | 2.833 | 734 ms (27%) | 680 ms (25%) | 435 ms (16%) | 462 ms (17%) | 303 ms (11%) |
+
+**DTLZ2 M=3, p=23, pop=512, gens=20** (larger `Run`, first binary after compile_s=9.862): `run_s=8.879`. `nds` 4314 ms (**50%**), `niche` 1524 ms (17%), `offspring` 1144 ms (13%), `associate` 695 ms (8%). NDS grows roughly with N²; it is divergent (guide: work like n-queens stays on CPU) and is not a first kernel target. Last-front `niche` fill is sequential and RNG-ordered, so a kernel that changed pick order would change fronts.
+
+The kernel target is **association**: perpendicular distance of each normalized point to each reference direction, strict-`<` argmin (ties keep the lower index), plus the O(N) niche-count histogram. That is the largest *uniform* numeric map at oracle size (16–18% of wall) and the Phase D family that had a measured 64× Metal win on inflated 32k × 3003 stand-ins.
+
+### Verdict
+
+A hand-written Metal kernel **beats the current Bend CPU associate map** at every size measured, including the realistic DTLZ2 combined pool, and **beats Bend's own `!` Metal backend** at the only size where that backend completed. Bend `!` is slower than Bend CPU at 184 × 91 (same story as Phase D) and **stack-overflows** at 2048 × 276 and 32k × 3003. A negative result for Bend's GPU bang on this map, and a real win for a flat MSL kernel against the cons-list CPU path, especially at big-N (277× at 32k × 3003).
+
+The kernel does **not** beat a flat Swift CPU reference end-to-end at 184 × 91: dispatch plus transfer (~0.25 ms) dominate a 0.116 ms kernel that is tied with the 0.108 ms Swift walk. That e2e crossing vs a C-like CPU is between 512 × 276 and 1024 × 276.
+
+### Method
+
+- Host: JMac-mini.local, Apple M4, Metal 4, 16 GB. `bend 2.0.40`. 2026-11-10 00:48–00:52 America/New_York. Load 2.20 2.23 2.12 at start.
+- Algorithm matches `src/survival.bend` `resid_sq` / `perp` / `nearest_go` (right-fold F32 dots and residuals, `ww < 1e-16` uses `||f||`, argmin on strict `<`).
+- Hand MSL: warmup then 5–7 timed reps; median and min–max. E2e includes buffer rewrite, dispatch, wait, readback touch. Kernel-only uses command-buffer GPU timestamps. Metal `fastMathEnabled = false`.
+- Bend: `IO.now()` (integer ms) around in-process repeats; per-rep = total / reps. `+RTS --threads 1`. GPU on vs `--gpu off` for the `!` binary.
+- Forced work: each Bend repeat folds argmin indices into `sum_idx_acc`, which must equal `reps * sum_idx`. Without that, unused repeats were dropped and wall time stopped scaling with `reps`.
+- Did it happen: MSL outputs matched the Swift CPU ref (0 index mismatches, counts equal, max abs distance ≤ 5.6e-8). Bend `sum_idx` matched MSL at every size that completed. Small fixture `n=8 r=5` also matched `ref_associate.py` on the box (max abs 5.4e-9).
+
+### Primary numbers
+
+| N × R × M | hand MSL e2e (median, range) | hand MSL kernel | Bend product (current path) | Bend Metal (`!`) | Bend `!` / `--gpu off` |
+|-----------|-----------------------------:|----------------:|----------------------------:|-----------------:|-----------------------:|
+| 184 × 91 × 3 (DTLZ2 pool) | **0.365 ms** (0.335–0.369), n=7 | **0.116 ms** (0.111–0.116) | **0.550 ms/rep** (44 ms / 80; acc=80×8819) | **6.55 ms/rep** (131 ms / 20; acc=20×8819) | 0.500 ms/rep (10 ms / 20) |
+| 2048 × 276 × 3 | **0.392 ms** (0.384–0.433), n=7 | **0.216 ms** (0.216–0.216) | **19.8 ms/rep** (198 ms / 10; acc=10×279827) | crash (stack overflow) | 18.6 ms/rep (149 ms / 8) |
+| 32000 × 3003 × 3 | **7.934 ms** (7.928–10.577), n=5 | **7.675 ms** (7.665–10.229) | **2195 ms/rep** (6586 ms / 3; acc=3×47882676) | crash (stack overflow) | 2236 ms/rep (4471 ms / 2) |
+
+MSL e2e vs Bend product: **1.51×**, **50.5×**, **277×**. MSL e2e vs Bend Metal: **17.9×** at 184 × 91; no Bend Metal number at the larger sizes.
+
+Limiter: at 184 × 91 the MSL e2e is dispatch-bound (e2e 0.365 ms vs kernel 0.116 ms vs Swift CPU 0.108 ms). At 32k × 3003, e2e ≈ kernel (7.93 vs 7.67 ms). That kernel is one thread per point walking all R refs (~96 million pairwise residuals, about 1.4e9 FLOP in 7.67 ms ≈ 180 GFLOP/s). A tiled layout could go faster; this prototype does not. Bend product at 32k is a sequential cons-list `nearest_go` on one CPU thread (~2.2 s/call), which is why 277× is plausible and still well under a “removed 100% of the 2.2 s” ceiling.
+
+`IO.now()` is 1 ms resolution, so the 184 × 91 Bend product 44 ms / 80 has about ±0.013 ms/rep quantization. That does not change the 1.51× call versus the 0.365 ms MSL median.
+
+### Break-even
+
+- **Vs current Bend CPU associate:** already ahead at the realistic DTLZ2 combined pool (184 × 91, 1.51× e2e). The gap widens with N×R.
+- **Vs Bend `!` Metal:** ahead at 184 × 91 (17.9×). Bend `!` does not currently run this fork-tree at 2048 × 276 or above (`memory fault (machine stack overflow?)`).
+- **Vs a flat Swift CPU walk of the same math:** MSL e2e is slower at 184 × 91 and 512 × 276, and faster from **1024 × 276** (1.54×) upward (17× at 8k × 1050, 61× at 32k × 3003). Extra MSL-only sizes are in `results/TABLES.md`.
+
+This is a different crossing than Phase D's Bend-`!` vs `--gpu off` rule (`pad(N)×M×(3+inflate) ≳ 3×10⁸`). A hand kernel has no ~80 ms process-wake tax and no cons-list fork tree, so it can win on residual-only (inflate=1) work that Bend `!` lost.
+
+### Correctness and IGD
+
+Product `src/` is unchanged, so the ZDT/DTLZ IGD yardstick is unchanged by construction (the `verify` job still runs seed-1 `oracle_multiseed.py` + `compare_yardstick.py`). Isolated map checks:
+
+- Box: Python `ref_associate.py` vs Bend product dump, n=8 r=5 seed=1: 0 index mismatches, counts match, max abs dist 5.4e-9.
+- Mac: same fixture vs hand MSL: 0 index mismatches, counts match, max abs dist 3.0e-8.
+- Every completed bench size: MSL `sum_idx` equals Bend `sum_idx`. Distance tolerance abs 1e-6 / rel 1e-5; observed max abs ≤ 5.6e-8.
+
+### CUDA
+
+`experiments/metal-kernels/associate.cu` is the same algorithm. `nvcc` is not on the box; a CUDA toolkit is not a cheap install (multi-GB), so the file was not compile-checked. No CUDA timings are claimed.
diff --git a/experiments/metal-kernels/.gitignore b/experiments/metal-kernels/.gitignore
new file mode 100644
index 0000000..d06ce3c
--- /dev/null
+++ b/experiments/metal-kernels/.gitignore
@@ -0,0 +1,5 @@
+associate_harness
+bin_*
+wrapper_tmp.bend
+__pycache__/
+*.pyc
diff --git a/experiments/metal-kernels/AssociateHarness.swift b/experiments/metal-kernels/AssociateHarness.swift
new file mode 100644
index 0000000..da8fdb7
--- /dev/null
+++ b/experiments/metal-kernels/AssociateHarness.swift
@@ -0,0 +1,362 @@
+// Experimental Metal host for associate + niche_count.
+// Not a product code path. Times include host<->device traffic for e2e;
+// kernel-only uses MTLCommandBuffer GPU timestamps.
+import Foundation
+import Metal
+
+let LCG_MUL: UInt32 = 1_664_525
+let LCG_ADD: UInt32 = 1_013_904_223
+let BEST_D0: Float = 1_000_000_000.0
+let TINY_WW: Float = 1.0e-16
+let DIST_ABS: Float = 1.0e-6
+let DIST_REL: Float = 1.0e-5
+
+func lcg(_ state: UInt32) -> UInt32 {
+    return state &* LCG_MUL &+ LCG_ADD
+}
+
+func u32ToUnit(_ state: UInt32) -> Float {
+    let u = (state >> 8) & 0x00FF_FFFF
+    return Float(u) / 16_777_216.0
+}
+
+func genMatrix(rows: Int, cols: Int, seed: UInt32) -> ([Float], UInt32) {
+    var s = seed
+    var out = [Float](repeating: 0, count: rows * cols)
+    var off = 0
+    for _ in 0..<rows {
+        for _ in 0..<cols {
+            s = lcg(s)
+            out[off] = u32ToUnit(s)
+            off += 1
+        }
+    }
+    return (out, s)
+}
+
+func rightDot(_ a: [Float], _ ao: Int, _ b: [Float], _ bo: Int, _ m: Int) -> Float {
+    var acc: Float = 0
+    var i = m
+    while i > 0 {
+        i -= 1
+        acc = a[ao + i] * b[bo + i] + acc
+    }
+    return acc
+}
+
+func rightResidSq(_ f: [Float], _ fo: Int, _ w: [Float], _ wo: Int, _ t: Float, _ m: Int) -> Float {
+    var acc: Float = 0
+    var i = m
+    while i > 0 {
+        i -= 1
+        let r = f[fo + i] - t * w[wo + i]
+        acc = r * r + acc
+    }
+    return acc
+}
+
+func cpuAssociate(points: [Float], refs: [Float], n: Int, r: Int, m: Int) -> ([UInt32], [Float]) {
+    var idx = [UInt32](repeating: 0, count: n)
+    var dist = [Float](repeating: 0, count: n)
+    for i in 0..<n {
+        let fo = i * m
+        var bestI: UInt32 = 0
+        var bestD = BEST_D0
+        for j in 0..<r {
+            let wo = j * m
+            let ww = rightDot(refs, wo, refs, wo, m)
+            let d: Float
+            if ww < TINY_WW {
+                d = rightDot(points, fo, points, fo, m).squareRoot()
+            } else {
+                let t = rightDot(points, fo, refs, wo, m) / ww
+                d = rightResidSq(points, fo, refs, wo, t, m).squareRoot()
+            }
+            if d < bestD {
+                bestD = d
+                bestI = UInt32(j)
+            }
+        }
+        idx[i] = bestI
+        dist[i] = bestD
+    }
+    return (idx, dist)
+}
+
+func cpuNicheCount(_ idx: [UInt32], r: Int) -> [UInt32] {
+    var counts = [UInt32](repeating: 0, count: r)
+    for v in idx {
+        let j = Int(v)
+        if j >= 0 && j < r { counts[j] += 1 }
+    }
+    return counts
+}
+
+func distClose(_ a: Float, _ b: Float) -> Bool {
+    let d = abs(a - b)
+    if d <= DIST_ABS { return true }
+    return d <= DIST_REL * max(abs(a), abs(b), 1e-30)
+}
+
+func median(_ xs: [Double]) -> Double {
+    let s = xs.sorted()
+    if s.isEmpty { return 0 }
+    return s[s.count / 2]
+}
+
+struct Args {
+    var n = 184
+    var r = 91
+    var m = 3
+    var seed: UInt32 = 1
+    var warmup = 1
+    var reps = 7
+    var metal = "associate.metal"
+    var json: String? = nil
+    var dump = false
+    var includeZeroRef = false
+}
+
+func parseArgs() -> Args {
+    var a = Args()
+    let argv = CommandLine.arguments
+    var i = 1
+    while i < argv.count {
+        let k = argv[i]
+        func take() -> String {
+            i += 1
+            return argv[i]
+        }
+        switch k {
+        case "--n": a.n = Int(take())!
+        case "--r": a.r = Int(take())!
+        case "--m": a.m = Int(take())!
+        case "--seed": a.seed = UInt32(take())!
+        case "--warmup": a.warmup = Int(take())!
+        case "--reps": a.reps = Int(take())!
+        case "--metal": a.metal = take()
+        case "--json": a.json = take()
+        case "--dump": a.dump = true
+        case "--include-zero-ref": a.includeZeroRef = true
+        default:
+            fputs("unknown arg \(k)\n", stderr)
+            exit(2)
+        }
+        i += 1
+    }
+    return a
+}
+
+func main() {
+    let args = parseArgs()
+    var (points, s1) = genMatrix(rows: args.n, cols: args.m, seed: args.seed)
+    var (refs, _) = genMatrix(rows: args.r, cols: args.m, seed: s1)
+    if args.includeZeroRef {
+        for k in 0..<args.m { refs[k] = 0 }
+    }
+
+    let tCPU0 = CFAbsoluteTimeGetCurrent()
+    let (cpuIdx, cpuDist) = cpuAssociate(points: points, refs: refs, n: args.n, r: args.r, m: args.m)
+    let cpuMs = (CFAbsoluteTimeGetCurrent() - tCPU0) * 1000.0
+    let cpuCounts = cpuNicheCount(cpuIdx, r: args.r)
+
+    guard let device = MTLCreateSystemDefaultDevice() else {
+        fputs("no Metal device\n", stderr)
+        exit(1)
+    }
+    let srcPath = URL(fileURLWithPath: args.metal)
+    guard let src = try? String(contentsOf: srcPath, encoding: .utf8) else {
+        fputs("cannot read \(args.metal)\n", stderr)
+        exit(1)
+    }
+    let opt = MTLCompileOptions()
+    opt.fastMathEnabled = false
+    let library: MTLLibrary
+    do {
+        library = try device.makeLibrary(source: src, options: opt)
+    } catch {
+        fputs("Metal compile failed: \(error)\n", stderr)
+        exit(1)
+    }
+    guard let assocFn = library.makeFunction(name: "associate"),
+          let nicheFn = library.makeFunction(name: "niche_count") else {
+        fputs("missing kernel function\n", stderr)
+        exit(1)
+    }
+    let assocPSO: MTLComputePipelineState
+    let nichePSO: MTLComputePipelineState
+    do {
+        assocPSO = try device.makeComputePipelineState(function: assocFn)
+        nichePSO = try device.makeComputePipelineState(function: nicheFn)
+    } catch {
+        fputs("pipeline failed: \(error)\n", stderr)
+        exit(1)
+    }
+    guard let queue = device.makeCommandQueue() else {
+        fputs("no command queue\n", stderr)
+        exit(1)
+    }
+
+    let pointBuf = device.makeBuffer(bytes: points, length: points.count * MemoryLayout<Float>.stride, options: .storageModeShared)!
+    let refBuf = device.makeBuffer(bytes: refs, length: refs.count * MemoryLayout<Float>.stride, options: .storageModeShared)!
+    let idxBuf = device.makeBuffer(length: args.n * MemoryLayout<UInt32>.stride, options: .storageModeShared)!
+    let distBuf = device.makeBuffer(length: args.n * MemoryLayout<Float>.stride, options: .storageModeShared)!
+    let countBuf = device.makeBuffer(length: args.r * MemoryLayout<UInt32>.stride, options: .storageModeShared)!
+
+    var n32 = UInt32(args.n)
+    var r32 = UInt32(args.r)
+    var m32 = UInt32(args.m)
+    let nBuf = device.makeBuffer(bytes: &n32, length: 4, options: .storageModeShared)!
+    let rBuf = device.makeBuffer(bytes: &r32, length: 4, options: .storageModeShared)!
+    let mBuf = device.makeBuffer(bytes: &m32, length: 4, options: .storageModeShared)!
+
+    func dispatch(_ pso: MTLComputePipelineState, _ encoder: MTLComputeCommandEncoder, count: Int) {
+        let tgWidth = min(pso.maxTotalThreadsPerThreadgroup, 256)
+        let tg = MTLSize(width: tgWidth, height: 1, depth: 1)
+        let grid = MTLSize(width: count, height: 1, depth: 1)
+        encoder.dispatchThreads(grid, threadsPerThreadgroup: tg)
+    }
+
+    func runOnce(includeTransfer: Bool) -> (e2eMs: Double, kernelMs: Double, nicheKernelMs: Double) {
+        let t0 = CFAbsoluteTimeGetCurrent()
+        if includeTransfer {
+            memcpy(pointBuf.contents(), points, points.count * MemoryLayout<Float>.stride)
+            memcpy(refBuf.contents(), refs, refs.count * MemoryLayout<Float>.stride)
+        }
+        memset(countBuf.contents(), 0, args.r * MemoryLayout<UInt32>.stride)
+
+        guard let cmd = queue.makeCommandBuffer(),
+              let enc = cmd.makeComputeCommandEncoder() else {
+            fputs("command buffer failed\n", stderr)
+            exit(1)
+        }
+        enc.setComputePipelineState(assocPSO)
+        enc.setBuffer(pointBuf, offset: 0, index: 0)
+        enc.setBuffer(refBuf, offset: 0, index: 1)
+        enc.setBuffer(idxBuf, offset: 0, index: 2)
+        enc.setBuffer(distBuf, offset: 0, index: 3)
+        enc.setBuffer(nBuf, offset: 0, index: 4)
+        enc.setBuffer(rBuf, offset: 0, index: 5)
+        enc.setBuffer(mBuf, offset: 0, index: 6)
+        dispatch(assocPSO, enc, count: args.n)
+
+        enc.setComputePipelineState(nichePSO)
+        enc.setBuffer(idxBuf, offset: 0, index: 0)
+        enc.setBuffer(countBuf, offset: 0, index: 1)
+        enc.setBuffer(nBuf, offset: 0, index: 2)
+        enc.setBuffer(rBuf, offset: 0, index: 3)
+        dispatch(nichePSO, enc, count: args.n)
+        enc.endEncoding()
+        cmd.commit()
+        cmd.waitUntilCompleted()
+
+        // Touch outputs so the readback is part of e2e.
+        let idxPtr = idxBuf.contents().bindMemory(to: UInt32.self, capacity: args.n)
+        let distPtr = distBuf.contents().bindMemory(to: Float.self, capacity: args.n)
+        var sink: UInt64 = 0
+        sink &+= UInt64(idxPtr[0])
+        sink &+= UInt64(idxPtr[args.n - 1])
+        sink &+= UInt64(distPtr[0].bitPattern)
+        _ = sink
+        let t1 = CFAbsoluteTimeGetCurrent()
+
+        let gpu0 = cmd.gpuStartTime
+        let gpu1 = cmd.gpuEndTime
+        let kernelMs = max(0, (gpu1 - gpu0) * 1000.0)
+        return ((t1 - t0) * 1000.0, kernelMs, 0)
+    }
+
+    for _ in 0..<max(args.warmup, 0) {
+        _ = runOnce(includeTransfer: true)
+    }
+
+    var e2e: [Double] = []
+    var kern: [Double] = []
+    for _ in 0..<max(args.reps, 1) {
+        let r = runOnce(includeTransfer: true)
+        e2e.append(r.e2eMs)
+        kern.append(r.kernelMs)
+    }
+
+    let outIdx = Array(UnsafeBufferPointer(
+        start: idxBuf.contents().bindMemory(to: UInt32.self, capacity: args.n),
+        count: args.n
+    ))
+    let outDist = Array(UnsafeBufferPointer(
+        start: distBuf.contents().bindMemory(to: Float.self, capacity: args.n),
+        count: args.n
+    ))
+    let outCounts = Array(UnsafeBufferPointer(
+        start: countBuf.contents().bindMemory(to: UInt32.self, capacity: args.r),
+        count: args.r
+    ))
+
+    var idxMismatches = 0
+    var distMismatches = 0
+    var maxAbs: Float = 0
+    for i in 0..<args.n {
+        if outIdx[i] != cpuIdx[i] { idxMismatches += 1 }
+        let d = abs(outDist[i] - cpuDist[i])
+        if d > maxAbs { maxAbs = d }
+        if !distClose(outDist[i], cpuDist[i]) { distMismatches += 1 }
+    }
+    let countsOk = outCounts == cpuCounts
+    let ok = idxMismatches == 0 && distMismatches == 0 && countsOk
+    let sumIdx = outIdx.reduce(0) { $0 + Int($1) }
+
+    var result: [String: Any] = [
+        "impl": "hand_msl",
+        "device": device.name,
+        "n": args.n,
+        "r": args.r,
+        "m": args.m,
+        "seed": args.seed,
+        "warmup": args.warmup,
+        "reps": args.reps,
+        "e2e_ms": e2e,
+        "kernel_ms": kern,
+        "e2e_median_ms": median(e2e),
+        "e2e_min_ms": e2e.min() ?? 0,
+        "e2e_max_ms": e2e.max() ?? 0,
+        "kernel_median_ms": median(kern),
+        "kernel_min_ms": kern.min() ?? 0,
+        "kernel_max_ms": kern.max() ?? 0,
+        "swift_cpu_ref_ms": cpuMs,
+        "ok": ok,
+        "idx_mismatches": idxMismatches,
+        "dist_mismatches": distMismatches,
+        "max_abs_dist_delta": maxAbs,
+        "counts_ok": countsOk,
+        "sum_idx": sumIdx,
+        "counts_sum": Int(outCounts.reduce(0, +)),
+        "tolerance": ["abs": DIST_ABS, "rel": DIST_REL],
+    ]
+    if args.dump {
+        result["idx"] = outIdx.map { Int($0) }
+        result["dist"] = outDist
+        result["counts"] = outCounts.map { Int($0) }
+    }
+
+    let data = try! JSONSerialization.data(withJSONObject: result, options: [.prettyPrinted, .sortedKeys])
+    if let path = args.json {
+        try! data.write(to: URL(fileURLWithPath: path))
+    }
+    FileHandle.standardOutput.write(data)
+    FileHandle.standardOutput.write(Data("\n".utf8))
+
+    fputs("kind=hand_msl n=\(args.n) r=\(args.r) m=\(args.m) seed=\(args.seed)\n", stderr)
+    fputs("ok=\(ok) idx_mismatches=\(idxMismatches) dist_mismatches=\(distMismatches) max_abs=\(maxAbs) counts_ok=\(countsOk)\n", stderr)
+    fputs(String(format: "e2e_median_ms=%.4f kernel_median_ms=%.4f swift_cpu_ref_ms=%.4f sum_idx=%d\n",
+                 median(e2e), median(kern), cpuMs, sumIdx), stderr)
+    if args.dump {
+        print("ASSIGNMENTS")
+        for i in 0..<args.n {
+            print("\(outIdx[i]) \(outDist[i])")
+        }
+        print("COUNTS")
+        print(outCounts.map { String($0) }.joined(separator: " "))
+    }
+    exit(ok ? 0 : 1)
+}
+
+main()
diff --git a/experiments/metal-kernels/README.md b/experiments/metal-kernels/README.md
new file mode 100644
index 0000000..28f81f8
--- /dev/null
+++ b/experiments/metal-kernels/README.md
@@ -0,0 +1,56 @@
+# Experimental hand-written associate kernels
+
+Isolated experiment. **Not** on the product `Run` path. Nothing here is
+imported by `src/algorithm.bend`.
+
+U-NSGA-III association (see `src/survival.bend` `perp` / `nearest_go`):
+for each of `N` normalized objective rows and each of `R` reference
+directions, compute the perpendicular distance (right-fold F32 residual)
+and take the strict-`<` argmin (ties keep the lower index). Niche count
+is the histogram of those argmin indices.
+
+## Files
+
+| File | What |
+|------|------|
+| `associate.metal` | Metal Shading Language kernels (`associate`, `niche_count`) |
+| `associate.cu` | CUDA source with the same algorithm (compile-check only; no timings) |
+| `AssociateHarness.swift` | Metal host: LCG data, CPU ref, dispatch, e2e + kernel-only times |
+| `ref_associate.py` | Python F32 reference matching Bend's right-fold |
+| `assoc_map.bend` | Isolated map calling product `Surv.nearest_one` |
+| `bench_product.bend` | Current-path bench (no `!`) |
+| `bench_bang.bend` | Same map with `map_nearest.go!` (Bend GPU backend) |
+| `run_mac_bench.sh` | Compile + run on an Apple GPU machine |
+
+## How to run (Mac, Metal)
+
+```bash
+# from a unsga3-bend checkout, current bend on PATH (never pin)
+./experiments/metal-kernels/run_mac_bench.sh
+```
+
+The script writes `results/msl_*.json`, `results/bend_*.txt`, and
+`results/bench.md`. End-to-end Metal times include shared-buffer write,
+dispatch, wait, and a readback touch. Kernel-only times use
+`MTLCommandBuffer` GPU timestamps. Bend times are `IO.now()` around the
+associate map (implicit transfer included; no separate kernel-only
+counter).
+
+## Correctness
+
+`ref_associate.py --compare DUMP` checks argmin indices for exact
+equality and distances against abs `1e-6` / rel `1e-5`. The product
+`Run` IGD path is unchanged: this folder is not imported by it.
+
+## Timing note
+
+`assoc_map.bend` folds each timed repeat's argmin indices into `sum_idx_acc`.
+That sum must equal `reps * sum_idx`. Without it, Bend can drop unused
+repeats and the wall clock stops scaling with `reps`. See
+`results/TABLES.md`.
+
+## CUDA
+
+`associate.cu` is the same algorithm as `associate.metal`. This tree does
+not claim CUDA timings. `nvcc` is not required to check the rest of the
+experiment.
diff --git a/experiments/metal-kernels/assoc_map.bend b/experiments/metal-kernels/assoc_map.bend
new file mode 100644
index 0000000..0edd5e5
--- /dev/null
+++ b/experiments/metal-kernels/assoc_map.bend
@@ -0,0 +1,283 @@
+# Isolated associate map used by the Metal-kernel experiment.
+# Calls Surv.nearest_one (the product perp + argmin) over generated points.
+# Not imported by src/algorithm.bend / Run.
+import Base
+import ../../src/survival.bend as Surv
+import ../../src/util.bend as Util
+
+def lcg(s: U32) -> U32:
+  U32.add(U32.mul(s, U32.from_nat(1664525n)), U32.from_nat(1013904223n))
+
+def u32_to_unit(s: U32) -> F32:
+  F32.div(F32.from_nat(U32.to_nat(U32.shrn(s, 8n))), 16777216.0)
+
+def skip_lcg(fuel: Nat, s: U32) -> U32:
+  match fuel:
+    case 0n:
+      s
+    case 1n+k:
+      skip_lcg(k, lcg(s))
+
+def gen_vec.cons(+s2: U32, pair: List<&2, F32> & U32) -> List<&2, F32> & U32:
+  (xs, s3) = pair
+  (u32_to_unit(s2) <> xs, s3)
+
+def gen_vec(fuel: Nat, s: U32) -> List<&2, F32> & U32:
+  match fuel:
+    case 0n:
+      (Nil{}, s)
+    case 1n+k:
+      +s2 = lcg(s)
+      gen_vec.cons(s2, gen_vec(k, s2))
+
+def gen_mat.join(pair_vec: List<&2, F32> & U32, pair_rest: List<&2, List<&2, F32>> & U32) -> List<&2, List<&2, F32>> & U32:
+  (v, s2) = pair_vec
+  (vs, s3) = pair_rest
+  (v <> vs, s3)
+
+def gen_mat(rows: Nat, +cols: Nat, +s: U32) -> List<&2, List<&2, F32>> & U32:
+  match rows:
+    case 0n:
+      (Nil{}, s)
+    case 1n+k:
+      gen_mat.join(gen_vec(cols, s), gen_mat(k, cols, skip_lcg(cols, s)))
+
+def map_nearest.go(+dirs: List<&2, List<&2, F32>>, +fuel: Nat, pts: List<&2, List<&2, F32>>) -> List<&2, Surv.Assignment>:
+  match fuel pts:
+    case 0n Nil{}:
+      Nil{}
+    case 0n p <> rest:
+      Surv.nearest_one(p, dirs) <> map_nearest.go(dirs, 0n, rest)
+    case 1n+p Nil{}:
+      Nil{}
+    case 1n+p h <> Nil{}:
+      Surv.nearest_one(h, dirs) <> Nil{}
+    case 1n+p h <> t:
+      +all = {h <> t : List<&2, List<&2, F32>>}
+      +mid = Nat.div(Util.list_len(List<&2, F32>, all), 2n)
+      lo hi = map_nearest.go(dirs, p, List.take(&2, List<&2, F32>, all, mid)) map_nearest.go(dirs, p, List.drop(&2, List<&2, F32>, all, mid))
+      List.append(&2, Surv.Assignment, lo, hi)
+
+def map_nearest(+pts: List<&2, List<&2, F32>>, +dirs: List<&2, List<&2, F32>>) -> List<&2, Surv.Assignment>:
+  map_nearest.go(dirs, Util.list_len(List<&2, F32>, pts), pts)
+
+def map_nearest_bang(+pts: List<&2, List<&2, F32>>, +dirs: List<&2, List<&2, F32>>) -> List<&2, Surv.Assignment>:
+  map_nearest.go!(dirs, Util.list_len(List<&2, F32>, pts), pts)
+
+def repeat.once(+bang: Bool, +pts: List<&2, List<&2, F32>>, +dirs: List<&2, List<&2, F32>>) -> List<&2, Surv.Assignment>:
+  match bang:
+    case True{}:
+      map_nearest_bang(pts, dirs)
+    case False{}:
+      map_nearest(pts, dirs)
+
+def sum_idx(asgs: List<&2, Surv.Assignment>) -> Nat:
+  match asgs:
+    case Nil{}:
+      0n
+    case a <> rest:
+      Nat.add(Surv.asg_ref(a), sum_idx(rest))
+
+def repeat.use(+asgs: List<&2, Surv.Assignment>) -> List<&2, Surv.Assignment> & Nat:
+  (asgs, sum_idx(asgs))
+
+def repeat.join(+asgs: List<&2, Surv.Assignment>, pair: List<&2, Surv.Assignment> & Nat) -> List<&2, Surv.Assignment> & Nat:
+  (rest_asgs, rest_acc) = pair
+  (rest_asgs, Nat.add(sum_idx(asgs), rest_acc))
+
+# Each iteration's assignments are folded into acc so the compiler cannot
+# drop repeated maps as unused. acc must equal reps * sum_idx(last).
+def repeat(+fuel: Nat, +bang: Bool, +pts: List<&2, List<&2, F32>>, +dirs: List<&2, List<&2, F32>>) -> List<&2, Surv.Assignment> & Nat:
+  match fuel:
+    case 0n:
+      (Nil{}, 0n)
+    case 1n:
+      repeat.use(repeat.once(bang, pts, dirs))
+    case 1n+k:
+      repeat.join(repeat.once(bang, pts, dirs), repeat(k, bang, pts, dirs))
+
+def warmup_repeat(+fuel: Nat, +bang: Bool, +pts: List<&2, List<&2, F32>>, +dirs: List<&2, List<&2, F32>>) -> List<&2, Surv.Assignment>:
+  match fuel:
+    case 0n:
+      Nil{}
+    case _:
+      repeat.once(bang, pts, dirs)
+
+
+def hist.go(asgs: List<&2, Surv.Assignment>, acc: List<&2, Nat>) -> List<&2, Nat>:
+  match asgs:
+    case Nil{}:
+      acc
+    case a <> rest:
+      hist.go(rest, Util.inc_at(acc, Surv.asg_ref(a)))
+
+def hist(+asgs: List<&2, Surv.Assignment>, +nref: Nat) -> List<&2, Nat>:
+  hist.go(asgs, Util.zeros_nat(nref))
+
+def sum_nats(xs: List<&2, Nat>) -> Nat:
+  match xs:
+    case Nil{}:
+      0n
+    case x <> rest:
+      Nat.add(x, sum_nats(rest))
+
+def show_asgs(asgs: List<&2, Surv.Assignment>) -> IO(Unit):
+  match asgs:
+    case Nil{}:
+      IO.pure(Unit, Unit{})
+    case +a <> rest:
+      do IO<Unit>:
+        IO.print(Nat.show(Surv.asg_ref(a)) ++ " " ++ F32.show(Surv.asg_dist(a)))
+        show_asgs(rest)
+
+def show_counts.go(xs: List<&2, Nat>, acc: String) -> String:
+  match xs:
+    case Nil{}:
+      acc
+    case x <> Nil{}:
+      acc ++ Nat.show(x)
+    case x <> rest:
+      show_counts.go(rest, acc ++ Nat.show(x) ++ " ")
+
+def show_counts(xs: List<&2, Nat>) -> IO(Unit):
+  IO.print(show_counts.go(xs, ""))
+
+def run.dump(+dump: Nat, +asgs: List<&2, Surv.Assignment>, +counts: List<&2, Nat>) -> IO(Unit):
+  match dump:
+    case 0n:
+      IO.pure(Unit, Unit{})
+    case _:
+      do IO<Unit>:
+        IO.print("ASSIGNMENTS")
+        show_asgs(asgs)
+        IO.print("COUNTS")
+        show_counts(counts)
+
+def run.report(
+  +n: Nat,
+  +r: Nat,
+  +m: Nat,
+  +seed: Nat,
+  +warmup: Nat,
+  +reps: Nat,
+  +dump: Nat,
+  +bang: Bool,
+  +keep: Nat,
+  +elapsed: Nat,
+  +acc: Nat,
+  +asgs: List<&2, Surv.Assignment>,
+  +counts: List<&2, Nat>
+) -> IO(Unit):
+  do IO<Unit>:
+    IO.print("kind=" ++ Bool.pick(String, bang, "bend_bang", "bend_product"))
+    IO.print(
+      "n=" ++ Nat.show(n)
+      ++ " r=" ++ Nat.show(r)
+      ++ " m=" ++ Nat.show(m)
+      ++ " seed=" ++ Nat.show(seed)
+      ++ " keep=" ++ Nat.show(keep)
+    )
+    IO.print("warmup=" ++ Nat.show(warmup) ++ " reps=" ++ Nat.show(reps))
+    IO.print("total_ms=" ++ Nat.show(elapsed))
+    IO.print("sum_idx=" ++ Nat.show(sum_idx(asgs)))
+    IO.print("sum_idx_acc=" ++ Nat.show(acc))
+    IO.print("n_asg=" ++ Nat.show(Util.list_len(Surv.Assignment, asgs)))
+    IO.print("counts_sum=" ++ Nat.show(sum_nats(counts)))
+    run.dump(dump, asgs, counts)
+
+def run.after_time(
+  +n: Nat,
+  +r: Nat,
+  +m: Nat,
+  +seed: Nat,
+  +warmup: Nat,
+  +reps: Nat,
+  +dump: Nat,
+  +bang: Bool,
+  +keep: Nat,
+  +t0: Nat,
+  +t1: Nat,
+  +acc: Nat,
+  +asgs: List<&2, Surv.Assignment>
+) -> IO(Unit):
+  run.report(n, r, m, seed, warmup, reps, dump, bang, keep, Nat.sub(t1, t0), acc, asgs, hist(asgs, r))
+
+def run.after_pair(
+  +n: Nat,
+  +r: Nat,
+  +m: Nat,
+  +seed: Nat,
+  +warmup: Nat,
+  +reps: Nat,
+  +dump: Nat,
+  +bang: Bool,
+  +keep: Nat,
+  +t0: Nat,
+  +t1: Nat,
+  pair: List<&2, Surv.Assignment> & Nat
+) -> IO(Unit):
+  (asgs, acc) = pair
+  run.after_time(n, r, m, seed, warmup, reps, dump, bang, keep, t0, t1, acc, asgs)
+
+def run.emit(
+  +n: Nat,
+  +r: Nat,
+  +m: Nat,
+  +seed: Nat,
+  +warmup: Nat,
+  +reps: Nat,
+  +dump: Nat,
+  +bang: Bool,
+  +pts: List<&2, List<&2, F32>>,
+  +dirs: List<&2, List<&2, F32>>
+) -> IO(Unit):
+  do IO<Unit>:
+    w : List<&2, Surv.Assignment> = warmup_repeat(warmup, bang, pts, dirs)
+    keep : Nat = Util.list_len(Surv.Assignment, w)
+    t0 : Nat <- IO.now()
+    pair : List<&2, Surv.Assignment> & Nat = repeat(reps, bang, pts, dirs)
+    t1 : Nat <- IO.now()
+    run.after_pair(n, r, m, seed, warmup, reps, dump, bang, keep, t0, t1, pair)
+
+
+def run.after_dirs(
+  +n: Nat,
+  +r: Nat,
+  +m: Nat,
+  +seed: Nat,
+  +warmup: Nat,
+  +reps: Nat,
+  +dump: Nat,
+  +bang: Bool,
+  +pts: List<&2, List<&2, F32>>,
+  pair: List<&2, List<&2, F32>> & U32
+) -> IO(Unit):
+  (dirs, s2) = pair
+  run.emit(n, r, m, seed, warmup, reps, dump, bang, pts, dirs)
+
+def run.after_pts(
+  +n: Nat,
+  +r: Nat,
+  +m: Nat,
+  +seed: Nat,
+  +warmup: Nat,
+  +reps: Nat,
+  +dump: Nat,
+  +bang: Bool,
+  pair: List<&2, List<&2, F32>> & U32
+) -> IO(Unit):
+  (pts, s1) = pair
+  run.after_dirs(n, r, m, seed, warmup, reps, dump, bang, pts, gen_mat(r, m, s1))
+
+def run(
+  +n: Nat,
+  +r: Nat,
+  +m: Nat,
+  +seed: Nat,
+  +warmup: Nat,
+  +reps: Nat,
+  +dump: Nat,
+  +bang: Bool
+) -> IO(Unit):
+  run.after_pts(n, r, m, seed, warmup, reps, dump, bang, gen_mat(n, m, U32.from_nat(seed)))
+
diff --git a/experiments/metal-kernels/associate.cu b/experiments/metal-kernels/associate.cu
new file mode 100644
index 0000000..183e0d7
--- /dev/null
+++ b/experiments/metal-kernels/associate.cu
@@ -0,0 +1,78 @@
+// Experimental CUDA associate + niche-count kernels.
+// Same algorithm as associate.metal / src/survival.bend perp+nearest_go.
+// Not a product code path. No CUDA timings are claimed in this tree:
+// this file exists so the kernel can be compile-checked with nvcc when
+// a toolkit is present.
+#include <cuda_runtime.h>
+#include <stdint.h>
+
+__device__ inline float right_dot(const float* a, const float* b, unsigned m) {
+    float acc = 0.0f;
+    for (unsigned k = m; k > 0; --k) {
+        unsigned i = k - 1;
+        acc = a[i] * b[i] + acc;
+    }
+    return acc;
+}
+
+__device__ inline float right_resid_sq(const float* f, const float* w, float t, unsigned m) {
+    float acc = 0.0f;
+    for (unsigned k = m; k > 0; --k) {
+        unsigned i = k - 1;
+        float r = f[i] - t * w[i];
+        acc = r * r + acc;
+    }
+    return acc;
+}
+
+extern "C" __global__ void associate(
+    const float* points,
+    const float* refs,
+    unsigned* out_idx,
+    float* out_dist,
+    unsigned N,
+    unsigned R,
+    unsigned M
+) {
+    unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
+    if (gid >= N) return;
+    const float* f = points + gid * M;
+    unsigned best_i = 0;
+    float best_d = 1000000000.0f;
+    for (unsigned j = 0; j < R; ++j) {
+        const float* w = refs + j * M;
+        float ww = right_dot(w, w, M);
+        float d;
+        if (ww < 1.0e-16f) {
+            d = sqrtf(right_dot(f, f, M));
+        } else {
+            float t = right_dot(f, w, M) / ww;
+            d = sqrtf(right_resid_sq(f, w, t, M));
+        }
+        if (d < best_d) {
+            best_d = d;
+            best_i = j;
+        }
+    }
+    out_idx[gid] = best_i;
+    out_dist[gid] = best_d;
+}
+
+extern "C" __global__ void niche_count(
+    const unsigned* idx,
+    unsigned* counts,
+    unsigned N,
+    unsigned R
+) {
+    unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
+    if (gid >= N) return;
+    unsigned r = idx[gid];
+    if (r < R) atomicAdd(&counts[r], 1u);
+}
+
+// Host stub so `nvcc -c associate.cu` has a translation unit that links
+// the kernels (no device runtime is invoked).
+void associate_launch_stub() {
+    (void)associate;
+    (void)niche_count;
+}
diff --git a/experiments/metal-kernels/associate.metal b/experiments/metal-kernels/associate.metal
new file mode 100644
index 0000000..c4f1706
--- /dev/null
+++ b/experiments/metal-kernels/associate.metal
@@ -0,0 +1,76 @@
+// Experimental associate + niche-count kernels.
+// Matches unsga3-bend src/survival.bend perp / nearest_go / resid_sq:
+//   ww = ||w||^2                     (right-fold dot)
+//   if ww < 1e-16: d = ||f||         (right-fold)
+//   else t = dot(f,w)/ww; d = sqrt(sum_k (f_k - t*w_k)^2)  (right-fold)
+//   argmin on strict < (ties keep the lower index)
+// Not a product code path.
+#include <metal_stdlib>
+using namespace metal;
+
+inline float right_dot(const device float* a, const device float* b, uint m) {
+    float acc = 0.0f;
+    // Right-fold: a0*b0 + (a1*b1 + (... + 0))
+    for (uint k = m; k > 0; --k) {
+        uint i = k - 1;
+        acc = a[i] * b[i] + acc;
+    }
+    return acc;
+}
+
+inline float right_resid_sq(const device float* f, const device float* w, float t, uint m) {
+    float acc = 0.0f;
+    for (uint k = m; k > 0; --k) {
+        uint i = k - 1;
+        float r = f[i] - t * w[i];
+        acc = r * r + acc;
+    }
+    return acc;
+}
+
+kernel void associate(
+    device const float* points [[buffer(0)]],
+    device const float* refs   [[buffer(1)]],
+    device uint*        out_idx [[buffer(2)]],
+    device float*       out_dist [[buffer(3)]],
+    constant uint&      N [[buffer(4)]],
+    constant uint&      R [[buffer(5)]],
+    constant uint&      M [[buffer(6)]],
+    uint gid [[thread_position_in_grid]]
+) {
+    if (gid >= N) return;
+    const device float* f = points + gid * M;
+    uint best_i = 0;
+    float best_d = 1000000000.0f;
+    for (uint j = 0; j < R; ++j) {
+        const device float* w = refs + j * M;
+        float ww = right_dot(w, w, M);
+        float d;
+        if (ww < 1.0e-16f) {
+            d = sqrt(right_dot(f, f, M));
+        } else {
+            float t = right_dot(f, w, M) / ww;
+            d = sqrt(right_resid_sq(f, w, t, M));
+        }
+        if (d < best_d) {
+            best_d = d;
+            best_i = j;
+        }
+    }
+    out_idx[gid] = best_i;
+    out_dist[gid] = best_d;
+}
+
+kernel void niche_count(
+    device const uint*        idx    [[buffer(0)]],
+    device atomic_uint*       counts [[buffer(1)]],
+    constant uint&            N      [[buffer(2)]],
+    constant uint&            R      [[buffer(3)]],
+    uint gid [[thread_position_in_grid]]
+) {
+    if (gid >= N) return;
+    uint r = idx[gid];
+    if (r < R) {
+        atomic_fetch_add_explicit(&counts[r], 1u, memory_order_relaxed);
+    }
+}
diff --git a/experiments/metal-kernels/bench_bang.bend b/experiments/metal-kernels/bench_bang.bend
new file mode 100644
index 0000000..1f164a3
--- /dev/null
+++ b/experiments/metal-kernels/bench_bang.bend
@@ -0,0 +1,7 @@
+# Bend GPU-bang associate bench (`map_nearest.go!`). Isolated experiment, not Run.
+# Default matches the small product fixture so outputs can be compared.
+import Base
+import ./assoc_map.bend as M
+
+def main() -> IO(Unit):
+  M.run(8n, 5n, 3n, 1n, 0n, 1n, 1n, True{})
diff --git a/experiments/metal-kernels/bench_product.bend b/experiments/metal-kernels/bench_product.bend
new file mode 100644
index 0000000..4bd9a01
--- /dev/null
+++ b/experiments/metal-kernels/bench_product.bend
@@ -0,0 +1,8 @@
+# Current-path associate bench: no GPU bang. Isolated experiment, not Run.
+# Default: small fixture (n=8 r=5 m=3 seed=1) with a full assignment dump.
+# Size-specific wrappers generated by run_mac_bench.sh call M.run directly.
+import Base
+import ./assoc_map.bend as M
+
+def main() -> IO(Unit):
+  M.run(8n, 5n, 3n, 1n, 0n, 1n, 1n, False{})
diff --git a/experiments/metal-kernels/combine_results.py b/experiments/metal-kernels/combine_results.py
new file mode 100755
index 0000000..47e8e74
--- /dev/null
+++ b/experiments/metal-kernels/combine_results.py
@@ -0,0 +1,100 @@
+#!/usr/bin/env python3
+"""Combine Swift/Metal and Bend bench logs into results JSON + markdown."""
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+
+def parse_bend(text: str) -> dict:
+    out: dict = {}
+    for line in text.splitlines():
+        if not line or line.startswith("ASSIGNMENTS") or line.startswith("COUNTS"):
+            continue
+        if "ALL PROOFS" in line or line.startswith("Use --verdict"):
+            continue
+        for part in line.split():
+            if "=" not in part:
+                continue
+            k, v = part.split("=", 1)
+            if k in {
+                "n", "r", "m", "seed", "warmup", "reps", "total_ms",
+                "sum_idx", "n_asg", "counts_sum", "keep",
+            }:
+                try:
+                    out[k] = int(v)
+                except ValueError:
+                    out[k] = v
+            else:
+                out[k] = v
+    if "total_ms" in out and out.get("reps"):
+        out["per_ms"] = out["total_ms"] / out["reps"]
+    return out
+
+
+def main() -> int:
+    here = Path(__file__).resolve().parent
+    res = here / "results"
+    res.mkdir(exist_ok=True)
+    rows = []
+    payload = {"cells": []}
+    for path in sorted(res.glob("msl_*.json")):
+        data = json.loads(path.read_text(encoding="utf-8"))
+        data["_file"] = path.name
+        payload["cells"].append(data)
+        size = f"{data.get('n')}x{data.get('r')}x{data.get('m')}"
+        rows.append({
+            "size": size,
+            "impl": data.get("impl", "hand_msl"),
+            "e2e_median_ms": data.get("e2e_median_ms"),
+            "e2e_min_ms": data.get("e2e_min_ms"),
+            "e2e_max_ms": data.get("e2e_max_ms"),
+            "kernel_median_ms": data.get("kernel_median_ms"),
+            "kernel_min_ms": data.get("kernel_min_ms"),
+            "kernel_max_ms": data.get("kernel_max_ms"),
+            "reps": data.get("reps"),
+            "sum_idx": data.get("sum_idx"),
+            "ok": data.get("ok"),
+        })
+    for path in sorted(res.glob("bend_*.txt")):
+        parsed = parse_bend(path.read_text(encoding="utf-8"))
+        parsed["_file"] = path.name
+        payload["cells"].append(parsed)
+        size = f"{parsed.get('n')}x{parsed.get('r')}x{parsed.get('m')}"
+        rows.append({
+            "size": size,
+            "impl": parsed.get("kind", path.stem),
+            "e2e_median_ms": parsed.get("per_ms"),
+            "e2e_min_ms": None,
+            "e2e_max_ms": None,
+            "kernel_median_ms": None,
+            "note": f"total_ms={parsed.get('total_ms')} over {parsed.get('reps')} in-process reps (Bend IO.now)",
+            "reps": parsed.get("reps"),
+            "sum_idx": parsed.get("sum_idx"),
+            "ok": None,
+        })
+    payload["table_rows"] = rows
+    (res / "bench.json").write_text(json.dumps(payload, indent=2) + "\n", encoding="utf-8")
+    lines = [
+        "| size (N x R x M) | impl | e2e (ms) | e2e min-max (ms) | kernel-only (ms) | reps | sum_idx | ok |",
+        "|---|---|---:|---|---:|---:|---:|---|",
+    ]
+    for r in rows:
+        e2e = r.get("e2e_median_ms")
+        e2e_s = f"{e2e:.4f}" if isinstance(e2e, float) else ("" if e2e is None else str(e2e))
+        if r.get("e2e_min_ms") is not None:
+            spread = f"{r['e2e_min_ms']:.4f}-{r['e2e_max_ms']:.4f}"
+        else:
+            spread = r.get("note", "")
+        kern = r.get("kernel_median_ms")
+        kern_s = f"{kern:.4f}" if isinstance(kern, float) else "n/a"
+        lines.append(
+            f"| {r.get('size')} | {r.get('impl')} | {e2e_s} | {spread} | {kern_s} | {r.get('reps')} | {r.get('sum_idx')} | {r.get('ok')} |"
+        )
+    (res / "bench.md").write_text("\n".join(lines) + "\n", encoding="utf-8")
+    print(f"wrote {res / 'bench.json'} and {res / 'bench.md'} ({len(rows)} rows)")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
diff --git a/experiments/metal-kernels/ref_associate.py b/experiments/metal-kernels/ref_associate.py
new file mode 100755
index 0000000..74feb6c
--- /dev/null
+++ b/experiments/metal-kernels/ref_associate.py
@@ -0,0 +1,225 @@
+#!/usr/bin/env python3
+"""CPU reference for the experimental associate kernel.
+
+Matches src/survival.bend `resid_sq` / `perp` / `nearest_go` F32 order:
+right-fold dots and residuals, ww < 1e-16 uses ||f||, argmin on strict `<`
+(ties keep the lower index). Also builds the niche-count histogram.
+
+This is not a product code path. Used to check the Metal/CUDA kernels and
+the isolated Bend bench dumps.
+"""
+from __future__ import annotations
+
+import argparse
+import json
+import struct
+import sys
+from pathlib import Path
+
+
+LCG_MUL = 1664525
+LCG_ADD = 1013904223
+BEST_D0 = 1_000_000_000.0
+TINY_WW = 1e-16
+DIST_ABS = 1e-6
+DIST_REL = 1e-5
+
+
+def lcg(state: int) -> int:
+    return (state * LCG_MUL + LCG_ADD) & 0xFFFFFFFF
+
+
+def u32_to_unit(state: int) -> float:
+    # (state >> 8) / 2^24 as binary32. Integers up to 2^24 are exact in F32.
+    u = (state >> 8) & 0xFFFFFF
+    return struct.unpack("f", struct.pack("f", u / 16777216.0))[0]
+
+
+def gen_matrix(rows: int, cols: int, seed: int) -> tuple[list[list[float]], int]:
+    s = seed & 0xFFFFFFFF
+    out: list[list[float]] = []
+    for _ in range(rows):
+        row: list[float] = []
+        for _ in range(cols):
+            s = lcg(s)
+            row.append(u32_to_unit(s))
+        out.append(row)
+    return out, s
+
+
+def f32(x: float) -> float:
+    return struct.unpack("f", struct.pack("f", x))[0]
+
+
+def right_dot(a: list[float], b: list[float]) -> float:
+    acc = 0.0
+    for i in range(len(a) - 1, -1, -1):
+        acc = f32(f32(a[i] * b[i]) + acc)
+    return acc
+
+
+def right_resid_sq(f: list[float], w: list[float], t: float) -> float:
+    acc = 0.0
+    for i in range(len(f) - 1, -1, -1):
+        r = f32(f[i] - f32(t * w[i]))
+        acc = f32(f32(r * r) + acc)
+    return acc
+
+
+def perp(f: list[float], w: list[float]) -> float:
+    ww = right_dot(w, w)
+    if ww < TINY_WW:
+        import math
+        return f32(math.sqrt(right_dot(f, f)))
+    t = f32(right_dot(f, w) / ww)
+    import math
+    return f32(math.sqrt(right_resid_sq(f, w, t)))
+
+
+def nearest_one(f: list[float], refs: list[list[float]]) -> tuple[int, float]:
+    best_i = 0
+    best_d = BEST_D0
+    for j, w in enumerate(refs):
+        d = perp(f, w)
+        if d < best_d:
+            best_d = d
+            best_i = j
+    return best_i, best_d
+
+
+def associate(points: list[list[float]], refs: list[list[float]]) -> tuple[list[int], list[float]]:
+    idx: list[int] = []
+    dist: list[float] = []
+    for f in points:
+        i, d = nearest_one(f, refs)
+        idx.append(i)
+        dist.append(d)
+    return idx, dist
+
+
+def niche_count(idx: list[int], n_ref: int) -> list[int]:
+    counts = [0] * n_ref
+    for r in idx:
+        if 0 <= r < n_ref:
+            counts[r] += 1
+    return counts
+
+
+def dist_close(a: float, b: float) -> bool:
+    if abs(a - b) <= DIST_ABS:
+        return True
+    scale = max(abs(a), abs(b), 1e-30)
+    return abs(a - b) <= DIST_REL * scale
+
+
+def compare(
+    idx_a: list[int],
+    dist_a: list[float],
+    idx_b: list[int],
+    dist_b: list[float],
+    label: str,
+) -> dict:
+    n = min(len(idx_a), len(idx_b))
+    mismatches = [i for i in range(n) if idx_a[i] != idx_b[i]]
+    dist_bad = [i for i in range(n) if not dist_close(dist_a[i], dist_b[i])]
+    max_abs = max((abs(dist_a[i] - dist_b[i]) for i in range(n)), default=0.0)
+    ok = (
+        len(idx_a) == len(idx_b)
+        and not mismatches
+        and not dist_bad
+    )
+    return {
+        "label": label,
+        "ok": ok,
+        "n": n,
+        "idx_mismatches": len(mismatches),
+        "idx_mismatch_examples": mismatches[:8],
+        "dist_mismatches": len(dist_bad),
+        "dist_mismatch_examples": dist_bad[:8],
+        "max_abs_dist_delta": max_abs,
+        "sum_idx": sum(idx_a),
+        "tolerance": {"abs": DIST_ABS, "rel": DIST_REL},
+    }
+
+
+def parse_dump(text: str) -> tuple[list[int], list[float], list[int]]:
+    idx: list[int] = []
+    dist: list[float] = []
+    counts: list[int] = []
+    mode = "head"
+    for raw in text.splitlines():
+        line = raw.strip()
+        if not line or line.startswith("#") or line.startswith("kind="):
+            continue
+        if line == "ASSIGNMENTS":
+            mode = "asg"
+            continue
+        if line == "COUNTS":
+            mode = "counts"
+            continue
+        if mode == "asg":
+            parts = line.split()
+            if len(parts) >= 2:
+                idx.append(int(parts[0]))
+                dist.append(float(parts[1]))
+        elif mode == "counts":
+            counts.extend(int(x) for x in line.split())
+    return idx, dist, counts
+
+
+def main() -> int:
+    p = argparse.ArgumentParser(description=__doc__)
+    p.add_argument("--n", type=int, default=8)
+    p.add_argument("--r", type=int, default=5)
+    p.add_argument("--m", type=int, default=3)
+    p.add_argument("--seed", type=int, default=1)
+    p.add_argument("--dump", type=Path, default=None)
+    p.add_argument("--compare", type=Path, default=None, help="Bend/Metal dump to check")
+    p.add_argument("--include-zero-ref", action="store_true")
+    args = p.parse_args()
+
+    points, s1 = gen_matrix(args.n, args.m, args.seed)
+    refs, _ = gen_matrix(args.r, args.m, s1)
+    if args.include_zero_ref:
+        refs[0] = [0.0] * args.m
+
+    idx, dist = associate(points, refs)
+    counts = niche_count(idx, len(refs))
+
+    payload = {
+        "n": args.n,
+        "r": len(refs),
+        "m": args.m,
+        "seed": args.seed,
+        "sum_idx": sum(idx),
+        "counts_sum": sum(counts),
+        "idx": idx,
+        "dist": dist,
+        "counts": counts,
+    }
+    if args.dump:
+        args.dump.write_text(json.dumps(payload, indent=2) + "\n", encoding="utf-8")
+        print(f"wrote {args.dump}", file=sys.stderr)
+
+    print(f"n={args.n} r={len(refs)} m={args.m} seed={args.seed}")
+    print(f"sum_idx={sum(idx)} counts_sum={sum(counts)}")
+    print("ASSIGNMENTS")
+    for i, d in zip(idx, dist):
+        print(f"{i} {d:.9g}")
+    print("COUNTS")
+    print(" ".join(str(c) for c in counts))
+
+    if args.compare:
+        text = args.compare.read_text(encoding="utf-8")
+        c_idx, c_dist, c_counts = parse_dump(text)
+        report = compare(idx, dist, c_idx, c_dist, args.compare.name)
+        if c_counts and c_counts != counts:
+            report["ok"] = False
+            report["count_mismatch"] = True
+        print(json.dumps(report, indent=2), file=sys.stderr)
+        return 0 if report["ok"] else 1
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
diff --git a/experiments/metal-kernels/results/TABLES.md b/experiments/metal-kernels/results/TABLES.md
new file mode 100644
index 0000000..d9e041f
--- /dev/null
+++ b/experiments/metal-kernels/results/TABLES.md
@@ -0,0 +1,70 @@
+# Associate kernel benches (2026-11-10, America/New_York)
+
+Host: JMac-mini.local, Apple M4, 10-core GPU, Metal 4, 16 GB unified.
+`bend version` = `bend 2.0.40` (updated via install.sh that night; started at 2.0.34).
+Apple clang 21.0.0 / Swift 6.4.2. Load averages at start: 2.20 2.23 2.12.
+Not an IGD claim. Not a product `Run` path.
+
+## What was timed
+
+- **hand MSL:** `AssociateHarness.swift` + `associate.metal`. End-to-end includes
+  shared-buffer rewrite, encode, commit, wait, and a readback touch.
+  Kernel-only is `MTLCommandBuffer.gpuEndTime - gpuStartTime` (associate +
+  niche_count in one command buffer). Warmup then individual reps; median and
+  min-max are reported. Library compile is outside the timed region.
+- **Bend product:** `assoc_map.bend` `map_nearest` calling product
+  `Surv.nearest_one` (no `!`), `+RTS --threads 1`. Time is `IO.now()` around
+  the repeated map (integer milliseconds). Per-rep = `total_ms / reps`.
+- **Bend Metal:** the same map with `map_nearest.go!`, GPU on (default).
+- **Bend bang CPU:** same `!` binary with `+RTS --threads 1 --gpu off`.
+
+An earlier Bend loop dropped unused repeats (total time did not scale with
+`reps`). The numbers below use the forced-work loop: each repeat's argmin
+indices are folded into `sum_idx_acc`, which must equal `reps * sum_idx`.
+Pre-fix logs remain in this folder as `*_hi.txt` / un-suffixed files and
+are not used for claims.
+
+## Primary table
+
+| N x R x M | impl | e2e | kernel-only | reps | sum_idx | work check |
+|-----------|------|-----|-------------|-----:|--------:|------------|
+| 184 x 91 x 3 (DTLZ2 combined pool) | hand MSL | median **0.365 ms** (0.335-0.369) | median **0.116 ms** (0.111-0.116) | 7 after 2 warm | 8819 | idx exact vs Swift CPU ref; max abs dist 5.6e-8 |
+| 184 x 91 x 3 | Bend product (current path) | **0.550 ms/rep** (44 ms / 80) | n/a | 80 after 1 warm | 8819 | sum_idx_acc=705520 = 80 x 8819 |
+| 184 x 91 x 3 | Bend Metal (`!`) | **6.55 ms/rep** (131 ms / 20) | n/a (Bend has no GPU timestamp) | 20 after 1 warm | 8819 | sum_idx_acc=176380 = 20 x 8819 |
+| 184 x 91 x 3 | Bend bang, `--gpu off` | **0.500 ms/rep** (10 ms / 20) | n/a | 20 after 1 warm | 8819 | sum_idx_acc=176380 = 20 x 8819 |
+| 2048 x 276 x 3 (mid) | hand MSL | median **0.392 ms** (0.384-0.433) | median **0.216 ms** (all 0.216-0.216) | 7 after 1 warm | 279827 | idx exact; max abs dist 4.8e-8 |
+| 2048 x 276 x 3 | Bend product | **19.8 ms/rep** (198 ms / 10) | n/a | 10 after 1 warm | 279827 | sum_idx_acc=2798270 = 10 x 279827 |
+| 2048 x 276 x 3 | Bend Metal (`!`) | **crash** | — | — | — | `bend: memory fault (machine stack overflow?)` |
+| 2048 x 276 x 3 | Bend bang, `--gpu off` | **18.6 ms/rep** (149 ms / 8) | n/a | 8 after 1 warm | 279827 | sum_idx_acc=2238616 = 8 x 279827 |
+| 32000 x 3003 x 3 (big-N) | hand MSL | median **7.934 ms** (7.928-10.577) | median **7.675 ms** (7.665-10.229) | 5 after 1 warm | 47882676 | idx exact; max abs dist 5.6e-8 |
+| 32000 x 3003 x 3 | Bend product | **2195 ms/rep** (6586 ms / 3) | n/a | 3 after 1 warm | 47882676 | sum_idx_acc=143648028 = 3 x 47882676 |
+| 32000 x 3003 x 3 | Bend Metal (`!`) | **crash** | — | — | — | same stack overflow |
+| 32000 x 3003 x 3 | Bend bang, `--gpu off` | **2236 ms/rep** (4471 ms / 2) | n/a | 2, no extra warm | 47882676 | sum_idx_acc=95765352 = 2 x 47882676 |
+
+Speedup of hand MSL e2e vs Bend product: **1.51x** at 184 x 91, **50.5x** at 2048 x 276, **277x** at 32000 x 3003.
+
+Speedup of hand MSL e2e vs Bend Metal: **17.9x** at 184 x 91 (the only size where Bend `!` completed). Larger sizes: Bend `!` did not run.
+
+Supporting one-shot Swift CPU ref (same process, used for correctness, not a 5-rep median): 0.108 ms at 184 x 91, 2.961 ms at 2048 x 276, 481.9 ms at 32000 x 3003.
+
+## Fixture
+
+`n=8 r=5 m=3 seed=1` (including Python `ref_associate.py` vs Bend product on the box, and vs hand MSL on the Mac): argmin indices identical (`4 2 3 3 4 3 3 1`, sum_idx=23), niche counts identical (`0 1 1 4 2`), max abs distance delta 5.4e-9 (Python vs Bend) and 3.0e-8 (Python vs MSL). The 8 x 5 MSL e2e of 3.32 ms had **no warmup** (first dispatch) and is not a performance claim.
+
+## Extra MSL sizes vs the Swift CPU ref (break-even vs a flat C-like CPU)
+
+Same harness, warmup 1, 7 reps, M=3, seed=1. Swift CPU is one-shot.
+
+| N x R | MSL e2e median (min-max) | MSL kernel median | Swift CPU (1-shot) | e2e vs Swift CPU |
+|-------|--------------------------|-------------------|--------------------|------------------|
+| 128 x 91 | 0.306 (0.293-0.341) | 0.108 | 0.089 | MSL slower (dispatch) |
+| 184 x 91 | 0.365 (0.335-0.369) | 0.116 | 0.108 | MSL slower |
+| 256 x 91 | 0.738 (0.403-0.793) | 0.122 | 0.131 | noisy; kernel ~tied |
+| 512 x 91 | 0.362 (0.317-0.743) | 0.114 | 0.262 | e2e still close; kernel 2.3x |
+| 512 x 276 | 0.997 (0.743-1.054) | 0.352 | 0.746 | kernel 2.1x; e2e still behind |
+| 1024 x 276 | 0.964 (0.923-1.027) | 0.351 | 1.488 | **e2e 1.54x** |
+| 2048 x 276 | 0.392 (0.384-0.433) | 0.216 | 2.961 | **e2e 7.6x** |
+| 8000 x 1050 | 2.507 (2.230-2.530) | 1.890 | 43.67 | **e2e 17.4x** |
+| 32000 x 3003 | 7.934 (7.928-10.577) | 7.675 | 481.9 | **e2e 60.7x** |
+
+Break-even vs the current Bend CPU path is already at the realistic DTLZ2 pool (184 x 91). Break-even vs a flat Swift CPU is between 512 x 276 (MSL e2e still slower) and 1024 x 276 (MSL e2e 1.54x).
diff --git a/experiments/metal-kernels/results/bench.json b/experiments/metal-kernels/results/bench.json
new file mode 100644
index 0000000..8a91ea0
--- /dev/null
+++ b/experiments/metal-kernels/results/bench.json
@@ -0,0 +1,835 @@
+{
+  "cells": [
+    {
+      "counts_ok": true,
+      "counts_sum": 184,
+      "device": "Apple M4",
+      "dist_mismatches": 0,
+      "e2e_max_ms": 0.36895275115966797,
+      "e2e_median_ms": 0.3650188446044922,
+      "e2e_min_ms": 0.3349781036376953,
+      "e2e_ms": [
+        0.3650188446044922,
+        0.3560781478881836,
+        0.3459453582763672,
+        0.3660917282104492,
+        0.36895275115966797,
+        0.36704540252685547,
+        0.3349781036376953
+      ],
+      "idx_mismatches": 0,
+      "impl": "hand_msl",
+      "kernel_max_ms": 0.11600001016631722,
+      "kernel_median_ms": 0.11575000826269388,
+      "kernel_min_ms": 0.11100003030151129,
+      "kernel_ms": [
+        0.11575000826269388,
+        0.11550000635907054,
+        0.11100003030151129,
+        0.1154583296738565,
+        0.11600001016631722,
+        0.11600001016631722,
+        0.11600001016631722
+      ],
+      "m": 3,
+      "max_abs_dist_delta": 5.587935447692871e-08,
+      "n": 184,
+      "ok": true,
+      "r": 91,
+      "reps": 7,
+      "seed": 1,
+      "sum_idx": 8819,
+      "swift_cpu_ref_ms": 0.10800361633300781,
+      "tolerance": {
+        "abs": 9.999999974752427e-07,
+        "rel": 9.999999747378752e-06
+      },
+      "warmup": 2,
+      "_file": "msl_n184r91.json"
+    },
+    {
+      "counts_ok": true,
+      "counts_sum": 2048,
+      "device": "Apple M4",
+      "dist_mismatches": 0,
+      "e2e_max_ms": 0.4329681396484375,
+      "e2e_median_ms": 0.39196014404296875,
+      "e2e_min_ms": 0.3839731216430664,
+      "e2e_ms": [
+        0.4329681396484375,
+        0.4119873046875,
+        0.38993358612060547,
+        0.3839731216430664,
+        0.39196014404296875,
+        0.3930330276489258,
+        0.3910064697265625
+      ],
+      "idx_mismatches": 0,
+      "impl": "hand_msl",
+      "kernel_max_ms": 0.21641666535288095,
+      "kernel_median_ms": 0.2162500168196857,
+      "kernel_min_ms": 0.2162500168196857,
+      "kernel_ms": [
+        0.2162500168196857,
+        0.21641666535288095,
+        0.2162500168196857,
+        0.2162500168196857,
+        0.2162500168196857,
+        0.2162500168196857,
+        0.2162500168196857
+      ],
+      "m": 3,
+      "max_abs_dist_delta": 4.842877388000488e-08,
+      "n": 2048,
+      "ok": true,
+      "r": 276,
+      "reps": 7,
+      "seed": 1,
+      "sum_idx": 279827,
+      "swift_cpu_ref_ms": 2.9610395431518555,
+      "tolerance": {
+        "abs": 9.999999974752427e-07,
+        "rel": 9.999999747378752e-06
+      },
+      "warmup": 1,
+      "_file": "msl_n2048r276.json"
+    },
+    {
+      "counts_ok": true,
+      "counts_sum": 32000,
+      "device": "Apple M4",
+      "dist_mismatches": 0,
+      "e2e_max_ms": 10.576963424682617,
+      "e2e_median_ms": 7.933974266052246,
+      "e2e_min_ms": 7.928013801574707,
+      "e2e_ms": [
+        10.576963424682617,
+        8.71896743774414,
+        7.930994033813477,
+        7.928013801574707,
+        7.933974266052246
+      ],
+      "idx_mismatches": 0,
+      "impl": "hand_msl",
+      "kernel_max_ms": 10.228624974843115,
+      "kernel_median_ms": 7.674624968785793,
+      "kernel_min_ms": 7.665375014767051,
+      "kernel_ms": [
+        10.228624974843115,
+        8.468249987345189,
+        7.665375014767051,
+        7.674624968785793,
+        7.670416671317071
+      ],
+      "m": 3,
+      "max_abs_dist_delta": 5.587935447692871e-08,
+      "n": 32000,
+      "ok": true,
+      "r": 3003,
+      "reps": 5,
+      "seed": 1,
+      "sum_idx": 47882676,
+      "swift_cpu_ref_ms": 481.8829298019409,
+      "tolerance": {
+        "abs": 9.999999974752427e-07,
+        "rel": 9.999999747378752e-06
+      },
+      "warmup": 1,
+      "_file": "msl_n32kr3003.json"
+    },
+    {
+      "counts": [
+        0,
+        1,
+        1,
+        4,
+        2
+      ],
+      "counts_ok": true,
+      "counts_sum": 8,
+      "device": "Apple M4",
+      "dist": [
+        0.13612471520900726,
+        0.14139918982982635,
+        0.4722365438938141,
+        0.05906146392226219,
+        0.47954317927360535,
+        0.5190921425819397,
+        0.308963805437088,
+        0.05956927686929703
+      ],
+      "dist_mismatches": 0,
+      "e2e_max_ms": 3.3190250396728516,
+      "e2e_median_ms": 3.3190250396728516,
+      "e2e_min_ms": 3.3190250396728516,
+      "e2e_ms": [
+        3.3190250396728516
+      ],
+      "idx": [
+        4,
+        2,
+        3,
+        3,
+        4,
+        3,
+        3,
+        1
+      ],
+      "idx_mismatches": 0,
+      "impl": "hand_msl",
+      "kernel_max_ms": 0.01087499549612403,
+      "kernel_median_ms": 0.01087499549612403,
+      "kernel_min_ms": 0.01087499549612403,
+      "kernel_ms": [
+        0.01087499549612403
+      ],
+      "m": 3,
+      "max_abs_dist_delta": 2.9802322387695312e-08,
+      "n": 8,
+      "ok": true,
+      "r": 5,
+      "reps": 1,
+      "seed": 1,
+      "sum_idx": 23,
+      "swift_cpu_ref_ms": 0.0010728836059570312,
+      "tolerance": {
+        "abs": 9.999999974752427e-07,
+        "rel": 9.999999747378752e-06
+      },
+      "warmup": 0,
+      "_file": "msl_n8r5.json"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 2,
+      "reps": 7,
+      "total_ms": 7,
+      "sum_idx": 8819,
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 1.0,
+      "_file": "bend_bang_n184r91.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 1,
+      "reps": 20,
+      "total_ms": 131,
+      "sum_idx": 8819,
+      "sum_idx_acc": "176380",
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 6.55,
+      "_file": "bend_bang_n184r91_fw.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 5,
+      "reps": 50,
+      "total_ms": 7,
+      "sum_idx": 8819,
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 0.14,
+      "_file": "bend_bang_n184r91_hi.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 2048,
+      "r": 276,
+      "m": 3,
+      "seed": 1,
+      "error": "memory",
+      "_file": "bend_bang_n2048r276.txt"
+    },
+    {
+      "_file": "bend_bang_n2048r276_fw.txt"
+    },
+    {
+      "rc": "0",
+      "_file": "bend_bang_n32kr3003.txt"
+    },
+    {
+      "_file": "bend_bang_n32kr3003_fw.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 2,
+      "reps": 7,
+      "total_ms": 0,
+      "sum_idx": 8819,
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 0.0,
+      "_file": "bend_bangcpu_n184r91.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 1,
+      "reps": 20,
+      "total_ms": 10,
+      "sum_idx": 8819,
+      "sum_idx_acc": "176380",
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 0.5,
+      "_file": "bend_bangcpu_n184r91_fw.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 5,
+      "reps": 50,
+      "total_ms": 0,
+      "sum_idx": 8819,
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 0.0,
+      "_file": "bend_bangcpu_n184r91_hi.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 2048,
+      "r": 276,
+      "m": 3,
+      "seed": 1,
+      "keep": 2048,
+      "warmup": 1,
+      "reps": 7,
+      "total_ms": 14,
+      "sum_idx": 279827,
+      "n_asg": 2048,
+      "counts_sum": 2048,
+      "per_ms": 2.0,
+      "_file": "bend_bangcpu_n2048r276.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 2048,
+      "r": 276,
+      "m": 3,
+      "seed": 1,
+      "keep": 2048,
+      "warmup": 1,
+      "reps": 8,
+      "total_ms": 149,
+      "sum_idx": 279827,
+      "sum_idx_acc": "2238616",
+      "n_asg": 2048,
+      "counts_sum": 2048,
+      "per_ms": 18.625,
+      "_file": "bend_bangcpu_n2048r276_fw.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 32000,
+      "r": 3003,
+      "m": 3,
+      "seed": 1,
+      "keep": 0,
+      "warmup": 0,
+      "reps": 1,
+      "total_ms": 2004,
+      "sum_idx": 47882676,
+      "n_asg": 32000,
+      "counts_sum": 32000,
+      "per_ms": 2004.0,
+      "_file": "bend_bangcpu_n32kr3003.txt"
+    },
+    {
+      "kind": "bend_bang",
+      "n": 32000,
+      "r": 3003,
+      "m": 3,
+      "seed": 1,
+      "keep": 0,
+      "warmup": 0,
+      "reps": 2,
+      "total_ms": 4471,
+      "sum_idx": 47882676,
+      "sum_idx_acc": "95765352",
+      "n_asg": 32000,
+      "counts_sum": 32000,
+      "per_ms": 2235.5,
+      "_file": "bend_bangcpu_n32kr3003_fw.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 2,
+      "reps": 7,
+      "total_ms": 0,
+      "sum_idx": 8819,
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 0.0,
+      "_file": "bend_product_n184r91.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 1,
+      "reps": 80,
+      "total_ms": 44,
+      "sum_idx": 8819,
+      "sum_idx_acc": "705520",
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 0.55,
+      "_file": "bend_product_n184r91_fw.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 184,
+      "r": 91,
+      "m": 3,
+      "seed": 1,
+      "keep": 184,
+      "warmup": 5,
+      "reps": 400,
+      "total_ms": 1,
+      "sum_idx": 8819,
+      "n_asg": 184,
+      "counts_sum": 184,
+      "per_ms": 0.0025,
+      "_file": "bend_product_n184r91_hi.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 2048,
+      "r": 276,
+      "m": 3,
+      "seed": 1,
+      "keep": 2048,
+      "warmup": 1,
+      "reps": 7,
+      "total_ms": 12,
+      "sum_idx": 279827,
+      "n_asg": 2048,
+      "counts_sum": 2048,
+      "per_ms": 1.7142857142857142,
+      "_file": "bend_product_n2048r276.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 2048,
+      "r": 276,
+      "m": 3,
+      "seed": 1,
+      "keep": 2048,
+      "warmup": 1,
+      "reps": 10,
+      "total_ms": 198,
+      "sum_idx": 279827,
+      "sum_idx_acc": "2798270",
+      "n_asg": 2048,
+      "counts_sum": 2048,
+      "per_ms": 19.8,
+      "_file": "bend_product_n2048r276_fw.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 2048,
+      "r": 276,
+      "m": 3,
+      "seed": 1,
+      "keep": 2048,
+      "warmup": 2,
+      "reps": 30,
+      "total_ms": 12,
+      "sum_idx": 279827,
+      "n_asg": 2048,
+      "counts_sum": 2048,
+      "per_ms": 0.4,
+      "_file": "bend_product_n2048r276_hi.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 32000,
+      "r": 3003,
+      "m": 3,
+      "seed": 1,
+      "keep": 32000,
+      "warmup": 1,
+      "reps": 3,
+      "total_ms": 2015,
+      "sum_idx": 47882676,
+      "n_asg": 32000,
+      "counts_sum": 32000,
+      "per_ms": 671.6666666666666,
+      "_file": "bend_product_n32kr3003.txt"
+    },
+    {
+      "kind": "bend_product",
+      "n": 32000,
+      "r": 3003,
+      "m": 3,
+      "seed": 1,
+      "keep": 32000,
+      "warmup": 1,
+      "reps": 3,
+      "total_ms": 6586,
+      "sum_idx": 47882676,
+      "sum_idx_acc": "143648028",
+      "n_asg": 32000,
+      "counts_sum": 32000,
+      "per_ms": 2195.3333333333335,
+      "_file": "bend_product_n32kr3003_fw.txt"
+    }
+  ],
+  "table_rows": [
+    {
+      "size": "184x91x3",
+      "impl": "hand_msl",
+      "e2e_median_ms": 0.3650188446044922,
+      "e2e_min_ms": 0.3349781036376953,
+      "e2e_max_ms": 0.36895275115966797,
+      "kernel_median_ms": 0.11575000826269388,
+      "kernel_min_ms": 0.11100003030151129,
+      "kernel_max_ms": 0.11600001016631722,
+      "reps": 7,
+      "sum_idx": 8819,
+      "ok": true
+    },
+    {
+      "size": "2048x276x3",
+      "impl": "hand_msl",
+      "e2e_median_ms": 0.39196014404296875,
+      "e2e_min_ms": 0.3839731216430664,
+      "e2e_max_ms": 0.4329681396484375,
+      "kernel_median_ms": 0.2162500168196857,
+      "kernel_min_ms": 0.2162500168196857,
+      "kernel_max_ms": 0.21641666535288095,
+      "reps": 7,
+      "sum_idx": 279827,
+      "ok": true
+    },
+    {
+      "size": "32000x3003x3",
+      "impl": "hand_msl",
+      "e2e_median_ms": 7.933974266052246,
+      "e2e_min_ms": 7.928013801574707,
+      "e2e_max_ms": 10.576963424682617,
+      "kernel_median_ms": 7.674624968785793,
+      "kernel_min_ms": 7.665375014767051,
+      "kernel_max_ms": 10.228624974843115,
+      "reps": 5,
+      "sum_idx": 47882676,
+      "ok": true
+    },
+    {
+      "size": "8x5x3",
+      "impl": "hand_msl",
+      "e2e_median_ms": 3.3190250396728516,
+      "e2e_min_ms": 3.3190250396728516,
+      "e2e_max_ms": 3.3190250396728516,
+      "kernel_median_ms": 0.01087499549612403,
+      "kernel_min_ms": 0.01087499549612403,
+      "kernel_max_ms": 0.01087499549612403,
+      "reps": 1,
+      "sum_idx": 23,
+      "ok": true
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 1.0,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=7 over 7 in-process reps (Bend IO.now)",
+      "reps": 7,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 6.55,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=131 over 20 in-process reps (Bend IO.now)",
+      "reps": 20,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 0.14,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=7 over 50 in-process reps (Bend IO.now)",
+      "reps": 50,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "2048x276x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": null,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=None over None in-process reps (Bend IO.now)",
+      "reps": null,
+      "sum_idx": null,
+      "ok": null
+    },
+    {
+      "size": "NonexNonexNone",
+      "impl": "bend_bang_n2048r276_fw",
+      "e2e_median_ms": null,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=None over None in-process reps (Bend IO.now)",
+      "reps": null,
+      "sum_idx": null,
+      "ok": null
+    },
+    {
+      "size": "NonexNonexNone",
+      "impl": "bend_bang_n32kr3003",
+      "e2e_median_ms": null,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=None over None in-process reps (Bend IO.now)",
+      "reps": null,
+      "sum_idx": null,
+      "ok": null
+    },
+    {
+      "size": "NonexNonexNone",
+      "impl": "bend_bang_n32kr3003_fw",
+      "e2e_median_ms": null,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=None over None in-process reps (Bend IO.now)",
+      "reps": null,
+      "sum_idx": null,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 0.0,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=0 over 7 in-process reps (Bend IO.now)",
+      "reps": 7,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 0.5,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=10 over 20 in-process reps (Bend IO.now)",
+      "reps": 20,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 0.0,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=0 over 50 in-process reps (Bend IO.now)",
+      "reps": 50,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "2048x276x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 2.0,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=14 over 7 in-process reps (Bend IO.now)",
+      "reps": 7,
+      "sum_idx": 279827,
+      "ok": null
+    },
+    {
+      "size": "2048x276x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 18.625,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=149 over 8 in-process reps (Bend IO.now)",
+      "reps": 8,
+      "sum_idx": 279827,
+      "ok": null
+    },
+    {
+      "size": "32000x3003x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 2004.0,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=2004 over 1 in-process reps (Bend IO.now)",
+      "reps": 1,
+      "sum_idx": 47882676,
+      "ok": null
+    },
+    {
+      "size": "32000x3003x3",
+      "impl": "bend_bang",
+      "e2e_median_ms": 2235.5,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=4471 over 2 in-process reps (Bend IO.now)",
+      "reps": 2,
+      "sum_idx": 47882676,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 0.0,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=0 over 7 in-process reps (Bend IO.now)",
+      "reps": 7,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 0.55,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=44 over 80 in-process reps (Bend IO.now)",
+      "reps": 80,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "184x91x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 0.0025,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=1 over 400 in-process reps (Bend IO.now)",
+      "reps": 400,
+      "sum_idx": 8819,
+      "ok": null
+    },
+    {
+      "size": "2048x276x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 1.7142857142857142,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=12 over 7 in-process reps (Bend IO.now)",
+      "reps": 7,
+      "sum_idx": 279827,
+      "ok": null
+    },
+    {
+      "size": "2048x276x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 19.8,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=198 over 10 in-process reps (Bend IO.now)",
+      "reps": 10,
+      "sum_idx": 279827,
+      "ok": null
+    },
+    {
+      "size": "2048x276x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 0.4,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=12 over 30 in-process reps (Bend IO.now)",
+      "reps": 30,
+      "sum_idx": 279827,
+      "ok": null
+    },
+    {
+      "size": "32000x3003x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 671.6666666666666,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=2015 over 3 in-process reps (Bend IO.now)",
+      "reps": 3,
+      "sum_idx": 47882676,
+      "ok": null
+    },
+    {
+      "size": "32000x3003x3",
+      "impl": "bend_product",
+      "e2e_median_ms": 2195.3333333333335,
+      "e2e_min_ms": null,
+      "e2e_max_ms": null,
+      "kernel_median_ms": null,
+      "note": "total_ms=6586 over 3 in-process reps (Bend IO.now)",
+      "reps": 3,
+      "sum_idx": 47882676,
+      "ok": null
+    }
+  ]
+}
diff --git a/experiments/metal-kernels/results/bench.md b/experiments/metal-kernels/results/bench.md
new file mode 100644
index 0000000..a1ae60e
--- /dev/null
+++ b/experiments/metal-kernels/results/bench.md
@@ -0,0 +1,28 @@
+| size (N x R x M) | impl | e2e (ms) | e2e min-max (ms) | kernel-only (ms) | reps | sum_idx | ok |
+|---|---|---:|---|---:|---:|---:|---|
+| 184x91x3 | hand_msl | 0.3650 | 0.3350-0.3690 | 0.1158 | 7 | 8819 | True |
+| 2048x276x3 | hand_msl | 0.3920 | 0.3840-0.4330 | 0.2163 | 7 | 279827 | True |
+| 32000x3003x3 | hand_msl | 7.9340 | 7.9280-10.5770 | 7.6746 | 5 | 47882676 | True |
+| 8x5x3 | hand_msl | 3.3190 | 3.3190-3.3190 | 0.0109 | 1 | 23 | True |
+| 184x91x3 | bend_bang | 1.0000 | total_ms=7 over 7 in-process reps (Bend IO.now) | n/a | 7 | 8819 | None |
+| 184x91x3 | bend_bang | 6.5500 | total_ms=131 over 20 in-process reps (Bend IO.now) | n/a | 20 | 8819 | None |
+| 184x91x3 | bend_bang | 0.1400 | total_ms=7 over 50 in-process reps (Bend IO.now) | n/a | 50 | 8819 | None |
+| 2048x276x3 | bend_bang |  | total_ms=None over None in-process reps (Bend IO.now) | n/a | None | None | None |
+| NonexNonexNone | bend_bang_n2048r276_fw |  | total_ms=None over None in-process reps (Bend IO.now) | n/a | None | None | None |
+| NonexNonexNone | bend_bang_n32kr3003 |  | total_ms=None over None in-process reps (Bend IO.now) | n/a | None | None | None |
+| NonexNonexNone | bend_bang_n32kr3003_fw |  | total_ms=None over None in-process reps (Bend IO.now) | n/a | None | None | None |
+| 184x91x3 | bend_bang | 0.0000 | total_ms=0 over 7 in-process reps (Bend IO.now) | n/a | 7 | 8819 | None |
+| 184x91x3 | bend_bang | 0.5000 | total_ms=10 over 20 in-process reps (Bend IO.now) | n/a | 20 | 8819 | None |
+| 184x91x3 | bend_bang | 0.0000 | total_ms=0 over 50 in-process reps (Bend IO.now) | n/a | 50 | 8819 | None |
+| 2048x276x3 | bend_bang | 2.0000 | total_ms=14 over 7 in-process reps (Bend IO.now) | n/a | 7 | 279827 | None |
+| 2048x276x3 | bend_bang | 18.6250 | total_ms=149 over 8 in-process reps (Bend IO.now) | n/a | 8 | 279827 | None |
+| 32000x3003x3 | bend_bang | 2004.0000 | total_ms=2004 over 1 in-process reps (Bend IO.now) | n/a | 1 | 47882676 | None |
+| 32000x3003x3 | bend_bang | 2235.5000 | total_ms=4471 over 2 in-process reps (Bend IO.now) | n/a | 2 | 47882676 | None |
+| 184x91x3 | bend_product | 0.0000 | total_ms=0 over 7 in-process reps (Bend IO.now) | n/a | 7 | 8819 | None |
+| 184x91x3 | bend_product | 0.5500 | total_ms=44 over 80 in-process reps (Bend IO.now) | n/a | 80 | 8819 | None |
+| 184x91x3 | bend_product | 0.0025 | total_ms=1 over 400 in-process reps (Bend IO.now) | n/a | 400 | 8819 | None |
+| 2048x276x3 | bend_product | 1.7143 | total_ms=12 over 7 in-process reps (Bend IO.now) | n/a | 7 | 279827 | None |
+| 2048x276x3 | bend_product | 19.8000 | total_ms=198 over 10 in-process reps (Bend IO.now) | n/a | 10 | 279827 | None |
+| 2048x276x3 | bend_product | 0.4000 | total_ms=12 over 30 in-process reps (Bend IO.now) | n/a | 30 | 279827 | None |
+| 32000x3003x3 | bend_product | 671.6667 | total_ms=2015 over 3 in-process reps (Bend IO.now) | n/a | 3 | 47882676 | None |
+| 32000x3003x3 | bend_product | 2195.3333 | total_ms=6586 over 3 in-process reps (Bend IO.now) | n/a | 3 | 47882676 | None |
diff --git a/experiments/metal-kernels/results/bend_bang_n184r91.err b/experiments/metal-kernels/results/bend_bang_n184r91.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bang_n184r91.txt b/experiments/metal-kernels/results/bend_bang_n184r91.txt
new file mode 100644
index 0000000..8ccfe03
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n184r91.txt
@@ -0,0 +1,7 @@
+kind=bend_bang
+n=184 r=91 m=3 seed=1 keep=184
+warmup=2 reps=7
+total_ms=7
+sum_idx=8819
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_bang_n184r91_fw.err b/experiments/metal-kernels/results/bend_bang_n184r91_fw.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bang_n184r91_fw.txt b/experiments/metal-kernels/results/bend_bang_n184r91_fw.txt
new file mode 100644
index 0000000..b0b5ec0
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n184r91_fw.txt
@@ -0,0 +1,8 @@
+kind=bend_bang
+n=184 r=91 m=3 seed=1 keep=184
+warmup=1 reps=20
+total_ms=131
+sum_idx=8819
+sum_idx_acc=176380
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_bang_n184r91_hi.err b/experiments/metal-kernels/results/bend_bang_n184r91_hi.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bang_n184r91_hi.txt b/experiments/metal-kernels/results/bend_bang_n184r91_hi.txt
new file mode 100644
index 0000000..862787a
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n184r91_hi.txt
@@ -0,0 +1,7 @@
+kind=bend_bang
+n=184 r=91 m=3 seed=1 keep=184
+warmup=5 reps=50
+total_ms=7
+sum_idx=8819
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_bang_n2048r276.err b/experiments/metal-kernels/results/bend_bang_n2048r276.err
new file mode 100644
index 0000000..71e857b
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n2048r276.err
@@ -0,0 +1 @@
+bend: memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bang_n2048r276.txt b/experiments/metal-kernels/results/bend_bang_n2048r276.txt
new file mode 100644
index 0000000..b3c349d
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n2048r276.txt
@@ -0,0 +1,2 @@
+kind=bend_bang n=2048 r=276 m=3 seed=1
+error=memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bang_n2048r276_fw.err b/experiments/metal-kernels/results/bend_bang_n2048r276_fw.err
new file mode 100644
index 0000000..71e857b
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n2048r276_fw.err
@@ -0,0 +1 @@
+bend: memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bang_n2048r276_fw.txt b/experiments/metal-kernels/results/bend_bang_n2048r276_fw.txt
new file mode 100644
index 0000000..15aa49a
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n2048r276_fw.txt
@@ -0,0 +1,2 @@
+FAILED
+bend: memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bang_n32kr3003.err b/experiments/metal-kernels/results/bend_bang_n32kr3003.err
new file mode 100644
index 0000000..71e857b
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n32kr3003.err
@@ -0,0 +1 @@
+bend: memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bang_n32kr3003.txt b/experiments/metal-kernels/results/bend_bang_n32kr3003.txt
new file mode 100644
index 0000000..5dc59be
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n32kr3003.txt
@@ -0,0 +1,2 @@
+FAILED rc=0
+bend: memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bang_n32kr3003_fw.err b/experiments/metal-kernels/results/bend_bang_n32kr3003_fw.err
new file mode 100644
index 0000000..71e857b
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n32kr3003_fw.err
@@ -0,0 +1 @@
+bend: memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bang_n32kr3003_fw.txt b/experiments/metal-kernels/results/bend_bang_n32kr3003_fw.txt
new file mode 100644
index 0000000..15aa49a
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bang_n32kr3003_fw.txt
@@ -0,0 +1,2 @@
+FAILED
+bend: memory fault (machine stack overflow?)
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n184r91.err b/experiments/metal-kernels/results/bend_bangcpu_n184r91.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n184r91.txt b/experiments/metal-kernels/results/bend_bangcpu_n184r91.txt
new file mode 100644
index 0000000..8c11c92
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bangcpu_n184r91.txt
@@ -0,0 +1,7 @@
+kind=bend_bang
+n=184 r=91 m=3 seed=1 keep=184
+warmup=2 reps=7
+total_ms=0
+sum_idx=8819
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n184r91_fw.err b/experiments/metal-kernels/results/bend_bangcpu_n184r91_fw.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n184r91_fw.txt b/experiments/metal-kernels/results/bend_bangcpu_n184r91_fw.txt
new file mode 100644
index 0000000..0b31618
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bangcpu_n184r91_fw.txt
@@ -0,0 +1,8 @@
+kind=bend_bang
+n=184 r=91 m=3 seed=1 keep=184
+warmup=1 reps=20
+total_ms=10
+sum_idx=8819
+sum_idx_acc=176380
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n184r91_hi.err b/experiments/metal-kernels/results/bend_bangcpu_n184r91_hi.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n184r91_hi.txt b/experiments/metal-kernels/results/bend_bangcpu_n184r91_hi.txt
new file mode 100644
index 0000000..3cfa1b6
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bangcpu_n184r91_hi.txt
@@ -0,0 +1,7 @@
+kind=bend_bang
+n=184 r=91 m=3 seed=1 keep=184
+warmup=5 reps=50
+total_ms=0
+sum_idx=8819
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n2048r276.err b/experiments/metal-kernels/results/bend_bangcpu_n2048r276.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n2048r276.txt b/experiments/metal-kernels/results/bend_bangcpu_n2048r276.txt
new file mode 100644
index 0000000..3540b25
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bangcpu_n2048r276.txt
@@ -0,0 +1,7 @@
+kind=bend_bang
+n=2048 r=276 m=3 seed=1 keep=2048
+warmup=1 reps=7
+total_ms=14
+sum_idx=279827
+n_asg=2048
+counts_sum=2048
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n2048r276_fw.err b/experiments/metal-kernels/results/bend_bangcpu_n2048r276_fw.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n2048r276_fw.txt b/experiments/metal-kernels/results/bend_bangcpu_n2048r276_fw.txt
new file mode 100644
index 0000000..5da2947
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bangcpu_n2048r276_fw.txt
@@ -0,0 +1,8 @@
+kind=bend_bang
+n=2048 r=276 m=3 seed=1 keep=2048
+warmup=1 reps=8
+total_ms=149
+sum_idx=279827
+sum_idx_acc=2238616
+n_asg=2048
+counts_sum=2048
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n32kr3003.err b/experiments/metal-kernels/results/bend_bangcpu_n32kr3003.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n32kr3003.txt b/experiments/metal-kernels/results/bend_bangcpu_n32kr3003.txt
new file mode 100644
index 0000000..6830994
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bangcpu_n32kr3003.txt
@@ -0,0 +1,7 @@
+kind=bend_bang
+n=32000 r=3003 m=3 seed=1 keep=0
+warmup=0 reps=1
+total_ms=2004
+sum_idx=47882676
+n_asg=32000
+counts_sum=32000
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n32kr3003_fw.err b/experiments/metal-kernels/results/bend_bangcpu_n32kr3003_fw.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_bangcpu_n32kr3003_fw.txt b/experiments/metal-kernels/results/bend_bangcpu_n32kr3003_fw.txt
new file mode 100644
index 0000000..125cef6
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_bangcpu_n32kr3003_fw.txt
@@ -0,0 +1,8 @@
+kind=bend_bang
+n=32000 r=3003 m=3 seed=1 keep=0
+warmup=0 reps=2
+total_ms=4471
+sum_idx=47882676
+sum_idx_acc=95765352
+n_asg=32000
+counts_sum=32000
diff --git a/experiments/metal-kernels/results/bend_product_n184r91.err b/experiments/metal-kernels/results/bend_product_n184r91.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n184r91.txt b/experiments/metal-kernels/results/bend_product_n184r91.txt
new file mode 100644
index 0000000..b81aab4
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n184r91.txt
@@ -0,0 +1,7 @@
+kind=bend_product
+n=184 r=91 m=3 seed=1 keep=184
+warmup=2 reps=7
+total_ms=0
+sum_idx=8819
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_product_n184r91_fw.err b/experiments/metal-kernels/results/bend_product_n184r91_fw.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n184r91_fw.txt b/experiments/metal-kernels/results/bend_product_n184r91_fw.txt
new file mode 100644
index 0000000..5851d6b
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n184r91_fw.txt
@@ -0,0 +1,8 @@
+kind=bend_product
+n=184 r=91 m=3 seed=1 keep=184
+warmup=1 reps=80
+total_ms=44
+sum_idx=8819
+sum_idx_acc=705520
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_product_n184r91_hi.err b/experiments/metal-kernels/results/bend_product_n184r91_hi.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n184r91_hi.txt b/experiments/metal-kernels/results/bend_product_n184r91_hi.txt
new file mode 100644
index 0000000..4a4a91f
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n184r91_hi.txt
@@ -0,0 +1,7 @@
+kind=bend_product
+n=184 r=91 m=3 seed=1 keep=184
+warmup=5 reps=400
+total_ms=1
+sum_idx=8819
+n_asg=184
+counts_sum=184
diff --git a/experiments/metal-kernels/results/bend_product_n2048r276.err b/experiments/metal-kernels/results/bend_product_n2048r276.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n2048r276.txt b/experiments/metal-kernels/results/bend_product_n2048r276.txt
new file mode 100644
index 0000000..37ee5f1
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n2048r276.txt
@@ -0,0 +1,7 @@
+kind=bend_product
+n=2048 r=276 m=3 seed=1 keep=2048
+warmup=1 reps=7
+total_ms=12
+sum_idx=279827
+n_asg=2048
+counts_sum=2048
diff --git a/experiments/metal-kernels/results/bend_product_n2048r276_fw.err b/experiments/metal-kernels/results/bend_product_n2048r276_fw.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n2048r276_fw.txt b/experiments/metal-kernels/results/bend_product_n2048r276_fw.txt
new file mode 100644
index 0000000..54b9074
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n2048r276_fw.txt
@@ -0,0 +1,8 @@
+kind=bend_product
+n=2048 r=276 m=3 seed=1 keep=2048
+warmup=1 reps=10
+total_ms=198
+sum_idx=279827
+sum_idx_acc=2798270
+n_asg=2048
+counts_sum=2048
diff --git a/experiments/metal-kernels/results/bend_product_n2048r276_hi.err b/experiments/metal-kernels/results/bend_product_n2048r276_hi.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n2048r276_hi.txt b/experiments/metal-kernels/results/bend_product_n2048r276_hi.txt
new file mode 100644
index 0000000..631f16e
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n2048r276_hi.txt
@@ -0,0 +1,7 @@
+kind=bend_product
+n=2048 r=276 m=3 seed=1 keep=2048
+warmup=2 reps=30
+total_ms=12
+sum_idx=279827
+n_asg=2048
+counts_sum=2048
diff --git a/experiments/metal-kernels/results/bend_product_n32kr3003.err b/experiments/metal-kernels/results/bend_product_n32kr3003.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n32kr3003.txt b/experiments/metal-kernels/results/bend_product_n32kr3003.txt
new file mode 100644
index 0000000..f4a41ac
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n32kr3003.txt
@@ -0,0 +1,7 @@
+kind=bend_product
+n=32000 r=3003 m=3 seed=1 keep=32000
+warmup=1 reps=3
+total_ms=2015
+sum_idx=47882676
+n_asg=32000
+counts_sum=32000
diff --git a/experiments/metal-kernels/results/bend_product_n32kr3003_fw.err b/experiments/metal-kernels/results/bend_product_n32kr3003_fw.err
new file mode 100644
index 0000000..e69de29
diff --git a/experiments/metal-kernels/results/bend_product_n32kr3003_fw.txt b/experiments/metal-kernels/results/bend_product_n32kr3003_fw.txt
new file mode 100644
index 0000000..5f57bd1
--- /dev/null
+++ b/experiments/metal-kernels/results/bend_product_n32kr3003_fw.txt
@@ -0,0 +1,8 @@
+kind=bend_product
+n=32000 r=3003 m=3 seed=1 keep=32000
+warmup=1 reps=3
+total_ms=6586
+sum_idx=47882676
+sum_idx_acc=143648028
+n_asg=32000
+counts_sum=32000
diff --git a/experiments/metal-kernels/results/fixture_python_vs_msl.stdout.txt b/experiments/metal-kernels/results/fixture_python_vs_msl.stdout.txt
new file mode 100644
index 0000000..ce92ab3
--- /dev/null
+++ b/experiments/metal-kernels/results/fixture_python_vs_msl.stdout.txt
@@ -0,0 +1,13 @@
+n=8 r=5 m=3 seed=1
+sum_idx=23 counts_sum=8
+ASSIGNMENTS
+4 0.13612473
+2 0.141399205
+3 0.472236544
+3 0.0590614639
+4 0.479543149
+3 0.519092143
+3 0.308963805
+1 0.0595692731
+COUNTS
+0 1 1 4 2
diff --git a/experiments/metal-kernels/results/mac_env.txt b/experiments/metal-kernels/results/mac_env.txt
new file mode 100644
index 0000000..593bc24
--- /dev/null
+++ b/experiments/metal-kernels/results/mac_env.txt
@@ -0,0 +1,13 @@
+host=Jason’s Mac mini
+date=Tue Nov 10 00:48:22 EST 2026
+uname=Darwin JMac-mini.local 27.0.0 Darwin Kernel Version 27.0.0: Tue Aug 11 21:02:57 PDT 2026; root:xnu-13432.1.9~1/RELEASE_ARM64_T8132 arm64
+Apple M4
+17179869184
+ 0:48  up 4 days,  9:58, 1 user, load averages: 2.20 2.23 2.12
+/Users/jkb/.bend/bin/bend
+bend 2.0.34
+/usr/bin/swiftc
+Apple Swift version 6.4.2 (swiftlang-6.4.2.1.9 clang-2100.4.1.1)
+/usr/bin/clang
+Apple clang version 21.0.0 (clang-2100.4.1.1)
+bend 2.0.40
diff --git a/experiments/metal-kernels/results/msl_n1024r276.json b/experiments/metal-kernels/results/msl_n1024r276.json
new file mode 100644
index 0000000..8c0dd72
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n1024r276.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 1024,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 1.0269880294799805,
+  "e2e_median_ms" : 0.96392631530761719,
+  "e2e_min_ms" : 0.92303752899169922,
+  "e2e_ms" : [
+    0.95796585083007812,
+    0.98097324371337891,
+    1.0039806365966797,
+    1.0269880294799805,
+    0.946044921875,
+    0.96392631530761719,
+    0.92303752899169922
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.35195838427171111,
+  "kernel_median_ms" : 0.35149999894201756,
+  "kernel_min_ms" : 0.34058332676067948,
+  "kernel_ms" : [
+    0.34374999813735485,
+    0.35187497269362211,
+    0.35149999894201756,
+    0.35149999894201756,
+    0.34058332676067948,
+    0.34775002859532833,
+    0.35195838427171111
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 4.4703483581542969e-08,
+  "n" : 1024,
+  "ok" : true,
+  "r" : 276,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 139707,
+  "swift_cpu_ref_ms" : 1.4879703521728516,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n128r91.json b/experiments/metal-kernels/results/msl_n128r91.json
new file mode 100644
index 0000000..7e14a81
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n128r91.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 128,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 0.34093856811523438,
+  "e2e_median_ms" : 0.30601024627685547,
+  "e2e_min_ms" : 0.29289722442626953,
+  "e2e_ms" : [
+    0.34093856811523438,
+    0.31805038452148438,
+    0.30601024627685547,
+    0.29706954956054688,
+    0.29289722442626953,
+    0.29504299163818359,
+    0.31101703643798828
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.10849995305761695,
+  "kernel_median_ms" : 0.10824995115399361,
+  "kernel_min_ms" : 0.10695832315832376,
+  "kernel_ms" : [
+    0.10825000936165452,
+    0.10849995305761695,
+    0.10841671610251069,
+    0.10799994925037026,
+    0.10699999984353781,
+    0.10695832315832376,
+    0.10824995115399361
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 2.9802322387695312e-08,
+  "n" : 128,
+  "ok" : true,
+  "r" : 91,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 5449,
+  "swift_cpu_ref_ms" : 0.088930130004882812,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n184r91.json b/experiments/metal-kernels/results/msl_n184r91.json
new file mode 100644
index 0000000..741cf13
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n184r91.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 184,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 0.36895275115966797,
+  "e2e_median_ms" : 0.36501884460449219,
+  "e2e_min_ms" : 0.33497810363769531,
+  "e2e_ms" : [
+    0.36501884460449219,
+    0.35607814788818359,
+    0.34594535827636719,
+    0.36609172821044922,
+    0.36895275115966797,
+    0.36704540252685547,
+    0.33497810363769531
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.11600001016631722,
+  "kernel_median_ms" : 0.11575000826269388,
+  "kernel_min_ms" : 0.11100003030151129,
+  "kernel_ms" : [
+    0.11575000826269388,
+    0.11550000635907054,
+    0.11100003030151129,
+    0.1154583296738565,
+    0.11600001016631722,
+    0.11600001016631722,
+    0.11600001016631722
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 5.5879354476928711e-08,
+  "n" : 184,
+  "ok" : true,
+  "r" : 91,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 8819,
+  "swift_cpu_ref_ms" : 0.10800361633300781,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 2
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n2048r276.json b/experiments/metal-kernels/results/msl_n2048r276.json
new file mode 100644
index 0000000..a980a66
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n2048r276.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 2048,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 0.4329681396484375,
+  "e2e_median_ms" : 0.39196014404296875,
+  "e2e_min_ms" : 0.38397312164306641,
+  "e2e_ms" : [
+    0.4329681396484375,
+    0.4119873046875,
+    0.38993358612060547,
+    0.38397312164306641,
+    0.39196014404296875,
+    0.39303302764892578,
+    0.3910064697265625
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.21641666535288095,
+  "kernel_median_ms" : 0.2162500168196857,
+  "kernel_min_ms" : 0.2162500168196857,
+  "kernel_ms" : [
+    0.2162500168196857,
+    0.21641666535288095,
+    0.2162500168196857,
+    0.2162500168196857,
+    0.2162500168196857,
+    0.2162500168196857,
+    0.2162500168196857
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 4.8428773880004883e-08,
+  "n" : 2048,
+  "ok" : true,
+  "r" : 276,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 279827,
+  "swift_cpu_ref_ms" : 2.9610395431518555,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n256r91.json b/experiments/metal-kernels/results/msl_n256r91.json
new file mode 100644
index 0000000..c3af75e
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n256r91.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 256,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 0.79298019409179688,
+  "e2e_median_ms" : 0.73802471160888672,
+  "e2e_min_ms" : 0.40304660797119141,
+  "e2e_ms" : [
+    0.78701972961425781,
+    0.79298019409179688,
+    0.73802471160888672,
+    0.77795982360839844,
+    0.69701671600341797,
+    0.44095516204833984,
+    0.40304660797119141
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.12275000335648656,
+  "kernel_median_ms" : 0.12254167813807726,
+  "kernel_min_ms" : 0.11495832586660981,
+  "kernel_ms" : [
+    0.12270838487893343,
+    0.12254167813807726,
+    0.12116663856431842,
+    0.12254167813807726,
+    0.12275000335648656,
+    0.11495832586660981,
+    0.11791661381721497
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 4.4703483581542969e-08,
+  "n" : 256,
+  "ok" : true,
+  "r" : 91,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 10252,
+  "swift_cpu_ref_ms" : 0.13101100921630859,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n32kr3003.json b/experiments/metal-kernels/results/msl_n32kr3003.json
new file mode 100644
index 0000000..cc2310e
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n32kr3003.json
@@ -0,0 +1,42 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 32000,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 10.576963424682617,
+  "e2e_median_ms" : 7.9339742660522461,
+  "e2e_min_ms" : 7.928013801574707,
+  "e2e_ms" : [
+    10.576963424682617,
+    8.7189674377441406,
+    7.9309940338134766,
+    7.928013801574707,
+    7.9339742660522461
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 10.228624974843115,
+  "kernel_median_ms" : 7.6746249687857926,
+  "kernel_min_ms" : 7.6653750147670507,
+  "kernel_ms" : [
+    10.228624974843115,
+    8.4682499873451889,
+    7.6653750147670507,
+    7.6746249687857926,
+    7.6704166713170707
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 5.5879354476928711e-08,
+  "n" : 32000,
+  "ok" : true,
+  "r" : 3003,
+  "reps" : 5,
+  "seed" : 1,
+  "sum_idx" : 47882676,
+  "swift_cpu_ref_ms" : 481.88292980194092,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n512r276.json b/experiments/metal-kernels/results/msl_n512r276.json
new file mode 100644
index 0000000..bdece7e
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n512r276.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 512,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 1.0540485382080078,
+  "e2e_median_ms" : 0.99694728851318359,
+  "e2e_min_ms" : 0.74303150177001953,
+  "e2e_ms" : [
+    0.99694728851318359,
+    0.98097324371337891,
+    0.96094608306884766,
+    0.74303150177001953,
+    1.0540485382080078,
+    0.99897384643554688,
+    1.0240077972412109
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.35241671139374375,
+  "kernel_median_ms" : 0.35158335231244564,
+  "kernel_min_ms" : 0.33962499583140016,
+  "kernel_ms" : [
+    0.34745829179883003,
+    0.35183329600840807,
+    0.3515416756272316,
+    0.33962499583140016,
+    0.35191664937883615,
+    0.35241671139374375,
+    0.35158335231244564
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 4.8428773880004883e-08,
+  "n" : 512,
+  "ok" : true,
+  "r" : 276,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 69997,
+  "swift_cpu_ref_ms" : 0.74601173400878906,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n512r91.json b/experiments/metal-kernels/results/msl_n512r91.json
new file mode 100644
index 0000000..57affff
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n512r91.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 512,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 0.74303150177001953,
+  "e2e_median_ms" : 0.36203861236572266,
+  "e2e_min_ms" : 0.31697750091552734,
+  "e2e_ms" : [
+    0.74303150177001953,
+    0.38492679595947266,
+    0.31697750091552734,
+    0.36406517028808594,
+    0.32007694244384766,
+    0.36203861236572266,
+    0.34701824188232422
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.1220832928083837,
+  "kernel_median_ms" : 0.11449999874457717,
+  "kernel_min_ms" : 0.10962499072775245,
+  "kernel_ms" : [
+    0.1220832928083837,
+    0.11449999874457717,
+    0.10962499072775245,
+    0.10987499263137579,
+    0.109999964479357,
+    0.11599995195865631,
+    0.116499955765903
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 4.4703483581542969e-08,
+  "n" : 512,
+  "ok" : true,
+  "r" : 91,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 23946,
+  "swift_cpu_ref_ms" : 0.26202201843261719,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n8kr1050.json b/experiments/metal-kernels/results/msl_n8kr1050.json
new file mode 100644
index 0000000..93922d4
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n8kr1050.json
@@ -0,0 +1,46 @@
+{
+  "counts_ok" : true,
+  "counts_sum" : 8000,
+  "device" : "Apple M4",
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 2.5299787521362305,
+  "e2e_median_ms" : 2.5069713592529297,
+  "e2e_min_ms" : 2.2299289703369141,
+  "e2e_ms" : [
+    2.5069713592529297,
+    2.406001091003418,
+    2.2299289703369141,
+    2.5299787521362305,
+    2.5290250778198242,
+    2.5060176849365234,
+    2.5110244750976562
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 1.8955416162498295,
+  "kernel_median_ms" : 1.890124985948205,
+  "kernel_min_ms" : 1.887249993160367,
+  "kernel_ms" : [
+    1.8955416162498295,
+    1.8928333302028477,
+    1.887249993160367,
+    1.8891249783337116,
+    1.890124985948205,
+    1.890874991659075,
+    1.8898749840445817
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 4.8428773880004883e-08,
+  "n" : 8000,
+  "ok" : true,
+  "r" : 1050,
+  "reps" : 7,
+  "seed" : 1,
+  "sum_idx" : 4149068,
+  "swift_cpu_ref_ms" : 43.668031692504883,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 1
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/msl_n8r5.dump.txt b/experiments/metal-kernels/results/msl_n8r5.dump.txt
new file mode 100644
index 0000000..02ff08b
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n8r5.dump.txt
@@ -0,0 +1,11 @@
+ASSIGNMENTS
+4 0.13612471520900726
+2 0.14139918982982635
+3 0.4722365438938141
+3 0.05906146392226219
+4 0.47954317927360535
+3 0.5190921425819397
+3 0.308963805437088
+1 0.05956927686929703
+COUNTS
+0 1 1 4 2
diff --git a/experiments/metal-kernels/results/msl_n8r5.json b/experiments/metal-kernels/results/msl_n8r5.json
new file mode 100644
index 0000000..5d66472
--- /dev/null
+++ b/experiments/metal-kernels/results/msl_n8r5.json
@@ -0,0 +1,61 @@
+{
+  "counts" : [
+    0,
+    1,
+    1,
+    4,
+    2
+  ],
+  "counts_ok" : true,
+  "counts_sum" : 8,
+  "device" : "Apple M4",
+  "dist" : [
+    0.13612471520900726,
+    0.14139918982982635,
+    0.47223654389381409,
+    0.059061463922262192,
+    0.47954317927360535,
+    0.5190921425819397,
+    0.30896380543708801,
+    0.059569276869297028
+  ],
+  "dist_mismatches" : 0,
+  "e2e_max_ms" : 3.3190250396728516,
+  "e2e_median_ms" : 3.3190250396728516,
+  "e2e_min_ms" : 3.3190250396728516,
+  "e2e_ms" : [
+    3.3190250396728516
+  ],
+  "idx" : [
+    4,
+    2,
+    3,
+    3,
+    4,
+    3,
+    3,
+    1
+  ],
+  "idx_mismatches" : 0,
+  "impl" : "hand_msl",
+  "kernel_max_ms" : 0.010874995496124029,
+  "kernel_median_ms" : 0.010874995496124029,
+  "kernel_min_ms" : 0.010874995496124029,
+  "kernel_ms" : [
+    0.010874995496124029
+  ],
+  "m" : 3,
+  "max_abs_dist_delta" : 2.9802322387695312e-08,
+  "n" : 8,
+  "ok" : true,
+  "r" : 5,
+  "reps" : 1,
+  "seed" : 1,
+  "sum_idx" : 23,
+  "swift_cpu_ref_ms" : 0.0010728836059570312,
+  "tolerance" : {
+    "abs" : 9.9999999747524271e-07,
+    "rel" : 9.9999997473787516e-06
+  },
+  "warmup" : 0
+}
\ No newline at end of file
diff --git a/experiments/metal-kernels/results/profile_box.txt b/experiments/metal-kernels/results/profile_box.txt
new file mode 100644
index 0000000..18ee7cb
--- /dev/null
+++ b/experiments/metal-kernels/results/profile_box.txt
@@ -0,0 +1,15 @@
+Host: shared Linux box, 8 cores, clang 22.1.8 (zig shim), bend 2.0.40
+Date: 2026-11-10 00:39-00:40 America/New_York
+Load average during runs: about 6.3-6.7 (busy; nds varied, associate/niche were stable)
+Protocol: ab/profile_bend_run.py, warm-native, --threads 1, seed=1
+IO.now() millisecond buckets. Not an IGD claim.
+
+=== DTLZ2 M=3 k=10 partitions=12 pop=92 gens=150 (oracle) ===
+Warm 1 (first invocation after compile_s=11.671): run_s=2.544 profile_sum_ms=2446 profile_wall_ms=2541
+  nds 566 ms (23%), niche 680 ms (27%), associate 448 ms (18%), normalize 379 ms (15%), offspring 290 ms (11%)
+Warm 2 (cache hit): run_s=2.833 profile_sum_ms=2709 profile_wall_ms=2830
+  nds 734 ms (27%), niche 680 ms (25%), associate 435 ms (16%), normalize 462 ms (17%), offspring 303 ms (11%)
+
+=== DTLZ2 M=3 partitions=23 pop=512 gens=20 (larger Run) ===
+First invocation after compile_s=9.862: run_s=8.879 profile_sum_ms=8574 profile_wall_ms=8871
+  nds 4314 ms (50%), niche 1524 ms (17%), offspring 1144 ms (13%), associate 695 ms (8%), normalize 574 ms (6%)
diff --git a/experiments/metal-kernels/run_mac_bench.sh b/experiments/metal-kernels/run_mac_bench.sh
new file mode 100755
index 0000000..6bb8813
--- /dev/null
+++ b/experiments/metal-kernels/run_mac_bench.sh
@@ -0,0 +1,153 @@
+#!/usr/bin/env bash
+# Run the experimental associate benches on the Mac mini (Metal).
+# Usage: ./run_mac_bench.sh
+# Expects this directory to sit inside a unsga3-bend checkout (../../src).
+set -euo pipefail
+HERE=$(cd "$(dirname "$0")" && pwd)
+cd "$HERE"
+mkdir -p results
+export BEND_NO_TELEMETRY=1
+export PATH="${HOME}/.bend/bin:/usr/bin:/bin:/usr/sbin:/sbin:/usr/local/bin:$PATH"
+
+{
+  echo "host=$(scutil --get ComputerName 2>/dev/null || hostname)"
+  echo "date=$(date)"
+  echo "uname=$(uname -a)"
+  sysctl -n machdep.cpu.brand_string 2>/dev/null || true
+  sysctl -n hw.memsize 2>/dev/null || true
+  uptime
+  command -v bend || true
+  bend version || true
+  command -v swiftc
+  swiftc --version | head -1
+  command -v clang
+  clang --version | head -1
+} | tee results/mac_env.txt
+
+if ! command -v bend >/dev/null 2>&1; then
+  echo "installing bend via install.sh"
+  curl --proto '=https' --tlsv1.2 -fsSL https://bend-lang.com/install.sh | sh
+  export PATH="${HOME}/.bend/bin:$PATH"
+fi
+# Use current Bend; never pin.
+bend update || true
+bend version | tee -a results/mac_env.txt
+
+echo "==== compile Swift harness ===="
+swiftc -O -o associate_harness AssociateHarness.swift -framework Metal -framework Foundation
+echo "swiftc ok"
+
+run_swift() {
+  local n=$1 r=$2 m=$3 seed=$4 warmup=$5 reps=$6 tag=$7
+  shift 7 || true
+  local json="results/msl_${tag}.json"
+  echo "==== hand MSL n=$n r=$r ===="
+  if [ "$#" -gt 0 ]; then
+    ./associate_harness --n "$n" --r "$r" --m "$m" --seed "$seed" \
+      --warmup "$warmup" --reps "$reps" --metal associate.metal --json "$json" "$@"
+  else
+    ./associate_harness --n "$n" --r "$r" --m "$m" --seed "$seed" \
+      --warmup "$warmup" --reps "$reps" --metal associate.metal --json "$json"
+  fi
+}
+
+compile_bend() {
+  local mode=$1 n=$2 r=$3 m=$4 seed=$5 warmup=$6 reps=$7 dump=$8 out=$9
+  local bangflag=()
+  if [ "$mode" = "bang" ]; then
+    bangflag=(--bang)
+  fi
+  python3 write_wrapper.py --out "$HERE/wrapper_tmp.bend" "${bangflag[@]}" \
+    --n "$n" --r "$r" --m "$m" --seed "$seed" --warmup "$warmup" --reps "$reps" --dump "$dump"
+  echo "==== bend -o $out (mode=$mode) n=$n r=$r ===="
+  bend "$HERE/wrapper_tmp.bend" -o "$out"
+}
+
+run_product() {
+  local bin=$1 log=$2
+  echo "==== run product $(basename "$bin") ===="
+  if "$bin" +RTS --threads 1 > "$log" 2> "${log%.txt}.err"; then
+    echo "product +RTS ok"
+  elif "$bin" --threads 1 > "$log" 2> "${log%.txt}.err"; then
+    echo "product legacy --threads ok"
+  else
+    echo "product run failed" >&2
+    cat "${log%.txt}.err" >&2 || true
+    return 1
+  fi
+  cat "$log"
+}
+
+run_bang_gpu() {
+  local bin=$1 log=$2
+  echo "==== run bang GPU $(basename "$bin") ===="
+  if "$bin" +RTS --threads 1 > "$log" 2> "${log%.txt}.err"; then
+    echo "bang GPU +RTS ok"
+  elif "$bin" --threads 1 > "$log" 2> "${log%.txt}.err"; then
+    echo "bang GPU legacy ok"
+  else
+    echo "bang GPU run failed" >&2
+    cat "${log%.txt}.err" >&2 || true
+    return 1
+  fi
+  cat "$log"
+}
+
+run_bang_cpu() {
+  local bin=$1 log=$2
+  echo "==== run bang --gpu off $(basename "$bin") ===="
+  if "$bin" +RTS --threads 1 --gpu off > "$log" 2> "${log%.txt}.err"; then
+    echo "bang --gpu off +RTS ok"
+  elif "$bin" --threads 1 --gpu off > "$log" 2> "${log%.txt}.err"; then
+    echo "bang --gpu off legacy ok"
+  else
+    echo "bang --gpu off failed (recorded, not fatal)" >&2
+    cat "${log%.txt}.err" >&2 || true
+    return 0
+  fi
+  cat "$log"
+}
+
+# Correctness fixture (matches ref_associate.py defaults).
+run_swift 8 5 3 1 0 1 n8r5 --dump
+python3 - << 'PY'
+import json
+from pathlib import Path
+d = json.loads(Path("results/msl_n8r5.json").read_text())
+dump = Path("results/msl_n8r5.dump.txt")
+lines = ["ASSIGNMENTS"]
+for i, x in zip(d["idx"], d["dist"]):
+    lines.append(f"{i} {x}")
+lines.append("COUNTS")
+lines.append(" ".join(str(c) for c in d["counts"]))
+dump.write_text("\n".join(lines) + "\n")
+print(dump)
+PY
+python3 ref_associate.py --n 8 --r 5 --m 3 --seed 1 --compare results/msl_n8r5.dump.txt \
+  > results/fixture_python_vs_msl.stdout.txt
+
+# Realistic, mid, big-N.
+#   184 x 91     : DTLZ2 combined parent+offspring pool (pop=92, p=12 -> 91 refs)
+#   2048 x 276   : mid (p=23 -> 276 refs)
+#   32000 x 3003 : Phase D go-big associate-shaped cell
+SIZES="184 91 3 1 2 7 n184r91
+2048 276 3 1 1 7 n2048r276
+32000 3003 3 1 1 5 n32kr3003"
+
+while read -r n r m seed warmup reps tag; do
+  [ -n "${n:-}" ] || continue
+  run_swift "$n" "$r" "$m" "$seed" "$warmup" "$reps" "$tag"
+
+  compile_bend product "$n" "$r" "$m" "$seed" "$warmup" "$reps" 0 "$HERE/bin_product_$tag"
+  run_product "$HERE/bin_product_$tag" "results/bend_product_${tag}.txt"
+
+  compile_bend bang "$n" "$r" "$m" "$seed" "$warmup" "$reps" 0 "$HERE/bin_bang_$tag"
+  run_bang_gpu "$HERE/bin_bang_$tag" "results/bend_bang_${tag}.txt"
+  run_bang_cpu "$HERE/bin_bang_$tag" "results/bend_bangcpu_${tag}.txt"
+done << EOF
+$SIZES
+EOF
+
+python3 combine_results.py
+echo "==== all benches done ===="
+ls -l results
diff --git a/experiments/metal-kernels/write_wrapper.py b/experiments/metal-kernels/write_wrapper.py
new file mode 100755
index 0000000..0b0bb33
--- /dev/null
+++ b/experiments/metal-kernels/write_wrapper.py
@@ -0,0 +1,48 @@
+#!/usr/bin/env python3
+"""Write a tiny Bend main that calls assoc_map.run with baked Nat literals."""
+from __future__ import annotations
+
+import argparse
+from pathlib import Path
+
+
+def main() -> int:
+    p = argparse.ArgumentParser()
+    p.add_argument("--out", type=Path, required=True)
+    p.add_argument("--bang", action="store_true")
+    p.add_argument("--n", type=int, required=True)
+    p.add_argument("--r", type=int, required=True)
+    p.add_argument("--m", type=int, default=3)
+    p.add_argument("--seed", type=int, default=1)
+    p.add_argument("--warmup", type=int, default=1)
+    p.add_argument("--reps", type=int, default=5)
+    p.add_argument("--dump", type=int, default=0)
+    args = p.parse_args()
+    bang = "True{}" if args.bang else "False{}"
+    kind = "bang" if args.bang else "product"
+    args.out.write_text(
+        "# Generated wrapper ({kind}). Isolated experiment, not Run.\n"
+        "import Base\n"
+        "import ./assoc_map.bend as M\n"
+        "\n"
+        "def main() -> IO(Unit):\n"
+        "  M.run({n}n, {r}n, {m}n, {seed}n, {warmup}n, {reps}n, {dump}n, {bang})\n"
+        .format(
+            kind=kind,
+            n=args.n,
+            r=args.r,
+            m=args.m,
+            seed=args.seed,
+            warmup=args.warmup,
+            reps=args.reps,
+            dump=args.dump,
+            bang=bang,
+        ),
+        encoding="utf-8",
+    )
+    print(args.out)
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())