Web demo: exact kernel gates + continuous cross-device kernel verification

web/public/app.js

@@ -565,19 +565,40 @@ function onFragment(peerId, buf) {               // returns full message when co
   return out.buffer;
 }
 
-// ---- gradient wire format: [int32 step][uint32 whash][f32 loss][f32 grad...]
-// whash = FNV-1a of the sender's weights BEFORE this step's update, so every
-// peer can verify the whole group is still training the same model.
-// loss = the sender's local batch loss, so everyone can show the true
-// cluster-average loss and log what each device contributes.
-function packGrad(step, whash, loss, grad) {
-  const buf = new ArrayBuffer(12 + grad.byteLength);
+// ---- gradient wire: [i32 step][u32 whash][u32 phash][f32 loss][f32 grad...]
+// whash = FNV-1a of the sender's weights BEFORE this step's update — proves the
+//   whole group is still applying the same updates to the same model.
+// phash = FNV-1a of the sender's canonical kernel probe — proves the sender's
+//   arithmetic still agrees with the fleet's. whash CANNOT do this: a device
+//   with a broken kernel broadcasts a bad gradient, everyone averages the same
+//   bad bytes, and every replica stays bit-identical. Replica agreement is not
+//   kernel correctness, so the two hashes check genuinely different things.
+// loss = the sender's local batch loss, for the true cluster-average display.
+function packGrad(step, whash, phash, loss, grad) {
+  const buf = new ArrayBuffer(16 + grad.byteLength);
   new Int32Array(buf, 0, 1)[0] = step;
-  new Uint32Array(buf, 4, 1)[0] = whash;
-  new Float32Array(buf, 8, 1)[0] = loss;
-  new Float32Array(buf, 12).set(grad);
+  new Uint32Array(buf, 4, 2).set([whash, phash]);
+  new Float32Array(buf, 12, 1)[0] = loss;
+  new Float32Array(buf, 16).set(grad);
   return buf;
 }
+// ---- continuous kernel verification ----------------------------------------
+// The boot gate only ever sees toy shapes at one moment. These two keep proving
+// it while the run is live: a random-cell audit against the units at the REAL
+// shapes, and the canonical probe re-run on a cadence so thermal/driver drift
+// shows up mid-run rather than never.
+const PROBE_EVERY = 25;                        // steps between probe refreshes
+let probeHash = 0, auditFailure = null;
+const gemmAudit = {
+  cells: 6,
+  due: () => Math.random() < 0.02,            // ~2% of GEMMs, a few cells each
+  fail: (msg) => { if (!auditFailure) { auditFailure = msg; log(`KERNEL AUDIT FAILED: ${msg}`); } },
+};
+async function refreshProbe() {
+  try { probeHash = await Compute.kernelProbe(compute, L); }
+  catch (e) { log(`kernel probe failed to run: ${e.message}`); }
+  return probeHash;
+}
 function hashWeights() {                        // FNV-1a over the flat param bytes
   const f = Transformer.getFlatParams(model);
   const b = new Uint8Array(f.buffer, f.byteOffset, f.byteLength);
@@ -618,16 +639,16 @@ function onGrad(peerId, buf) {
     rosters.set(iv[1], [...iv.slice(3, 3 + iv[2])].map(n => "p" + n));
     wake(); return;
   }
-  const whash = new Uint32Array(buf, 4, 1)[0];
-  const loss = new Float32Array(buf, 8, 1)[0];
-  const grad = new Float32Array(buf.slice(12));
+  const [whash, phash] = new Uint32Array(buf, 4, 2);
+  const loss = new Float32Array(buf, 12, 1)[0];
+  const grad = new Float32Array(buf.slice(16));
   if (!incoming.has(step)) incoming.set(step, new Map());
   incoming.get(step).set(peerId, grad);
   if (!peerHashes.has(step)) peerHashes.set(step, new Map());
-  peerHashes.get(step).set(peerId, { hash: whash, loss });
+  peerHashes.get(step).set(peerId, { hash: whash, phash, loss });
   wake();                                        // resolve waits immediately (no polling)
 }
-function broadcastGrad(step, whash, loss, grad) { return broadcast(packGrad(step, whash, loss, grad)); }
+function broadcastGrad(step, whash, phash, loss, grad) { return broadcast(packGrad(step, whash, phash, loss, grad)); }
 // Event-driven: re-checked on every gradient arrival and peer departure, plus a
 // coarse fallback timer (background tabs throttle timers to ~1s, so the old
 // 15ms poll was the bottleneck there). Peers that left are dropped from the
@@ -680,6 +701,8 @@ async function train(cfg, resume) {
   const opt = TrainCore.makeAdam(model.nParams, { lr: cfg.lr });
   if (resume) opt.setState(resume.adam);               // bit-identical moments from the leader
   pendingJoins.clear();
+  auditFailure = null;
+  await refreshProbe();                                // fresh kernel proof for this run
   log(`training started — cohort ${cohort.length} peer(s), world ${cohort.length + 1}, ` +
       `width=${cfg.c} seq=${cfg.t} batch=${cfg.b}×${cohort.length + 1}, optimizer ${opt.name}` +
       (cohort.length ? `, sync ${iLead ? "led by me" : "led by " + nmeOf(leaderId)}` : "") +
@@ -712,9 +735,14 @@ async function train(cfg, resume) {
       if (lastLocal && lastLocal.s === s) {
         ({ loss, grad } = lastLocal);
       } else {
+        // re-prove this device's kernel against the units on a cadence, so
+        // drift that appears mid-run (thermal, driver, shape) is caught while
+        // it runs rather than only at boot on toy inputs
+        if (s % PROBE_EVERY === 0) await refreshProbe();
         ({ loss, grad } = await localStep());
+        if (auditFailure) { halted = auditFailure; break; }
         lastLocal = { s, loss, grad };
-        await broadcastGrad(s, whash, loss, grad);
+        await broadcastGrad(s, whash, probeHash, loss, grad);
       }
     }
     let all;
@@ -782,6 +810,16 @@ async function train(cfg, resume) {
     if (hs) {
       const bad = roster.find(id => id !== myId && hs.has(id) && hs.get(id).hash !== whash);
       if (bad) { halted = `weights diverged from ${nmeOf(bad)} (detected at step ${s + 1})`; break; }
+      // kernel disagreement: same seeded probe, same exact integer arithmetic
+      // expected on every backend. This is the check the weight hash can't make.
+      const badK = roster.find(id => id !== myId && hs.has(id) && hs.get(id).phash &&
+                                     probeHash && hs.get(id).phash !== probeHash);
+      if (badK) {
+        halted = `kernel disagreement with ${nmeOf(badK)} at step ${s + 1} — their verified-unit probe ` +
+                 `hashes ${hs.get(badK).phash} vs mine ${probeHash}. One of us is computing different ` +
+                 `arithmetic; the shared weights would look fine either way, so stopping.`;
+        break;
+      }
     }
     // cluster-average loss over this step's actual contributors
     let lossSum = 0, lossN = 0;
@@ -842,7 +880,7 @@ function cfgSliders() { return [ui.cfgC, ui.cfgT, ui.cfgB, ui.cfgSteps, ui.cfgLr
 // (re)build the mini transformer — deterministic shared init (same seeds on
 // every peer); each device samples its own batch windows from the shared corpus
 function buildModel(cfg) {
-  model = Transformer.init(cfg, L, compute);       // compute.bgemm = fused batched GPU path
+  model = Transformer.init(cfg, L, compute, gemmAudit);  // compute.bgemm = fused batched GPU path
   trainedSteps = 0;
 }
 
@@ -894,6 +932,11 @@ function buildModel(cfg) {
     return;                                              // no signaling, no training
   }
   ui.backend.textContent = `${compute.backend.toUpperCase()} — ${compute.label} · block-scaled INT8, batched + fused epilogue, through verified units`;
+  // canonical kernel probe: a fixed seeded GEMM through this device's live
+  // kernel. Every device that agrees with the units produces the same number,
+  // GPU or CPU — so it is comparable across the fleet, and re-run mid-training.
+  await refreshProbe();
+  log(`kernel probe: ${probeHash} (exact int32 through the units — every honest device gets this same number)`);
   // tokenizer must load BEFORE the model is built (it sets the vocab size)
   try {
     const name = await Transformer.loadTokenizer();

web/public/transformer.js

@@ -117,7 +117,8 @@
   // the kernel epilogue (ex. 12). Replaces the per-tensor 3-pass: same outlier
   // robustness at one third of the unit ops.
   async function vmm(Xf, Wf, m, k, n, ctx, relu) {
-    return V.vgemmBlock(Xf, Wf, { m, k, n, batch: 1, relu: !!relu }, ctx.L, ctx.bgemm);
+    // ctx.audit re-checks random cells of this LIVE GEMM against the units
+    return V.vgemmBlock(Xf, Wf, { m, k, n, batch: 1, relu: !!relu }, ctx.L, ctx.bgemm, ctx.audit);
   }
 
   // ---- layernorm (no affine) -------------------------------------------------
@@ -146,7 +147,7 @@
   // cfg: { c: width, t: seq len, b: batch/device, layers, heads, steps, lr }
   // engine: the Compute backend object ({bgemm} for the fused WebGPU path) or a
   // legacy matmulInt8 function (Node tests, inference kit) -> CPU LUT mirror
-  function init(cfg, L, engine) {
+  function init(cfg, L, engine, audit) {
     const c = cfg.c, layers = cfg.layers || 2, heads = cfg.heads || 2, hidden = 2 * c;
     let seed = 100;
     const mk = (nEl, scale) => { const w = randn(nEl, mulberry32(seed++)); for (let i = 0; i < nEl; i++) w[i] *= scale; return w; };
@@ -155,7 +156,8 @@
     const m = {
       cfg: { ...cfg, layers, heads, hidden, vocab: vocabSize() },
       ctx: { L, bgemm: (engine && engine.bgemm) || null,
-             att: (engine && engine.att) || null, fgemm: (engine && engine.fgemm) || null },
+             att: (engine && engine.att) || null, fgemm: (engine && engine.fgemm) || null,
+             audit: audit || null },
       emb: add("emb", mk(vocabSize() * c, 0.08)),
       pos: add("pos", mk(cfg.t * c, 0.02)),
       blocks: [], params, names,

web/public/verified_core.js

@@ -100,12 +100,23 @@
     return { q, s };
   }
 
+  // ---- epilogue mirror -------------------------------------------------------
+  // BIT-EXACT mirror of the WGSL epilogue `f32(s) * a * b`. WGSL rounds to f32
+  // after the int->float conversion and after EACH multiply; plain JS would do
+  // the whole chain in f64 and round once, which differs in the last ulp. That
+  // last-ulp gap is what used to force a tolerance into the kernel gates —
+  // mirroring the rounding exactly is what lets the gates compare with `!==`.
+  const f32 = Math.fround;
+  function epi(s, a, b) { return f32(f32(f32(s) * a) * b); }
+
   // CPU mirror of the fused GPU kernel: batched int8 GEMM through the LUT with
   // the epilogue (block dequant + optional ReLU) applied before returning —
-  // exactly what the WGSL kernel does on-device.
+  // exactly what the WGSL kernel does on-device. d.acc=true returns the raw
+  // int32 accumulator instead (the exact oracle the fused kernel normally hides).
   function bgemmJS(Xq, Wq, rs, cs, d, L) {
     const { m, k, n } = d, batch = d.batch || 1, relu = !!d.relu, mul = L.mul;
-    const out = new Float32Array(batch * m * n);
+    const raw = !!d.acc;
+    const out = raw ? new Int32Array(batch * m * n) : new Float32Array(batch * m * n);
     const acc = new Int32Array(n);
     for (let bz = 0; bz < batch; bz++) {
       const xo = bz * m * k, wo = bz * k * n, oo = bz * m * n, co = bz * n;
@@ -116,9 +127,11 @@
           const au = (Xq[xrow + p] & 0xFF) * 256, wrow = wo + p * n;
           for (let j = 0; j < n; j++) acc[j] += mul[au + (Wq[wrow + j] & 0xFF)];
         }
-        const rscale = rs[bz * m + i], orow = oo + i * n;
+        const orow = oo + i * n;
+        if (raw) { for (let j = 0; j < n; j++) out[orow + j] = acc[j]; continue; }
+        const rscale = rs[bz * m + i];
         for (let j = 0; j < n; j++) {
-          const v = acc[j] * rscale * cs[co + j];
+          const v = epi(acc[j], rscale, cs[co + j]);
           out[orow + j] = relu && v < 0 ? 0 : v;
         }
       }
@@ -126,11 +139,32 @@
     return out;
   }
 
+  // Recompute a handful of RANDOM output cells of a live GEMM through the LUT
+  // mirror and compare against what the kernel produced. Sampling cells instead
+  // of whole matrices makes this cheap enough to run continuously, at the real
+  // shapes training uses — not once at boot on toy inputs.
+  function auditTile(Xq, Wq, rs, cs, d, got, L, nCells) {
+    const { m, k, n } = d, batch = d.batch || 1, relu = !!d.relu, mul = L.mul;
+    for (let t = 0; t < (nCells || 8); t++) {
+      const bz = (Math.random() * batch) | 0;
+      const i = (Math.random() * m) | 0, j = (Math.random() * n) | 0;
+      let acc = 0;
+      const xrow = bz * m * k + i * k, wo = bz * k * n;
+      for (let p = 0; p < k; p++) acc += mul[(Xq[xrow + p] & 0xFF) * 256 + (Wq[wo + p * n + j] & 0xFF)];
+      let v = epi(acc, rs[bz * m + i], cs[bz * n + j]);
+      if (relu && v < 0) v = 0;
+      const idx = (bz * m + i) * n + j;
+      if (got[idx] !== v)
+        return `GEMM audit failed at [b${bz},${i},${j}] shape ${m}x${k}x${n}: kernel ${got[idx]} vs units ${v}`;
+    }
+    return null;
+  }
+
   // block-scaled verified GEMM, float in → float out.
   //   d = { m, k, n, batch=1, relu=false }; X is (batch·m)×k, W is batch×(k×n)
   // gpuBgemm (from webgpu.js) runs the batched kernel with the fused epilogue;
   // without it the CPU LUT mirror runs. Every product goes through the units.
-  async function vgemmBlock(Xf, Wf, d, L, gpuBgemm) {
+  async function vgemmBlock(Xf, Wf, d, L, gpuBgemm, audit) {
     const { m, k, n } = d, batch = d.batch || 1;
     const x = quantizeRows(Xf, batch * m, k);
     let wq, ws;
@@ -143,7 +177,16 @@
         wq.set(w.q, bz * k * n); ws.set(w.s, bz * n);
       }
     }
-    if (gpuBgemm) return gpuBgemm(x.q, wq, x.s, ws, d);
+    if (gpuBgemm) {
+      const out = await gpuBgemm(x.q, wq, x.s, ws, d);
+      // continuous re-verification at LIVE shapes: the boot gate only ever saw
+      // toy inputs, so sample a few real cells against the units as we go
+      if (audit && audit.due()) {
+        const bad = auditTile(x.q, wq, x.s, ws, d, out, L, audit.cells);
+        if (bad) audit.fail(bad);
+      }
+      return out;
+    }
     return bgemmJS(x.q, wq, x.s, ws, d, L);
   }
 
@@ -176,9 +219,10 @@
     return { q, s };
   }
   // scores S[bz,ti,tj] = q_row(bi,ti,h) · k_row(bi,tj,h), every product via the LUT
+  // d.acc=true returns the raw int32 accumulator (exact oracle for the kernel gate)
   function attScoresJS(qq, kq, qs, ks, d, L) {
-    const { B, T, heads, hd } = d, C = heads * hd, mul = L.mul;
-    const out = new Float32Array(B * heads * T * T);
+    const { B, T, heads, hd } = d, C = heads * hd, mul = L.mul, raw = !!d.acc;
+    const out = raw ? new Int32Array(B * heads * T * T) : new Float32Array(B * heads * T * T);
     for (let bi = 0; bi < B; bi++) for (let h = 0; h < heads; h++) {
       const bz = bi * heads + h;
       for (let ti = 0; ti < T; ti++) {
@@ -187,7 +231,7 @@
           const ko = (bi * T + tj) * C + h * hd;
           let acc = 0;
           for (let p = 0; p < hd; p++) acc += mul[(qq[qo + p] & 0xFF) * 256 + (kq[ko + p] & 0xFF)];
-          out[(bz * T + ti) * T + tj] = acc * rscale * ks[(bi * T + tj) * heads + h];
+          out[(bz * T + ti) * T + tj] = raw ? acc : epi(acc, rscale, ks[(bi * T + tj) * heads + h]);
         }
       }
     }
@@ -195,8 +239,8 @@
   }
   // ctx[(bi,ti),(h,j)] = Σ_tj a[bz,ti,tj]·v[(bi,tj),(h,j)] — scatter fused into BT×C
   function attCtxJS(aq, vq, as, vs, d, L) {
-    const { B, T, heads, hd } = d, C = heads * hd, mul = L.mul;
-    const out = new Float32Array(B * T * C);
+    const { B, T, heads, hd } = d, C = heads * hd, mul = L.mul, raw = !!d.acc;
+    const out = raw ? new Int32Array(B * T * C) : new Float32Array(B * T * C);
     for (let bi = 0; bi < B; bi++) for (let h = 0; h < heads; h++) {
       const bz = bi * heads + h;
       for (let ti = 0; ti < T; ti++) {
@@ -205,7 +249,7 @@
           let acc = 0;
           for (let tj = 0; tj < T; tj++)
             acc += mul[(aq[ao + tj] & 0xFF) * 256 + (vq[(bi * T + tj) * C + h * hd + j] & 0xFF)];
-          out[(bi * T + ti) * C + h * hd + j] = acc * rscale * vs[(bi * heads + h) * hd + j];
+          out[(bi * T + ti) * C + h * hd + j] = raw ? acc : epi(acc, rscale, vs[(bi * heads + h) * hd + j]);
         }
       }
     }
@@ -263,7 +307,7 @@
   }
 
   const api = { quantize, quantize2, quantizeRows, quantizeCols, quantizeHeadCols, lutMatmulJS, lutMatmul3JS, lutMatmul3,
-                bgemmJS, vgemmBlock, attScoresJS, attCtxJS, linearFwd, forward, backward, splitApply };
+                bgemmJS, vgemmBlock, auditTile, epi, attScoresJS, attCtxJS, linearFwd, forward, backward, splitApply };
   if (typeof module !== "undefined" && module.exports) { TC = require("./traincore.js"); module.exports = api; }
   else { TC = root.TrainCore; root.Verified = api; }
 })(typeof self !== "undefined" ? self : this);

web/public/webgpu.js

@@ -35,36 +35,32 @@
       C[row * n + col] = s;
     }`;
 
-  // X packed 4×int8 along k; W transposed then packed the same way, so each
-  // thread streams two contiguous rows of u32 words through dot4I8Packed.
-  const WGSL_DP4 = `
-    @group(0) @binding(0) var<storage, read> Xp : array<u32>;
-    @group(0) @binding(1) var<storage, read> Wp : array<u32>;    // Wᵀ, packed
-    @group(0) @binding(2) var<storage, read_write> C : array<i32>;
-    @group(0) @binding(3) var<uniform> dims : vec3<u32>;         // m, kw=ceil(k/4), n
-    @compute @workgroup_size(8, 8)
-    fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
-      let m = dims.x; let kw = dims.y; let n = dims.z;
-      let row = gid.x; let col = gid.y;
-      if (row >= m || col >= n) { return; }
-      var s : i32 = 0;
-      for (var p = 0u; p < kw; p = p + 1u) {
-        s = s + dot4I8Packed(Xp[row * kw + p], Wp[col * kw + p]);
-      }
-      C[row * n + col] = s;
-    }`;
+  // NOTE: the un-batched DP4A matmul that used to live here was removed. It was
+  // the only kernel with an exact gate, but the transformer stopped calling it
+  // when the block-scaled path landed — so it sat here passing its own gate
+  // while verifying nothing that ran. A gate on a kernel nobody calls is worse
+  // than no gate: it reads like coverage. The batched kernels below are the
+  // ones training uses, and they now carry that exact gate instead.
 
   // Batched block-scaled GEMM with a FUSED EPILOGUE (CUTLASS ex. 05/24 + 12):
   // grid z = batch index, so ALL attention heads run in ONE dispatch, and the
   // epilogue (block dequant rs·cs + optional ReLU) happens before the data
   // leaves the GPU — f32 out, no int32 readback, no second pass in JS.
-  const WGSL_BG_LUT = `
+  // Each kernel is emitted in two variants from ONE source string: the live one
+  // (fused epilogue, f32 out) and a `verify` one that writes the raw int32
+  // accumulator instead. The indexing — the part that actually goes wrong — is
+  // textually identical, so gating the verify variant genuinely gates the live
+  // kernel, and the comparison is EXACT (int8xint8->int32 has no rounding to
+  // hide in) instead of an allclose that whole bug classes walk through.
+  const OUT_DECL = (v, b) => `@group(0) @binding(${b}) var<storage, read_write> O : array<${v ? "i32" : "f32"}>;`;
+
+  const WGSL_BG_LUT = (verify) => `
     @group(0) @binding(0) var<storage, read> Xq  : array<i32>;   // int8 byte per elem
     @group(0) @binding(1) var<storage, read> Wq  : array<i32>;
     @group(0) @binding(2) var<storage, read> lut : array<i32>;
     @group(0) @binding(3) var<storage, read> rs  : array<f32>;   // per (batch,row)
     @group(0) @binding(4) var<storage, read> cs  : array<f32>;   // per (batch,col)
-    @group(0) @binding(5) var<storage, read_write> O : array<f32>;
+    ${OUT_DECL(verify, 5)}
     @group(0) @binding(6) var<uniform> dims : vec4<u32>;         // m, k, n, flags(1=relu)
     @compute @workgroup_size(8, 8, 1)
     fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
@@ -79,18 +75,19 @@
         let bu = u32(Wq[wo + p * n] & 255);
         s = s + lut[au * 256u + bu];
       }
+      ${verify ? `O[(bz * m + row) * n + col] = s;` : `
       var v = f32(s) * rs[bz * m + row] * cs[bz * n + col];
       if ((dims.w & 1u) == 1u && v < 0.0) { v = 0.0; }
-      O[(bz * m + row) * n + col] = v;
+      O[(bz * m + row) * n + col] = v;`}
     }`;
 
   // same fused/batched kernel on the DP4A hardware path (Wᵀ packed per batch)
-  const WGSL_BG_DP4 = `
+  const WGSL_BG_DP4 = (verify) => `
     @group(0) @binding(0) var<storage, read> Xp : array<u32>;
     @group(0) @binding(1) var<storage, read> Wp : array<u32>;    // per-batch Wᵀ, packed
     @group(0) @binding(2) var<storage, read> rs : array<f32>;
     @group(0) @binding(3) var<storage, read> cs : array<f32>;
-    @group(0) @binding(4) var<storage, read_write> O : array<f32>;
+    ${OUT_DECL(verify, 4)}
     @group(0) @binding(5) var<uniform> dims : vec4<u32>;         // m, kw, n, flags
     @compute @workgroup_size(8, 8, 1)
     fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
@@ -103,21 +100,22 @@
       for (var p = 0u; p < kw; p = p + 1u) {
         s = s + dot4I8Packed(Xp[xo + p], Wp[wo + p]);
       }
+      ${verify ? `O[(bz * m + row) * n + col] = s;` : `
       var v = f32(s) * rs[bz * m + row] * cs[bz * n + col];
       if ((dims.w & 1u) == 1u && v < 0.0) { v = 0.0; }
-      O[(bz * m + row) * n + col] = v;
+      O[(bz * m + row) * n + col] = v;`}
     }`;
 
   // Gather-fused attention (CUTLASS ex. 36/52): kernels index q/k/v directly in
   // their BT×C layout (head-strided) — no gather copies, no kᵀ transpose — and
   // the ctx kernel scatters straight back into BT×C. int8×int8→i32 is exact, so
   // these are bit-identical to the LUT mirrors (proved at init before use).
-  const WGSL_ATT_SCORES = `
+  const WGSL_ATT_SCORES = (verify) => `
     @group(0) @binding(0) var<storage, read> Q : array<i32>;     // int8 per elem, BT×C
     @group(0) @binding(1) var<storage, read> K : array<i32>;
     @group(0) @binding(2) var<storage, read> qs : array<f32>;    // per (token,head)
     @group(0) @binding(3) var<storage, read> ks : array<f32>;
-    @group(0) @binding(4) var<storage, read_write> O : array<f32>;
+    ${OUT_DECL(verify, 4)}
     @group(0) @binding(5) var<uniform> dims : vec4<u32>;         // T, heads, hd, _
     @compute @workgroup_size(8, 8, 1)
     fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
@@ -130,14 +128,15 @@
       let ko = (bi * T + tj) * C + h * hd;
       var s : i32 = 0;
       for (var p = 0u; p < hd; p = p + 1u) { s = s + Q[qo + p] * K[ko + p]; }
-      O[(bz * T + ti) * T + tj] = f32(s) * qs[(bi * T + ti) * heads + h] * ks[(bi * T + tj) * heads + h];
+      ${verify ? `O[(bz * T + ti) * T + tj] = s;`
+               : `O[(bz * T + ti) * T + tj] = f32(s) * qs[(bi * T + ti) * heads + h] * ks[(bi * T + tj) * heads + h];`}
     }`;
-  const WGSL_ATT_CTX = `
+  const WGSL_ATT_CTX = (verify) => `
     @group(0) @binding(0) var<storage, read> A : array<i32>;     // int8, BH×T×T
     @group(0) @binding(1) var<storage, read> V : array<i32>;     // int8, BT×C
     @group(0) @binding(2) var<storage, read> as_ : array<f32>;   // per (bz,row)
     @group(0) @binding(3) var<storage, read> vs : array<f32>;    // per (batch,head,chan)
-    @group(0) @binding(4) var<storage, read_write> O : array<f32>;  // BT×C (scatter fused)
+    ${OUT_DECL(verify, 4)}                                       // BT×C (scatter fused)
     @group(0) @binding(5) var<uniform> dims : vec4<u32>;         // T, heads, hd, _
     @compute @workgroup_size(8, 8, 1)
     fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
@@ -149,7 +148,8 @@
       let ao = (bz * T + ti) * T;
       var s : i32 = 0;
       for (var tj = 0u; tj < T; tj = tj + 1u) { s = s + A[ao + tj] * V[(bi * T + tj) * C + h * hd + j]; }
-      O[(bi * T + ti) * C + h * hd + j] = f32(s) * as_[bz * T + ti] * vs[(bi * heads + h) * hd + j];
+      ${verify ? `O[(bi * T + ti) * C + h * hd + j] = s;`
+               : `O[(bi * T + ti) * C + h * hd + j] = f32(s) * as_[bz * T + ti] * vs[(bi * heads + h) * hd + j];`}
     }`;
 
   // Split-K f32 GEMM (CUTLASS ex. 06) for the STE BACKWARD only — the backward
@@ -237,32 +237,58 @@
       const lut32 = new Int32Array(L.mul);   // widen int16 -> int32
       const lutBuf = device.createBuffer({ size: lut32.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST });
       device.queue.writeBuffer(lutBuf, 0, lut32);
-      const bgLutModule = device.createShaderModule({ code: WGSL_BG_LUT });
-      const bgLutPipe = device.createComputePipeline({ layout: "auto", compute: { module: bgLutModule, entryPoint: "main" } });
-
-      // gather-fused attention kernels — must reproduce the LUT mirrors exactly
       const mkPipe = (code) => device.createComputePipeline({ layout: "auto",
         compute: { module: device.createShaderModule({ code }), entryPoint: "main" } });
-      const scoresPipe = mkPipe(WGSL_ATT_SCORES), ctxPipe = mkPipe(WGSL_ATT_CTX);
-      let att = { scores: (qq, kq, qs, ks, d) => gpuAttScores(device, scoresPipe, qq, kq, qs, ks, d),
-                  ctx: (aq, vq, as, vs, d) => gpuAttCtx(device, ctxPipe, aq, vq, as, vs, d) };
+      // The verify variant doesn't reference the scale buffers, so `layout:auto`
+      // would strip those bindings and the bind group would silently mismatch.
+      // An EXPLICIT layout keeps both variants binding-compatible — which is the
+      // point: they must differ only in the final write, nothing else.
+      const mkLayout = (spec) => device.createBindGroupLayout({
+        entries: spec.map((t, i) => ({ binding: i, visibility: GPUShaderStage.COMPUTE,
+          buffer: { type: t === "u" ? "uniform" : t === "rw" ? "storage" : "read-only-storage" } })) });
+      const mkPipeL = (code, layout) => device.createComputePipeline({
+        layout: device.createPipelineLayout({ bindGroupLayouts: [layout] }),
+        compute: { module: device.createShaderModule({ code }), entryPoint: "main" } });
+      const bgLutLayout = mkLayout(["r", "r", "r", "r", "r", "rw", "u"]);
+      const bgDp4Layout = mkLayout(["r", "r", "r", "r", "rw", "u"]);
+      const attLayout = mkLayout(["r", "r", "r", "r", "rw", "u"]);
+      // live + verify variants, compiled from the same source (see WGSL_* above)
+      const bgLutPipe = mkPipeL(WGSL_BG_LUT(false), bgLutLayout), bgLutVPipe = mkPipeL(WGSL_BG_LUT(true), bgLutLayout);
+      const scoresPipe = mkPipeL(WGSL_ATT_SCORES(false), attLayout), scoresVPipe = mkPipeL(WGSL_ATT_SCORES(true), attLayout);
+      const ctxPipe = mkPipeL(WGSL_ATT_CTX(false), attLayout), ctxVPipe = mkPipeL(WGSL_ATT_CTX(true), attLayout);
+
+      // gather-fused attention kernels. The gate runs the VERIFY variant of the
+      // same source and compares the int32 accumulator with `!==` — exact, no
+      // tolerance — then checks the fused epilogue against a bit-exact JS mirror
+      // of the WGSL rounding. Swept over several shapes, incl. odd/ragged ones,
+      // because head-strided addressing is where these kernels can go wrong.
+      let att = { scores: (qq, kq, qs, ks, d) => gpuAttScores(device, d.acc ? scoresVPipe : scoresPipe, qq, kq, qs, ks, d),
+                  ctx: (aq, vq, as, vs, d) => gpuAttCtx(device, d.acc ? ctxVPipe : ctxPipe, aq, vq, as, vs, d) };
       try {
-        const d0 = { B: 2, T: 8, heads: 2, hd: 8 };
-        const nQ = d0.B * d0.T * d0.heads * d0.hd;
-        const qq = new Int8Array(nQ), kq = new Int8Array(nQ), vq = new Int8Array(nQ);
-        for (let i = 0; i < nQ; i++) { qq[i] = (Math.random() * 256 - 128) | 0; kq[i] = (Math.random() * 256 - 128) | 0; vq[i] = (Math.random() * 256 - 128) | 0; }
-        const qs = Float32Array.from({ length: d0.B * d0.T * d0.heads }, () => Math.random() + 0.5);
-        const ks = Float32Array.from(qs, () => Math.random() + 0.5);
-        const aq = new Int8Array(d0.B * d0.heads * d0.T * d0.T);
-        for (let i = 0; i < aq.length; i++) aq[i] = (Math.random() * 127) | 0;
-        const as = Float32Array.from({ length: d0.B * d0.heads * d0.T }, () => Math.random() + 0.5);
-        const vs = Float32Array.from({ length: d0.B * d0.heads * d0.hd }, () => Math.random() + 0.5);
-        const [hwS, hwC] = await Promise.all([att.scores(qq, kq, qs, ks, d0), att.ctx(aq, vq, as, vs, d0)]);
-        const refS = root.Verified.attScoresJS(qq, kq, qs, ks, d0, L);
-        const refC = root.Verified.attCtxJS(aq, vq, as, vs, d0, L);
-        const close = (a, b) => Math.abs(a - b) <= Math.abs(b) * 1e-6 + 1e-6;
-        for (let i = 0; i < refS.length; i++) if (!close(hwS[i], refS[i])) throw new Error("scores mismatch");
-        for (let i = 0; i < refC.length; i++) if (!close(hwC[i], refC[i])) throw new Error("ctx mismatch");
+        for (const d0 of [{ B: 2, T: 8, heads: 2, hd: 8 }, { B: 1, T: 32, heads: 2, hd: 16 },
+                          { B: 3, T: 7, heads: 3, hd: 5 }, { B: 2, T: 33, heads: 4, hd: 8 }]) {
+          const nQ = d0.B * d0.T * d0.heads * d0.hd;
+          const qq = new Int8Array(nQ), kq = new Int8Array(nQ), vq = new Int8Array(nQ);
+          for (let i = 0; i < nQ; i++) { qq[i] = (Math.random() * 256 - 128) | 0; kq[i] = (Math.random() * 256 - 128) | 0; vq[i] = (Math.random() * 256 - 128) | 0; }
+          const qs = Float32Array.from({ length: d0.B * d0.T * d0.heads }, () => Math.random() + 0.5);
+          const ks = Float32Array.from({ length: d0.B * d0.T * d0.heads }, () => Math.random() + 0.5);
+          const aq = new Int8Array(d0.B * d0.heads * d0.T * d0.T);
+          for (let i = 0; i < aq.length; i++) aq[i] = (Math.random() * 127) | 0;
+          const as = Float32Array.from({ length: d0.B * d0.heads * d0.T }, () => Math.random() + 0.5);
+          const vs = Float32Array.from({ length: d0.B * d0.heads * d0.hd }, () => Math.random() + 0.5);
+          const dv = { ...d0, acc: true };
+          const [accS, accC, hwS, hwC] = await Promise.all([
+            att.scores(qq, kq, qs, ks, dv), att.ctx(aq, vq, as, vs, dv),
+            att.scores(qq, kq, qs, ks, d0), att.ctx(aq, vq, as, vs, d0)]);
+          const refAccS = root.Verified.attScoresJS(qq, kq, qs, ks, dv, L);
+          const refAccC = root.Verified.attCtxJS(aq, vq, as, vs, dv, L);
+          for (let i = 0; i < refAccS.length; i++) if (accS[i] !== refAccS[i]) throw new Error(`scores accumulator mismatch @${i} shape ${JSON.stringify(d0)}`);
+          for (let i = 0; i < refAccC.length; i++) if (accC[i] !== refAccC[i]) throw new Error(`ctx accumulator mismatch @${i} shape ${JSON.stringify(d0)}`);
+          const refS = root.Verified.attScoresJS(qq, kq, qs, ks, d0, L);
+          const refC = root.Verified.attCtxJS(aq, vq, as, vs, d0, L);
+          for (let i = 0; i < refS.length; i++) if (hwS[i] !== refS[i]) throw new Error(`scores epilogue mismatch @${i}`);
+          for (let i = 0; i < refC.length; i++) if (hwC[i] !== refC[i]) throw new Error(`ctx epilogue mismatch @${i}`);
+        }
       } catch (e) {
         console.warn("fused attention kernels failed verification — using CPU LUT mirrors:", e.message);
         att = null;
@@ -284,51 +310,55 @@
         fgemm = null;
       }
 
-      const viaLUT = { backend: "webgpu", label: `${gpuName} (LUT shader)`,
+      // Shared exact gate for a bgemm implementation. Sweeps shapes (including
+      // ragged ones and a k long enough to matter), compares the int32
+      // accumulator from the verify variant with `!==`, then compares the fused
+      // f32 epilogue against the bit-exact JS mirror. Returns null if clean.
+      async function gateBgemm(bgFn) {
+        for (const d0 of [{ m: 5, k: 9, n: 6, batch: 3, relu: true },
+                          { m: 32, k: 64, n: 32, batch: 1, relu: false },
+                          { m: 7, k: 253, n: 5, batch: 2, relu: true },
+                          { m: 1, k: 4, n: 1, batch: 1, relu: false },
+                          { m: 17, k: 33, n: 9, batch: 1, relu: true }]) {
+          const Xq = new Int8Array(d0.batch * d0.m * d0.k), Wq = new Int8Array(d0.batch * d0.k * d0.n);
+          for (let i = 0; i < Xq.length; i++) Xq[i] = (Math.random() * 256 - 128) | 0;
+          for (let i = 0; i < Wq.length; i++) Wq[i] = (Math.random() * 256 - 128) | 0;
+          const rs = Float32Array.from({ length: d0.batch * d0.m }, () => Math.random() + 0.5);
+          const cs = Float32Array.from({ length: d0.batch * d0.n }, () => Math.random() + 0.5);
+          const shape = `${d0.m}x${d0.k}x${d0.n}b${d0.batch}`;
+          const accHw = await bgFn(Xq, Wq, rs, cs, { ...d0, acc: true });
+          const accRef = root.Verified.bgemmJS(Xq, Wq, rs, cs, { ...d0, acc: true }, L);
+          for (let i = 0; i < accRef.length; i++)
+            if (accHw[i] !== accRef[i]) return `accumulator mismatch @${i} (${shape}): ${accHw[i]} vs ${accRef[i]}`;
+          const hw = await bgFn(Xq, Wq, rs, cs, d0);
+          const ref = root.Verified.bgemmJS(Xq, Wq, rs, cs, d0, L);
+          for (let i = 0; i < ref.length; i++)
+            if (hw[i] !== ref[i]) return `epilogue mismatch @${i} (${shape}): ${hw[i]} vs ${ref[i]}`;
+        }
+        return null;
+      }
+
+      const bgLut = (Xq, Wq, rs, cs, d) => gpuBgemmLUT(device, d.acc ? bgLutVPipe : bgLutPipe, lutBuf, Xq, Wq, rs, cs, d);
+      // the LUT bgemm is the fallback AND the oracle's shader twin — gate it too
+      const lutBad = await gateBgemm(bgLut);
+      if (lutBad) { console.warn("LUT bgemm shader failed verification — CPU mirrors only:", lutBad); return cpu; }
+      const viaLUT = { backend: "webgpu", label: `${gpuName} (LUT shader · exact-gated)`,
                        matmulInt8: (Xq, Wq, m, k, n) => gpuMatmulLUT(device, lutPipe, lutBuf, Xq, Wq, m, k, n),
-                       bgemm: (Xq, Wq, rs, cs, d) => gpuBgemmLUT(device, bgLutPipe, lutBuf, Xq, Wq, rs, cs, d),
-                       att, fgemm };
+                       bgemm: bgLut, att, fgemm };
 
-      // DP4A pipeline — only if the WGSL feature exists AND it reproduces the
-      // verified units exactly on random self-tests
+      // DP4A pipeline — only if the WGSL feature exists AND its batched kernel
+      // reproduces the verified units exactly across the shape sweep
       if (!(navigator.gpu.wgslLanguageFeatures && navigator.gpu.wgslLanguageFeatures.has("packed_4x8_integer_dot_product")))
         return viaLUT;
-      const dp4Module = device.createShaderModule({ code: WGSL_DP4 });
-      const dp4Pipe = device.createComputePipeline({ layout: "auto", compute: { module: dp4Module, entryPoint: "main" } });
-      const dp4mm = (Xq, Wq, m, k, n) => gpuMatmulDP4(device, dp4Pipe, Xq, Wq, m, k, n);
-      for (let trial = 0; trial < 3; trial++) {         // hardware vs verified units
-        const m0 = 5 + trial, k0 = 7 + 3 * trial, n0 = 6 + trial;
-        const Xq = new Int8Array(m0 * k0), Wq = new Int8Array(k0 * n0);
-        for (let i = 0; i < Xq.length; i++) Xq[i] = (Math.random() * 256 - 128) | 0;
-        for (let i = 0; i < Wq.length; i++) Wq[i] = (Math.random() * 256 - 128) | 0;
-        const hw = await dp4mm(Xq, Wq, m0, k0, n0);
-        const ref = root.Verified.lutMatmulJS(Xq, Wq, m0, k0, n0, L);
-        for (let i = 0; i < ref.length; i++)
-          if (hw[i] !== ref[i]) { console.warn("DP4A disagreed with the verified units — using LUT shader"); return viaLUT; }
-      }
-      // fused/batched DP4A kernel, gated by the same verification: it must
-      // reproduce the CPU LUT mirror bit-for-bit on the integer accumulate
-      const bgDp4Module = device.createShaderModule({ code: WGSL_BG_DP4 });
-      const bgDp4Pipe = device.createComputePipeline({ layout: "auto", compute: { module: bgDp4Module, entryPoint: "main" } });
-      const bg = (Xq, Wq, rs, cs, d) => gpuBgemmDP4(device, bgDp4Pipe, Xq, Wq, rs, cs, d);
-      {
-        const d0 = { m: 5, k: 9, n: 6, batch: 3, relu: true };
-        const Xq = new Int8Array(d0.batch * d0.m * d0.k), Wq = new Int8Array(d0.batch * d0.k * d0.n);
-        for (let i = 0; i < Xq.length; i++) Xq[i] = (Math.random() * 256 - 128) | 0;
-        for (let i = 0; i < Wq.length; i++) Wq[i] = (Math.random() * 256 - 128) | 0;
-        const rs = new Float32Array(d0.batch * d0.m).fill(0).map(() => Math.random() + 0.5);
-        const cs = new Float32Array(d0.batch * d0.n).fill(0).map(() => Math.random() + 0.5);
-        const hw = await bg(Xq, Wq, rs, cs, d0);
-        const ref = root.Verified.bgemmJS(Xq, Wq, rs, cs, d0, L);
-        for (let i = 0; i < ref.length; i++)
-          if (Math.abs(hw[i] - ref[i]) > Math.abs(ref[i]) * 1e-6 + 1e-6) {
-            console.warn("batched DP4A disagreed with the verified units — using LUT bgemm");
-            return { backend: "webgpu", label: `${gpuName} (DP4A int8 dot — HW verified vs units)`,
-                     matmulInt8: dp4mm, bgemm: viaLUT.bgemm, att, fgemm };
-          }
+      const bgDp4Pipe = mkPipeL(WGSL_BG_DP4(false), bgDp4Layout), bgDp4VPipe = mkPipeL(WGSL_BG_DP4(true), bgDp4Layout);
+      const bg = (Xq, Wq, rs, cs, d) => gpuBgemmDP4(device, d.acc ? bgDp4VPipe : bgDp4Pipe, Xq, Wq, rs, cs, d);
+      const dp4Bad = await gateBgemm(bg);
+      if (dp4Bad) {
+        console.warn("batched DP4A disagreed with the verified units — using LUT bgemm:", dp4Bad);
+        return viaLUT;
       }
-      return { backend: "webgpu", label: `${gpuName} (DP4A int8 dot — HW verified vs units)`,
-               matmulInt8: dp4mm, bgemm: bg, att, fgemm };
+      return { backend: "webgpu", label: `${gpuName} (DP4A int8 dot · exact-gated vs units)`,
+               bgemm: bg, att, fgemm };
     } catch (e) { console.warn("WebGPU init failed, CPU fallback:", e); return cpu; }
   }
 
@@ -364,20 +394,6 @@
     return r.out;
   }
 
-  async function gpuMatmulDP4(device, pipeline, Xq, Wq, m, k, n) {
-    const kw = Math.ceil(k / 4);
-    const Xp = packRows(Xq, m, k, kw);
-    const Wp = packRows(transposeI8(Wq, k, n), n, k, kw);   // Wᵀ so col j is a row
-    const bufX = up(device, Xp, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
-    const bufW = up(device, Wp, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
-    const bufD = up(device, new Uint32Array([m, kw, n, 0]), GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST);
-    const r = await runPass(device, pipeline, (bufC) => [
-      { binding: 0, resource: { buffer: bufX } }, { binding: 1, resource: { buffer: bufW } },
-      { binding: 2, resource: { buffer: bufC } }, { binding: 3, resource: { buffer: bufD } } ], m, n);
-    [bufX, bufW, bufD, r.bufC, r.read].forEach(b => b.destroy());
-    return r.out;
-  }
-
   // attention kernels: int8 (widened i32) in, f32 out, strided head indexing
   async function gpuAttScores(device, pipeline, qq, kq, qs, ks, d) {
     const { B, T, heads } = d;
@@ -389,7 +405,7 @@
     const r = await runBgPass(device, pipeline, (bufO) => [
       { binding: 0, resource: { buffer: bufQ } }, { binding: 1, resource: { buffer: bufK } },
       { binding: 2, resource: { buffer: bufQs } }, { binding: 3, resource: { buffer: bufKs } },
-      { binding: 4, resource: { buffer: bufO } }, { binding: 5, resource: { buffer: bufD } } ], T, T, B * heads);
+      { binding: 4, resource: { buffer: bufO } }, { binding: 5, resource: { buffer: bufD } } ], T, T, B * heads, d.acc);
     [bufQ, bufK, bufQs, bufKs, bufD, r.bufO, r.read].forEach(b => b.destroy());
     return r.out;
   }
@@ -403,7 +419,7 @@
     const r = await runBgPass(device, pipeline, (bufO) => [
       { binding: 0, resource: { buffer: bufA } }, { binding: 1, resource: { buffer: bufV } },
       { binding: 2, resource: { buffer: bufAs } }, { binding: 3, resource: { buffer: bufVs } },
-      { binding: 4, resource: { buffer: bufO } }, { binding: 5, resource: { buffer: bufD } } ], T, hd, B * heads);
+      { binding: 4, resource: { buffer: bufO } }, { binding: 5, resource: { buffer: bufD } } ], T, hd, B * heads, d.acc);
     [bufA, bufV, bufAs, bufVs, bufD, r.bufO, r.read].forEach(b => b.destroy());
     return r.out;
   }
@@ -441,8 +457,9 @@
     return out;
   }
 
-  // fused batched dispatch: f32 out, epilogue done on-device
-  async function runBgPass(device, pipeline, entries, m, n, batch) {
+  // fused batched dispatch: f32 out, epilogue done on-device (raw=true reads the
+  // verify variant's int32 accumulator instead)
+  async function runBgPass(device, pipeline, entries, m, n, batch, raw) {
     const bytesO = batch * m * n * 4;
     const bufO = mk(device, bytesO, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
     const bind = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: entries(bufO) });
@@ -454,7 +471,8 @@
     enc.copyBufferToBuffer(bufO, 0, read, 0, bytesO);
     device.queue.submit([enc.finish()]);
     await read.mapAsync(GPUMapMode.READ);
-    const out = new Float32Array(read.getMappedRange().slice(0));
+    const buf = read.getMappedRange().slice(0);
+    const out = raw ? new Int32Array(buf) : new Float32Array(buf);
     read.unmap();
     return { out, bufO, read };
   }
@@ -470,7 +488,7 @@
       { binding: 0, resource: { buffer: bufX } }, { binding: 1, resource: { buffer: bufW } },
       { binding: 2, resource: { buffer: lutBuf } }, { binding: 3, resource: { buffer: bufR } },
       { binding: 4, resource: { buffer: bufS } }, { binding: 5, resource: { buffer: bufO } },
-      { binding: 6, resource: { buffer: bufD } } ], m, n, batch);
+      { binding: 6, resource: { buffer: bufD } } ], m, n, batch, d.acc);
     [bufX, bufW, bufR, bufS, bufD, r.bufO, r.read].forEach(b => b.destroy());
     return r.out;
   }
@@ -492,7 +510,7 @@
     const r = await runBgPass(device, pipeline, (bufO) => [
       { binding: 0, resource: { buffer: bufX } }, { binding: 1, resource: { buffer: bufW } },
       { binding: 2, resource: { buffer: bufR } }, { binding: 3, resource: { buffer: bufS } },
-      { binding: 4, resource: { buffer: bufO } }, { binding: 5, resource: { buffer: bufD } } ], m, n, batch);
+      { binding: 4, resource: { buffer: bufO } }, { binding: 5, resource: { buffer: bufD } } ], m, n, batch, d.acc);
     [bufX, bufW, bufR, bufS, bufD, r.bufO, r.read].forEach(b => b.destroy());
     return r.out;
   }
@@ -504,5 +522,37 @@
     return b;
   }
 
-  root.Compute = { initCompute, loadLUTs };
+  // ---- canonical kernel probe -------------------------------------------------
+  // The weight hash CANNOT catch a device whose kernel is wrong: weights only
+  // depend on the gradient bytes everyone receives, so a fleet averaging one
+  // device's bad gradient stays bit-identical and perfectly happy. This is the
+  // check that can. Every device runs the SAME seeded int8 GEMM through its own
+  // live kernel (verify variant -> raw int32 accumulator, exact on every
+  // backend — GPU DP4A, GPU LUT, CPU mirror alike) and hashes the result. Same
+  // input + correct kernels => same hash, regardless of hardware. A device that
+  // disagrees is computing different arithmetic than the rest of the fleet.
+  const PROBE = { m: 24, k: 96, n: 24, batch: 2 };
+  function probeInputs() {
+    let seed = 0x5EED;                       // fixed: identical on every device
+    const rnd = () => { seed = (Math.imul(seed, 1103515245) + 12345) & 0x7fffffff; return seed / 0x7fffffff; };
+    const { m, k, n, batch } = PROBE;
+    const Xq = new Int8Array(batch * m * k), Wq = new Int8Array(batch * k * n);
+    for (let i = 0; i < Xq.length; i++) Xq[i] = Math.round(rnd() * 254 - 127);
+    for (let i = 0; i < Wq.length; i++) Wq[i] = Math.round(rnd() * 254 - 127);
+    const rs = Float32Array.from({ length: batch * m }, () => 1);
+    const cs = Float32Array.from({ length: batch * n }, () => 1);
+    return { Xq, Wq, rs, cs, d: { ...PROBE, acc: true } };
+  }
+  async function kernelProbe(compute, L) {
+    const { Xq, Wq, rs, cs, d } = probeInputs();
+    const out = compute && compute.bgemm
+      ? await compute.bgemm(Xq, Wq, rs, cs, d)
+      : root.Verified.bgemmJS(Xq, Wq, rs, cs, d, L);
+    let h = 0x811c9dc5;                      // FNV-1a over the exact int32 results
+    const b = new Uint8Array(out.buffer, out.byteOffset, out.byteLength);
+    for (let i = 0; i < b.length; i++) { h ^= b[i]; h = Math.imul(h, 0x01000193); }
+    return h >>> 0;
+  }
+
+  root.Compute = { initCompute, loadLUTs, kernelProbe };
 })(self);

web/test_gates.js

@@ -0,0 +1,118 @@
+// Mutation test for the kernel admission gate.
+//
+// A gate that has never rejected anything is decoration. This injects deliberately
+// buggy "kernels" and asserts the gate rejects each one — and, separately, shows
+// two things the OLD gate (single tiny shape, allclose) let through.
+//
+// Runs in Node against the CPU mirrors, so the buggy kernels stand in for what a
+// wrong GPU shader would produce.
+const fs = require("fs");
+const path = require("path");
+const V = require("./public/verified_core.js");
+
+const mul = new Int16Array(fs.readFileSync(path.join(__dirname, "public", "mul_lut.bin")).buffer.slice(0));
+const L = { mul };
+
+// ---- the gate under test (mirrors webgpu.js gateBgemm) ----------------------
+const SWEEP = [
+  { m: 5, k: 9, n: 6, batch: 3, relu: true },
+  { m: 32, k: 64, n: 32, batch: 1, relu: false },
+  { m: 7, k: 253, n: 5, batch: 2, relu: true },
+  { m: 1, k: 4, n: 1, batch: 1, relu: false },
+  { m: 17, k: 33, n: 9, batch: 1, relu: true },
+];
+const OLD_SHAPE = [{ m: 5, k: 9, n: 6, batch: 3, relu: true }];   // what we used to test
+const rnd = (n, f) => { const a = new Int8Array(n); for (let i = 0; i < n; i++) a[i] = f(); return a; };
+
+function inputs(d) {
+  const Xq = rnd(d.batch * d.m * d.k, () => (Math.random() * 256 - 128) | 0);
+  const Wq = rnd(d.batch * d.k * d.n, () => (Math.random() * 256 - 128) | 0);
+  const rs = Float32Array.from({ length: d.batch * d.m }, () => Math.random() + 0.5);
+  const cs = Float32Array.from({ length: d.batch * d.n }, () => Math.random() + 0.5);
+  return { Xq, Wq, rs, cs };
+}
+
+// exact gate: int32 accumulator with !==, then f32 epilogue with !==
+function gateExact(kernel, shapes) {
+  for (const d of shapes) {
+    const { Xq, Wq, rs, cs } = inputs(d);
+    const accHw = kernel(Xq, Wq, rs, cs, { ...d, acc: true }, L);
+    const accRef = V.bgemmJS(Xq, Wq, rs, cs, { ...d, acc: true }, L);
+    for (let i = 0; i < accRef.length; i++) if (accHw[i] !== accRef[i]) return `acc @${i}`;
+    const hw = kernel(Xq, Wq, rs, cs, d, L);
+    const ref = V.bgemmJS(Xq, Wq, rs, cs, d, L);
+    for (let i = 0; i < ref.length; i++) if (hw[i] !== ref[i]) return `epilogue @${i}`;
+  }
+  return null;
+}
+// the old gate: one shape, allclose on the fused f32 output only
+function gateOld(kernel) {
+  for (const d of OLD_SHAPE) {
+    const { Xq, Wq, rs, cs } = inputs(d);
+    const hw = kernel(Xq, Wq, rs, cs, d, L);
+    const ref = V.bgemmJS(Xq, Wq, rs, cs, d, L);
+    for (let i = 0; i < ref.length; i++)
+      if (Math.abs(hw[i] - ref[i]) > Math.abs(ref[i]) * 1e-6 + 1e-6) return `allclose @${i}`;
+  }
+  return null;
+}
+
+// ---- buggy kernels ----------------------------------------------------------
+// Each is bgemmJS with exactly one realistic defect injected.
+function mutant(bug) {
+  return function (Xq, Wq, rs, cs, d, LL) {
+    const { m, k, n } = d, batch = d.batch || 1, relu = !!d.relu, mulT = LL.mul;
+    const raw = !!d.acc;
+    const out = raw ? new Int32Array(batch * m * n) : new Float32Array(batch * m * n);
+    const acc = new Int32Array(n);
+    for (let bz = 0; bz < batch; bz++) {
+      const xo = bz * m * k, oo = bz * m * n, co = bz * n;
+      // BUG: batch stride dropped on W — invisible unless batch > 1
+      const wo = bug === "batchStride" ? 0 : bz * k * n;
+      for (let i = 0; i < m; i++) {
+        acc.fill(0);
+        const xrow = xo + i * k;
+        // BUG: last element of the dot product skipped
+        const kEnd = bug === "shortK" ? k - 1 : k;
+        for (let p = 0; p < kEnd; p++) {
+          const au = (Xq[xrow + p] & 0xFF) * 256, wrow = wo + p * n;
+          for (let j = 0; j < n; j++) acc[j] += mulT[au + (Wq[wrow + j] & 0xFF)];
+        }
+        const orow = oo + i * n;
+        if (raw) { for (let j = 0; j < n; j++) out[orow + j] = acc[j]; continue; }
+        const rscale = rs[bz * m + i];
+        for (let j = 0; j < n; j++) {
+          // BUG: column scale ignored / ReLU dropped / f64 rounding (the real one)
+          let v;
+          if (bug === "missingColScale") v = V.epi(acc[j], rscale, 1);
+          else if (bug === "f64Rounding") v = acc[j] * rscale * cs[co + j];   // rounds once, not thrice
+          else v = V.epi(acc[j], rscale, cs[co + j]);
+          out[orow + j] = (relu && bug !== "noRelu" && v < 0) ? 0 : v;
+        }
+      }
+    }
+    return out;
+  };
+}
+
+let pass = true;
+const ok = (cond, msg) => { console.log(`${cond ? "  ok  " : "  FAIL"}  ${msg}`); if (!cond) pass = false; };
+
+console.log("\nthe gate must reject every injected bug:");
+for (const bug of ["batchStride", "shortK", "missingColScale", "noRelu", "f64Rounding"])
+  ok(gateExact(mutant(bug), SWEEP) !== null, `rejects ${bug}  (${gateExact(mutant(bug), SWEEP) || "NOT CAUGHT"})`);
+
+console.log("\nthe gate must accept the correct kernel:");
+ok(gateExact(V.bgemmJS, SWEEP) === null, "accepts the real mirror across the shape sweep");
+
+console.log("\nwhat the OLD gate (one tiny shape, allclose) let through:");
+const oldMissedRounding = gateOld(mutant("f64Rounding")) === null;
+const newCatchesRounding = gateExact(mutant("f64Rounding"), SWEEP) !== null;
+ok(oldMissedRounding && newCatchesRounding,
+   "f64-vs-f32 epilogue rounding: allclose PASSES it, exact gate catches it");
+const oldMissedStride = gateOld(mutant("batchStride")) === null;
+console.log(`  note   batchStride under the old gate: ${oldMissedStride ? "MISSED" : "caught (batch>1 was in the shape)"}`);
+
+console.log(pass ? "\nGATE TEST PASSED — the gate rejects real bugs and accepts the real kernel."
+                 : "\nGATE TEST FAILED");
+process.exit(pass ? 0 : 1);