Web trainer upgrade: sync guard (leader roster + weight-hash divergence check), 3xINT8 fast-accurate GEMM, DP4A hardware path (verified vs units), Spikewhale tokenizer (16.5k vocab), FineWeb-Edu streaming, gradient/checkpoint fragmentation, inference kit, generation tester, contribution logs

web/public/app.js

@@ -29,6 +29,12 @@ const ui = {
   cfgB: document.getElementById("cfgB"),
   cfgSteps: document.getElementById("cfgSteps"),
   cfgLr: document.getElementById("cfgLr"),
+  dataset: document.getElementById("dataset"),
+  tokenizer: document.getElementById("tokenizer"),
+  kit: document.getElementById("kit"),
+  genPrompt: document.getElementById("genPrompt"),
+  genBtn: document.getElementById("genBtn"),
+  genOut: document.getElementById("genOut"),
 };
 function log(m) { ui.log.textContent = `${new Date().toLocaleTimeString()}  ${m}\n` + ui.log.textContent; }
 function setStatus(s) { ui.status.textContent = s; }
@@ -52,10 +58,24 @@ const names = new Map();                       // peerId -> device name
 const incoming = new Map();                    // step -> Map(peerId -> Float32Array)
 let model = null, training = false;            // the mini transformer (transformer.js)
 let trainedSteps = 0;                          // steps baked into the current weights
+// ---- sync guard state -------------------------------------------------------
+// The peer that presses Start is the round leader. Every step the leader
+// publishes the exact set of contributors to average (the roster); followers
+// apply that set verbatim or stop — no device may silently average a
+// different set, which is what used to fork the weights.
+let leaderId = null;                           // null => I lead (I pressed Start)
+const rosters = new Map();                     // step -> [peerId,...] from the leader
+const peerHashes = new Map();                  // step -> Map(peerId -> uint32 pre-step weight hash)
 function nmeOf(id) { return names.get(id) || id; }
 
 function room() { return new URLSearchParams(location.search).get("room"); }   // null -> group by network
-function updatePeers() { ui.peers.textContent = chans.size ? [...chans.keys()].map(nmeOf).join(", ") : "(none yet)"; }
+// show EVERY device the room knows about, not just open data channels —
+// a peer whose WebRTC is still connecting (or failed) must stay visible
+function updatePeers() {
+  if (!names.size) { ui.peers.textContent = "(none yet)"; return; }
+  ui.peers.textContent = [...names.entries()]
+    .map(([id, n]) => `${n} ${chans.has(id) ? "✓" : "(connecting…)"}`).join(", ");
+}
 
 // ---- signaling + WebRTC ----------------------------------------------------
 function connectSignaling() {
@@ -75,6 +95,7 @@ function connectSignaling() {
       else if (room()) setStatus(`accepted into private room "${room()}"`);
       // I'm newest: initiate to everyone already here
       for (const p of msg.peers) { names.set(p.id, p.name); initiatePeer(p.id); }
+      updatePeers();
     } else if (msg.type === "waiting") {
       setStatus("knocking — waiting for the room's host to let you in…");
     } else if (msg.type === "denied") {
@@ -88,6 +109,7 @@ function connectSignaling() {
       addJoinRequest(msg.id, msg.name);
     } else if (msg.type === "peer-joined") {
       names.set(msg.id, msg.name);
+      updatePeers();
       log(`${msg.name} joined (they will connect to me)`);
     } else if (msg.type === "peer-left") {
       log(`${nmeOf(msg.id)} left`); cleanupPeer(msg.id); names.delete(msg.id); updatePeers();
@@ -173,8 +195,10 @@ function parseCheckpoint(buf) {
   const magic = String.fromCharCode(...new Uint8Array(buf, 0, 8));
   if (magic !== CKPT_MAGIC) throw new Error("not a DaisyChain v2 checkpoint");
   const [c, t, vocab, steps] = new Int32Array(buf, 8, 4);
-  if (vocab !== Transformer.VOCAB) throw new Error(`vocab mismatch (file ${vocab}, this build ${Transformer.VOCAB})`);
-  if (c < 16 || c > 64 || t < 16 || t > 64) throw new Error(`bad dims in checkpoint (width ${c}, seq ${t})`);
+  if (vocab !== Transformer.vocabSize())
+    throw new Error(`tokenizer mismatch — checkpoint has a ${vocab}-token vocab, this device loaded ` +
+                    `${Transformer.vocabSize()} (${Transformer.tokenizerName()}); use matching builds`);
+  if (c < 16 || c > 128 || t < 16 || t > 128) throw new Error(`bad dims in checkpoint (width ${c}, seq ${t})`);
   return { c, t, steps, flat: new Float32Array(buf.slice(24)) };
 }
 function applyCheckpoint(ck, from) {
@@ -185,7 +209,7 @@ function applyCheckpoint(ck, from) {
   if (ck.flat.length !== model.nParams) throw new Error("truncated checkpoint");
   Transformer.setFlatParams(model, ck.flat);
   trainedSteps = ck.steps;
-  ui.save.disabled = false;
+  modelReady();
   ui.step.textContent = `${ck.steps} baked in`;
   log(`checkpoint loaded (${ck.steps} steps) ${from ? "from " + from : "from file"} — all set to resume`);
 }
@@ -195,9 +219,65 @@ function broadcastCheckpoint() {
   new Int32Array(msg, 0, 1)[0] = CKPT_SENTINEL;
   new Uint8Array(msg, 4).set(new Uint8Array(ck));
   let n = 0;
-  for (const dc of chans.values()) if (dc.readyState === "open") { dc.send(msg); n++; }
-  log(`checkpoint pushed to ${n} device(s)`);
+  for (const dc of chans.values()) if (dc.readyState === "open") { dcSend(dc, msg); n++; }
+  log(`checkpoint pushed to ${n} device(s) (${(msg.byteLength / 1048576).toFixed(1)} MB)`);
+}
+// ---- inference kit: one self-contained HTML file with the trained weights --
+// Bundles the model code + current checkpoint (base64) + a prompt box into a
+// single file that runs generations offline — no server, nothing to install.
+// The mul8 LUT is rebuilt in the file itself (it is exactly a×b for int8).
+async function downloadInferenceKit() {
+  const srcs = await Promise.all(["traincore.js", "verified_core.js", "transformer.js"]
+    .map(f => fetch(f).then(r => r.text())));
+  let tokData = "null";                          // embed the tokenizer (offline file)
+  try { tokData = await (await fetch("tokenizer.json")).text(); } catch (e) {}
+  const ck = new Uint8Array(packCheckpoint());
+  let b64 = ""; const CH = 0x8000;
+  for (let i = 0; i < ck.length; i += CH) b64 += String.fromCharCode(...ck.subarray(i, i + CH));
+  b64 = btoa(b64);
+  const esc = (s) => s.replace(/<\/script/gi, "<\\/script");
+  const html = `<!doctype html><html><head><meta charset="utf-8">
+<meta name="viewport" content="width=device-width, initial-scale=1">
+<title>DaisyChain inference — ${trainedSteps} steps</title>
+<style>body{font-family:sans-serif;max-width:680px;margin:0 auto;padding:24px 16px;background:#efe4c9;color:#2a1d0a}
+input{width:100%;padding:10px;border-radius:8px;border:1px solid #8b6f47;font-family:monospace;box-sizing:border-box}
+button{margin-top:10px;padding:10px 24px;border:0;border-radius:8px;background:#4a7c2e;color:#f5ecd9;font-weight:700;cursor:pointer}
+pre{background:#fbf6e8;border-radius:8px;padding:12px;white-space:pre-wrap;min-height:60px}</style></head><body>
+<h2>🌼 DaisyChain model — inference</h2>
+<p>Trained ${trainedSteps} steps · width ${model.cfg.c} · seq ${model.cfg.t} · dataset: ${Transformer.datasetName()}.
+Runs entirely in this file through the verified INT8 units (CPU LUT).</p>
+<input id="p" value="the "><button id="g">Generate</button><pre id="o"></pre>
+<script>${esc(srcs[0])}<\/script><script>${esc(srcs[1])}<\/script><script>${esc(srcs[2])}<\/script>
+<script>
+const mul = new Int16Array(65536);
+for (let a = 0; a < 256; a++) for (let b = 0; b < 256; b++)
+  mul[a*256+b] = ((a>127?a-256:a) * (b>127?b-256:b));
+const L = { mul };
+const TOKDATA = ${esc(tokData)};
+if (TOKDATA) Transformer.loadTokenizerData(TOKDATA);
+const bytes = Uint8Array.from(atob("${b64}"), c => c.charCodeAt(0));
+const buf = bytes.buffer;
+const dims = new Int32Array(buf, 8, 4);            // c, t, vocab, steps
+const m = Transformer.init({ c: dims[0], t: dims[1], b: 1, steps: 0, lr: 0 }, L,
+  (Xq, Wq, mm, kk, nn, LL) => Verified.lutMatmulJS(Xq, Wq, mm, kk, nn, LL));
+Transformer.setFlatParams(m, new Float32Array(buf.slice(24)));
+document.getElementById("g").onclick = async () => {
+  const b = document.getElementById("g"), o = document.getElementById("o");
+  b.disabled = true; o.textContent = "generating…";
+  try { o.textContent = await Transformer.generate(m, document.getElementById("p").value || "the ", 150); }
+  catch (e) { o.textContent = "error: " + e.message; }
+  b.disabled = false;
+};
+<\/script></body></html>`;
+  const blob = new Blob([html], { type: "text/html" });
+  const a = document.createElement("a");
+  a.href = URL.createObjectURL(blob);
+  a.download = `daisychain-inference-step${trainedSteps}.html`;
+  a.click();
+  URL.revokeObjectURL(a.href);
+  log("inference kit downloaded — a single HTML file: open it anywhere, type a prompt, generate");
 }
+
 function saveCheckpoint() {
   const blob = new Blob([packCheckpoint()], { type: "application/octet-stream" });
   const a = document.createElement("a");
@@ -210,8 +290,10 @@ function saveCheckpoint() {
 // ---- training config: whoever presses Start sets it for the whole group ----
 const CFG_SENTINEL = -3;                       // wire: [int32 -3][int32 c,t,b,steps][f32 lr]
 function readCfgFromUI() {
+  // lr crosses the wire as float32 — fround here so the leader trains with
+  // the exact same value the followers decode (else Adam forks the weights)
   return { c: +ui.cfgC.value, t: +ui.cfgT.value, b: +ui.cfgB.value,
-           steps: +ui.cfgSteps.value, lr: +ui.cfgLr.value / 1000 };
+           steps: +ui.cfgSteps.value, lr: Math.fround(+ui.cfgLr.value / 1000) };
 }
 function showCfgInUI(cfg) {
   const set = (el, vid, val) => { el.value = val; document.getElementById(vid).textContent = val; };
@@ -220,37 +302,109 @@ function showCfgInUI(cfg) {
   set(ui.cfgLr, "vcfgLr", Math.round(cfg.lr * 1000));
 }
 function broadcastConfig(cfg) {
-  const buf = new ArrayBuffer(24);
+  const buf = new ArrayBuffer(28);
   new Int32Array(buf, 0, 5).set([CFG_SENTINEL, cfg.c, cfg.t, cfg.b, cfg.steps]);
   new Float32Array(buf, 20, 1)[0] = cfg.lr;
-  for (const dc of chans.values()) if (dc.readyState === "open") dc.send(buf);
+  new Int32Array(buf, 24, 1)[0] = Transformer.vocabSize();   // tokenizer must match
+  broadcast(buf);
 }
 function onConfig(peerId, buf) {
   if (training) { log(`ignored settings from ${nmeOf(peerId)} (already training)`); return; }
   const [, c, t, b, steps] = new Int32Array(buf, 0, 5);
   const lr = new Float32Array(buf, 20, 1)[0];
-  if (!(c >= 16 && c <= 64 && c % 2 === 0 && t >= 16 && t <= 64 && b >= 1 && b <= 32 &&
+  const vocab = buf.byteLength >= 28 ? new Int32Array(buf, 24, 1)[0] : -1;
+  if (vocab !== Transformer.vocabSize()) {
+    log(`NOT JOINING: ${nmeOf(peerId)} uses a ${vocab}-token tokenizer, this device has ` +
+        `${Transformer.vocabSize()} (${Transformer.tokenizerName()}) — refresh so all devices match`);
+    return;
+  }
+  if (!(c >= 16 && c <= 128 && c % 2 === 0 && t >= 16 && t <= 128 && b >= 1 && b <= 32 &&
         steps >= 1 && steps <= 10000 && lr > 0 && lr <= 0.2)) {
     log(`rejected bad settings from ${nmeOf(peerId)}`); return;
   }
   const cfg = { c, t, b, steps, lr };
+  leaderId = peerId;                           // the starter leads sync for this run
   showCfgInUI(cfg);
   log(`${nmeOf(peerId)} started the group: width=${c}, seq=${t}, batch=${b}, ${steps} steps, lr=${lr}`);
   buildModel(cfg);
   train(cfg);                                  // follow automatically
 }
 
-// ---- gradient wire format:  [int32 step][float32 grad...] -----------------
-function packGrad(step, grad) {
-  const buf = new ArrayBuffer(4 + grad.byteLength);
+// ---- big-message fragmentation ----------------------------------------------
+// WebRTC data channels cap a single message (~256KB Chrome, 64KB elsewhere).
+// Gradients with the Spikewhale vocab are multi-MB, and checkpoints always
+// were at width ≥ 64 — so anything large goes out in 48KB chunks:
+//   [int32 -5][int32 msgId][int32 seq][int32 total][bytes...]
+// Channels are ordered+reliable, so chunks arrive in order per peer.
+const FRAG_SENTINEL = -5, FRAG_CHUNK = 48 * 1024;
+let fragSeq = 1;
+const fragIn = new Map();                        // peerId -> {id, parts, got, total}
+function dcSend(dc, buf) {
+  if (buf.byteLength <= FRAG_CHUNK) { dc.send(buf); return; }
+  const id = fragSeq++, src = new Uint8Array(buf);
+  const total = Math.ceil(src.length / FRAG_CHUNK);
+  for (let s = 0; s < total; s++) {
+    const part = src.subarray(s * FRAG_CHUNK, Math.min((s + 1) * FRAG_CHUNK, src.length));
+    const msg = new ArrayBuffer(16 + part.length);
+    new Int32Array(msg, 0, 4).set([FRAG_SENTINEL, id, s, total]);
+    new Uint8Array(msg, 16).set(part);
+    dc.send(msg);
+  }
+}
+function broadcast(buf) { for (const dc of chans.values()) if (dc.readyState === "open") dcSend(dc, buf); }
+function onFragment(peerId, buf) {               // returns full message when complete
+  const [, id, seq, total] = new Int32Array(buf, 0, 4);
+  let st = fragIn.get(peerId);
+  if (!st || st.id !== id) { st = { id, parts: [], got: 0, total }; fragIn.set(peerId, st); }
+  st.parts[seq] = new Uint8Array(buf, 16).slice(0);
+  st.got++;
+  if (st.got < st.total) return null;
+  fragIn.delete(peerId);
+  let len = 0; for (const p of st.parts) len += p.length;
+  const out = new Uint8Array(len);
+  let off = 0; for (const p of st.parts) { out.set(p, off); off += p.length; }
+  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);
   new Int32Array(buf, 0, 1)[0] = step;
-  new Float32Array(buf, 4).set(grad);
+  new Uint32Array(buf, 4, 1)[0] = whash;
+  new Float32Array(buf, 8, 1)[0] = loss;
+  new Float32Array(buf, 12).set(grad);
   return buf;
 }
+function hashWeights() {                        // FNV-1a over the flat param bytes
+  const f = Transformer.getFlatParams(model);
+  const b = new Uint8Array(f.buffer, f.byteOffset, f.byteLength);
+  let h = 0x811c9dc5;
+  for (let i = 0; i < b.length; i++) { h ^= b[i]; h = Math.imul(h, 0x01000193); }
+  return h >>> 0;
+}
+// roster wire: [int32 -4][int32 step][int32 n][int32 peerNums...]
+const ROSTER_SENTINEL = -4;
+function broadcastRoster(step, ids) {
+  const nums = ids.map(id => +id.slice(1));     // "p12" -> 12
+  const buf = new ArrayBuffer(12 + 4 * nums.length);
+  const iv = new Int32Array(buf);
+  iv[0] = ROSTER_SENTINEL; iv[1] = step; iv[2] = nums.length;
+  iv.set(nums, 3);
+  broadcast(buf);
+}
 const waiters = new Set();                       // pending waitForGrads checkers
 function wake() { for (const w of waiters) w(); }
 function onGrad(peerId, buf) {
   const step = new Int32Array(buf, 0, 1)[0];
+  if (step === FRAG_SENTINEL) {                  // chunk of a large message
+    const whole = onFragment(peerId, buf);
+    if (whole) onGrad(peerId, whole);
+    return;
+  }
   if (step === CFG_SENTINEL) { onConfig(peerId, buf); return; }
   if (step === CKPT_SENTINEL) {                  // a peer pushed a checkpoint
     if (training) { log(`ignored checkpoint from ${nmeOf(peerId)} (training in progress)`); return; }
@@ -258,27 +412,35 @@ function onGrad(peerId, buf) {
     catch (e) { log(`bad checkpoint from ${nmeOf(peerId)}: ${e.message}`); }
     return;
   }
-  const grad = new Float32Array(buf.slice(4));
+  if (step === ROSTER_SENTINEL) {                // the leader's contributor set for a step
+    if (training && peerId !== leaderId) return;  // only the run's leader may steer sync
+    const iv = new Int32Array(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));
   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 });
   wake();                                        // resolve waits immediately (no polling)
 }
-function broadcastGrad(step, grad) { const b = packGrad(step, grad); for (const dc of chans.values()) if (dc.readyState === "open") dc.send(b); }
+function broadcastGrad(step, whash, loss, grad) { broadcast(packGrad(step, whash, 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
 // wait — a device dying no longer costs the full timeout every step.
-function waitForGrads(step, cohort, timeoutMs = 8000) {
+function waitFor(pred, timeoutMs) {              // resolves true if pred held, false on timeout
   return new Promise((resolve) => {
     const t0 = Date.now();
     let timer = null;
     const check = () => {
-      const live = cohort.filter(id => chans.has(id));   // prune departed peers
-      const got = incoming.get(step) || new Map();
-      const have = live.filter(id => got.has(id));
-      if (have.length === live.length || Date.now() - t0 > timeoutMs) {
+      const ok = pred();
+      if (ok || Date.now() - t0 > timeoutMs) {
         waiters.delete(check); clearInterval(timer);
-        resolve(have.map(id => got.get(id)));
+        resolve(!!ok);
       }
     };
     waiters.add(check);
@@ -286,6 +448,15 @@ function waitForGrads(step, cohort, timeoutMs = 8000) {
     check();
   });
 }
+async function waitForGradIds(step, cohort, timeoutMs = 8000) {
+  await waitFor(() => {
+    const live = cohort.filter(id => chans.has(id));     // prune departed peers
+    const got = incoming.get(step) || new Map();
+    return live.every(id => got.has(id));
+  }, timeoutMs);
+  const got = incoming.get(step) || new Map();
+  return cohort.filter(id => chans.has(id) && got.has(id));
+}
 
 // ---- compute: one async training step THROUGH the verified units -----------
 async function localStep() {
@@ -297,40 +468,117 @@ async function localStep() {
 async function train(cfg) {
   if (training) return; training = true; ui.start.disabled = true;
   cfgSliders().forEach(el => el.disabled = true);
+  incoming.clear(); rosters.clear(); peerHashes.clear();  // no stale grads from a previous run
+  const iLead = leaderId === null;                     // set by Start (null) or onConfig (starter's id)
   const cohort = [...chans.keys()];                    // lock the cohort (departed peers are pruned per-step)
   const steps = cfg.steps;
   const opt = TrainCore.makeAdam(model.nParams, { lr: cfg.lr });
   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}`);
+      `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)}` : "") +
+      ` · data: ${Transformer.datasetName()}` +
+      ` · grad payload ${(model.nParams * 4 / 1048576).toFixed(1)} MB/step/peer`);
+  if ((cfg.c > 64 || cfg.t > 64) && cohort.length + 1 < 4)
+    log(`⚠ large model (width ${cfg.c}, seq ${cfg.t}) with only ${cohort.length + 1} device(s) — ` +
+        `steps will be slow and the effective batch small. 4+ devices recommended at this size.`);
+  let halted = null;                                   // sync-guard stop reason
+  const contrib = new Map([[myId, 0]]);                // peerId -> grads contributed
+  for (const id of cohort) contrib.set(id, 0);
+  try {
   for (let s = 0; s < steps; s++) {
+    const whash = hashWeights();                       // pre-step fingerprint, sent with the grad
     const { loss, grad } = await localStep();
-    broadcastGrad(s, grad);
-    const remote = cohort.length ? await waitForGrads(s, cohort) : [];
-    const all = [grad, ...remote];
+    broadcastGrad(s, whash, loss, grad);
+    let all;
+    if (!cohort.length) {
+      all = [grad];                                    // solo run
+    } else if (iLead) {
+      const ids = await waitForGradIds(s, cohort);
+      // sync guard: publish the exact contributor set so every device
+      // averages the same gradients (or stops) — never a silent fork
+      broadcastRoster(s, [myId, ...ids]);
+      all = [grad, ...ids.map(id => incoming.get(s).get(id))];
+      rosters.set(s, [myId, ...ids]);
+    } else {
+      // follower: apply the leader's roster verbatim, or stop
+      await waitFor(() => rosters.has(s) || !chans.has(leaderId), 15000);
+      if (!rosters.has(s)) {
+        halted = chans.has(leaderId)
+          ? `no roster from ${nmeOf(leaderId)} for step ${s + 1}`
+          : `the sync leader (${nmeOf(leaderId)}) disconnected`;
+        break;
+      }
+      const roster = rosters.get(s);
+      const need = roster.filter(id => id !== myId);
+      const ok = await waitFor(() => need.every(id => (incoming.get(s) || new Map()).has(id)), 10000);
+      if (!ok) {
+        halted = `missing a roster gradient at step ${s + 1} — applying a partial average would fork the weights`;
+        break;
+      }
+      all = need.map(id => incoming.get(s).get(id));
+      if (roster.includes(myId)) all.unshift(grad);    // leader may have dropped my late grad — then I skip it too
+    }
+    // divergence check: every contributor's pre-step weight hash must match mine
+    const hs = peerHashes.get(s);
+    const roster = rosters.get(s) || [myId];
+    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; }
+    }
+    // cluster-average loss over this step's actual contributors
+    let lossSum = 0, lossN = 0;
+    for (const id of roster) {
+      if (id === myId) { lossSum += loss; lossN++; contrib.set(myId, (contrib.get(myId) || 0) + 1); }
+      else if (hs && hs.has(id)) { lossSum += hs.get(id).loss; lossN++; contrib.set(id, (contrib.get(id) || 0) + 1); }
+    }
+    const clusterLoss = lossSum / Math.max(1, lossN);
     const avg = TrainCore.averageGrads(all);
     const upd = opt.step(avg);                      // DaisyAdam on the cluster-avg grad
     Transformer.applyUpdate(model, upd);            // W -= upd (lr folded into upd)
-    incoming.delete(s);
+    incoming.delete(s); rosters.delete(s); peerHashes.delete(s);
     trainedSteps++;
     if (s % 10 === 0 || s === steps - 1) {
-      ui.loss.textContent = loss.toFixed(5);
+      ui.loss.textContent = clusterLoss.toFixed(5);
       ui.step.textContent = `${s + 1} / ${steps}`;
       ui.bar.style.width = `${Math.round(100 * (s + 1) / steps)}%`;
       await new Promise(r => setTimeout(r, 0));         // yield to UI
     }
+    if (cohort.length && (s + 1) % 50 === 0)          // who is contributing what
+      log(`step ${s + 1}: contributions — ` +
+          [...contrib.entries()].map(([id, n]) => `${id === myId ? "me" : nmeOf(id)} ${n}/${s + 1}`).join(", "));
   }
-  log(`training done — final loss ${ui.loss.textContent}`);
-  try {
-    const sample = await Transformer.generate(model, "the ", 70);
-    ui.diff.textContent = `the model speaks: “${sample.trim()}”`;
-  } catch (e) {
-    ui.diff.textContent = "done — trained through the verified units; all peers share one model.";
+  } catch (e) {                                       // a device that errors REPORTS it
+    halted = `error during training: ${e.message}`;   // (instead of freezing at "—")
+    console.error(e);
+  }
+  if (cohort.length)
+    log(`contribution totals — ` +
+        [...contrib.entries()].map(([id, n]) => `${id === myId ? "me" : nmeOf(id)} ${n} grad(s)`).join(", "));
+  if (halted) {
+    log(`SYNC GUARD: stopped — ${halted}. The model here is intact (${trainedSteps} steps); ` +
+        `to re-sync the group, load/push a checkpoint and start again.`);
+    ui.diff.textContent = "stopped by the sync guard — see log";
+  } else {
+    log(`training done — final loss ${ui.loss.textContent}`);
+    try {
+      const sample = await Transformer.generate(model, "the ", 70);
+      ui.diff.textContent = `the model speaks: “${sample.trim()}”`;
+    } catch (e) {
+      ui.diff.textContent = "done — trained through the verified units; all peers share one model.";
+    }
   }
+  incoming.clear(); rosters.clear(); peerHashes.clear();
+  leaderId = null;
   training = false;
-  ui.save.disabled = false;
+  modelReady();
   ui.start.disabled = false;
   cfgSliders().forEach(el => el.disabled = false);
 }
+function modelReady() {                 // trained weights exist: enable save/kit/generate
+  ui.save.disabled = false;
+  ui.kit.disabled = false;
+  ui.genBtn.disabled = false;
+}
 function cfgSliders() { return [ui.cfgC, ui.cfgT, ui.cfgB, ui.cfgSteps, ui.cfgLr]; }
 
 // (re)build the mini transformer — deterministic shared init (same seeds on
@@ -386,12 +634,25 @@ function buildModel(cfg) {
     ui.start.disabled = true;
     return;                                              // no signaling, no training
   }
-  ui.backend.textContent = `${compute.backend.toUpperCase()} — ${compute.label} · through verified INT8 units`;
+  ui.backend.textContent = `${compute.backend.toUpperCase()} — ${compute.label} · 3×INT8 fast-accurate through verified units`;
+  // tokenizer must load BEFORE the model is built (it sets the vocab size)
+  try {
+    const name = await Transformer.loadTokenizer();
+    ui.tokenizer.textContent = name;
+    log(`tokenizer: ${name}`);
+  } catch (e) {
+    ui.tokenizer.textContent = Transformer.tokenizerName();
+    log(`tokenizer.json unavailable (${e.message}) — using the ${Transformer.tokenizerName()} vocab`);
+  }
+  if (compute.backend === "cpu" && Transformer.vocabSize() > 1000)
+    log(`⚠ CPU backend with a ${Transformer.vocabSize()}-token vocab — steps will take seconds; ` +
+        `a WebGPU-capable browser will be much faster`);
   // slider readouts
   for (const [el, v] of [[ui.cfgC, "vcfgC"], [ui.cfgT, "vcfgT"], [ui.cfgB, "vcfgB"],
                          [ui.cfgSteps, "vcfgSteps"], [ui.cfgLr, "vcfgLr"]])
     el.oninput = () => document.getElementById(v).textContent = el.value;
   buildModel(readCfgFromUI());
+  ui.start.disabled = false;                             // solo training works too
   ui.me.textContent = deviceName;
   if (room()) {
     ui.roomInfo.style.display = "";
@@ -404,13 +665,32 @@ function buildModel(cfg) {
   }
   updatePeers();
   connectSignaling();
+  // dataset: stream FineWeb-Edu from HuggingFace; built-in corpus if offline
+  ui.dataset.textContent = "streaming FineWeb-Edu…";
+  Transformer.streamFineWebEdu().then(({ name, chars }) => {
+    ui.dataset.textContent = "FineWeb-Edu (streamed)";
+    log(`dataset: ${name} — ${(chars / 1000).toFixed(0)}k chars loaded (each device streams its own slice)`);
+  }).catch((e) => {
+    ui.dataset.textContent = "built-in corpus (offline)";
+    log(`dataset: FineWeb-Edu stream unavailable (${e.message}) — using the built-in corpus`);
+  });
   ui.start.onclick = () => {
     const cfg = readCfgFromUI();
+    leaderId = null;                                     // I pressed Start: I lead sync
     buildModel(cfg);
     broadcastConfig(cfg);                                // everyone follows these settings
     train(cfg);
   };
   ui.save.onclick = saveCheckpoint;
+  ui.kit.onclick = downloadInferenceKit;
+  ui.genBtn.onclick = async () => {
+    if (training) { log("can't generate mid-training"); return; }
+    ui.genBtn.disabled = true;
+    ui.genOut.textContent = "generating…";
+    try { ui.genOut.textContent = await Transformer.generate(model, ui.genPrompt.value || "the ", 150); }
+    catch (e) { ui.genOut.textContent = `error: ${e.message}`; }
+    ui.genBtn.disabled = false;
+  };
   ui.loadBtn.onclick = () => { if (training) { log("can't load a checkpoint mid-training"); return; } ui.load.click(); };
   ui.load.onchange = async () => {
     const f = ui.load.files[0]; ui.load.value = "";

web/public/index.html

@@ -90,6 +90,8 @@
     <div class="device" id="me">—</div>
     <div class="row" style="margin-top:8px"><span class="k">Status</span><span class="v" id="status">starting…</span></div>
     <div class="row"><span class="k">Compute</span><span class="v" id="backend">detecting…</span></div>
+    <div class="row"><span class="k">Dataset</span><span class="v" id="dataset">—</span></div>
+    <div class="row"><span class="k">Tokenizer</span><span class="v" id="tokenizer">loading…</span></div>
   </div>
 
   <div class="card">
@@ -106,21 +108,22 @@
   <div class="card">
     <div class="lbl">🎛 Training settings</div>
     <label class="slbl">Model width <span class="sval" id="vcfgC">32</span></label>
-    <input type="range" id="cfgC" min="16" max="64" step="16" value="32">
+    <input type="range" id="cfgC" min="16" max="128" step="16" value="32">
     <label class="slbl">Sequence length <span class="sval" id="vcfgT">32</span></label>
-    <input type="range" id="cfgT" min="16" max="64" step="16" value="32">
+    <input type="range" id="cfgT" min="16" max="128" step="16" value="32">
+    <p class="note" style="margin:.4rem 0 0"><b>Width or sequence above 64?</b> Bring more devices — big settings on 1–3 devices mean slow steps and a small effective batch. 4+ devices recommended.</p>
     <label class="slbl">Batch per device <span class="sval" id="vcfgB">8</span></label>
     <input type="range" id="cfgB" min="2" max="32" step="2" value="8">
     <label class="slbl">Steps <span class="sval" id="vcfgSteps">300</span></label>
     <input type="range" id="cfgSteps" min="100" max="10000" step="100" value="300">
     <label class="slbl">Learning rate ×1000 <span class="sval" id="vcfgLr">20</span></label>
     <input type="range" id="cfgLr" min="5" max="50" step="5" value="20">
-    <p class="note" style="margin:.5rem 0 0">A mini transformer language model — attention, MLP blocks, next-character prediction — every multiply through the verified INT8 units. Whoever presses Start sets the settings for the whole group; total batch scales with every device that joins.</p>
+    <p class="note" style="margin:.5rem 0 0">A mini transformer language model — attention, MLP blocks, next-character prediction — every multiply through the verified INT8 units. Training text is streamed from <b>FineWeb-Edu</b> (HuggingFace); offline devices fall back to a built-in corpus. Whoever presses Start sets the settings for the whole group; total batch scales with every device that joins.</p>
   </div>
 
   <div class="card" style="text-align:center">
     <button id="start" disabled>Start training</button>
-    <p class="note" style="margin:.6rem 0 0">Enabled once another device joins. (Or open a second tab to try it.)</p>
+    <p class="note" style="margin:.6rem 0 0">Works solo — every device that joins adds its batch to the group.</p>
   </div>
 
   <div class="card">
@@ -131,10 +134,22 @@
     <div class="row" style="margin-top:10px"><span class="diff" id="diff"></span></div>
   </div>
 
+  <div class="card">
+    <div class="lbl">🗣 Test the model</div>
+    <div style="display:flex;gap:8px;flex-wrap:wrap">
+      <input type="text" id="genPrompt" value="the " spellcheck="false"
+             style="flex:1;min-width:160px;padding:10px 12px;border-radius:8px;border:1px solid var(--card-border);background:transparent;color:inherit;font-family:'Courier New',monospace">
+      <button id="genBtn" disabled style="padding:10px 18px">Generate</button>
+    </div>
+    <pre id="genOut" style="margin-top:10px;min-height:44px"></pre>
+    <p class="note" style="margin:.5rem 0 0">Enabled after training finishes or a checkpoint is loaded. The “inference kit” below downloads a single HTML file with these weights baked in — open it anywhere to run generations offline.</p>
+  </div>
+
   <div class="card">
     <div class="lbl">💾 Model checkpoint</div>
     <div style="display:flex;gap:10px;flex-wrap:wrap;justify-content:center">
       <button id="save" disabled>Download model (.pt)</button>
+      <button id="kit" disabled>Download inference kit</button>
       <button id="loadBtn">Load checkpoint…</button>
       <input type="file" id="load" accept=".pt" style="display:none">
     </div>

web/public/transformer.js

@@ -28,29 +28,96 @@
       s += `the ${pick(W_ADJ)} ${pick(W_NOUN)} ${pick(W_VERB)} ${pick(W_PREP)} the ${pick(W_ADJ)} ${pick(W_NOUN)}. `;
     return s;
   }
-  const CORPUS = buildCorpus();
-  const CHARS = [...new Set(CORPUS)].sort();
-  const VOCAB = CHARS.length;
-  const STOI = Object.fromEntries(CHARS.map((c, i) => [c, i]));
-  const IDS = new Int32Array(CORPUS.length);
-  for (let i = 0; i < CORPUS.length; i++) IDS[i] = STOI[CORPUS[i]];
+  // ---- tokenizer -------------------------------------------------------------
+  // Spikewhale tokenizer (tokenizer.json): byte-level greedy longest-match
+  // ("length-max"), ~16.5k tokens. Until it loads (or if the file is missing)
+  // a 96-char byte-level vocab keeps the app working — but ALL devices in a
+  // group must use the same tokenizer (the config broadcast enforces it).
+  const FALLBACK_CHARS = [...Array(95)].map((_, i) => String.fromCharCode(32 + i)).concat(["\n"]);
+  let tok = {
+    name: "char-96 (fallback)",
+    vocab: Object.fromEntries(FALLBACK_CHARS.map((c, i) => [c, i])),
+    ids: FALLBACK_CHARS, maxLen: 1, size: FALLBACK_CHARS.length,
+    unk: 0, specials: new Set(),
+  };
+  tok.unk = tok.vocab[" "];
+  function vocabSize() { return tok.size; }
+  function tokenizerName() { return tok.name; }
+  function loadTokenizerData(d) {              // plain {vocab, vocab_size, max_token_len}
+    const ids = new Array(d.vocab_size);
+    for (const [t, i] of Object.entries(d.vocab)) ids[i] = t;
+    tok = { name: `Spikewhale length-max (${d.vocab_size} tokens)`,
+            vocab: d.vocab, ids, maxLen: d.max_token_len || 24, size: d.vocab_size,
+            unk: d.vocab["<unk>"] ?? 1,
+            specials: new Set(["<pad>", "<unk>", "<bos>", "<eos>", ...(d.special_tokens || [])]) };
+    IDS = encode(CORPUS);                      // re-tokenize whatever corpus is loaded
+    return tok.name;
+  }
+  async function loadTokenizer(url) {
+    const r = await fetch(url || "tokenizer.json");
+    if (!r.ok) throw new Error(`tokenizer.json HTTP ${r.status}`);
+    return loadTokenizerData(await r.json());
+  }
+  function toLatin1(s) { const b = new TextEncoder().encode(s); let o = ""; for (const x of b) o += String.fromCharCode(x); return o; }
+  function encode(text) {                      // greedy longest match over bytes
+    const s = toLatin1(text), out = [];
+    let i = 0;
+    while (i < s.length) {
+      let m = null;
+      for (let L = Math.min(tok.maxLen, s.length - i); L > 0; L--) {
+        const sub = s.substr(i, L);
+        if (sub in tok.vocab) { m = sub; break; }
+      }
+      if (m === null) { out.push(tok.unk); i++; continue; }
+      out.push(tok.vocab[m]); i += m.length;
+    }
+    return Int32Array.from(out);
+  }
+  function decode(idArr) {
+    let s = "";
+    for (const id of idArr) {
+      const t = tok.ids[id];
+      if (t === undefined || tok.specials.has(t)) continue;
+      s += t;
+    }
+    const bytes = Uint8Array.from([...s].map(c => c.charCodeAt(0)));
+    return new TextDecoder().decode(bytes);
+  }
+  let CORPUS = buildCorpus();
+  let IDS = encode(CORPUS);
+  let DATASET = "built-in corpus";
+
+  // Stream real training text: FineWeb-Edu, via the public HuggingFace
+  // datasets-server rows API. Each device pulls its own random slice (that's
+  // data parallelism — batches were always per-device anyway). Offline or on
+  // API failure the built-in corpus stays in place.
+  async function streamFineWebEdu() {
+    const offset = Math.floor(Math.random() * 9900);
+    const url = "https://datasets-server.huggingface.co/rows?dataset=HuggingFaceFW%2Ffineweb-edu" +
+                `&config=default&split=train&offset=${offset}&length=100`;
+    const r = await fetch(url);
+    if (!r.ok) throw new Error(`datasets-server HTTP ${r.status}`);
+    const j = await r.json();
+    const text = j.rows.map(x => x.row.text).join("\n").replace(/[^\x20-\x7e\n]/g, " ");
+    if (text.length < 10000) throw new Error("too little text returned");
+    CORPUS = text.slice(0, 500000);
+    IDS = encode(CORPUS);
+    DATASET = "FineWeb-Edu (streamed from HuggingFace)";
+    return { name: DATASET, chars: CORPUS.length };
+  }
+  function datasetName() { return DATASET; }
 
   // ---- verified matmul: float in -> quantize -> LUT multiply -> dequant -----
   // weights/projections may go through WebGPU; attention (per-head, many small
   // matmuls) uses the CPU LUT path — same verified units, no dispatch overhead.
+  // 3xINT8 fast-accurate GEMM (CUTLASS 3xTF32 scheme through the verified
+  // units): three exact LUT GEMMs recover ~14-bit precision from the 8-bit
+  // units — see Verified.lutMatmul3.
   async function vmm(Xf, Wf, m, k, n, ctx) {
-    const xq = V.quantize(Xf), wq = V.quantize(Wf);
-    const acc = await ctx.matmulInt8(xq.q, wq.q, m, k, n, ctx.L);
-    const dq = xq.scale * wq.scale, out = new Float32Array(m * n);
-    for (let i = 0; i < out.length; i++) out[i] = acc[i] * dq;
-    return out;
+    return V.lutMatmul3(Xf, Wf, m, k, n, ctx.L, ctx.matmulInt8);
   }
   function vmmCPU(Xf, Wf, m, k, n, ctx) {
-    const xq = V.quantize(Xf), wq = V.quantize(Wf);
-    const acc = V.lutMatmulJS(xq.q, wq.q, m, k, n, ctx.L);
-    const dq = xq.scale * wq.scale, out = new Float32Array(m * n);
-    for (let i = 0; i < out.length; i++) out[i] = acc[i] * dq;
-    return out;
+    return V.lutMatmul3JS(Xf, Wf, m, k, n, ctx.L);
   }
 
   // ---- layernorm (no affine) -------------------------------------------------
@@ -84,9 +151,9 @@
     const params = [], names = [];
     const add = (name, w) => { params.push(w); names.push(name); return w; };
     const m = {
-      cfg: { ...cfg, layers, heads, hidden, vocab: VOCAB },
+      cfg: { ...cfg, layers, heads, hidden, vocab: vocabSize() },
       ctx: { L, matmulInt8 },
-      emb: add("emb", mk(VOCAB * c, 0.08)),
+      emb: add("emb", mk(vocabSize() * c, 0.08)),
       pos: add("pos", mk(cfg.t * c, 0.02)),
       blocks: [], params, names,
     };
@@ -96,7 +163,7 @@
         Wv: add(`b${l}.Wv`, mk(c * c, 0.08)), Wo: add(`b${l}.Wo`, mk(c * c, 0.08)),
         W1: add(`b${l}.W1`, mk(c * hidden, 0.08)), W2: add(`b${l}.W2`, mk(hidden * c, 0.08)),
       });
-    m.Wu = add("Wu", mk(c * VOCAB, 0.08));
+    m.Wu = add("Wu", mk(c * vocabSize(), 0.08));
     m.nParams = params.reduce((a, p) => a + p.length, 0);
     return m;
   }
@@ -289,7 +356,7 @@
   // greedy sampling — watch the model actually speak
   async function generate(m, prompt, nChars) {
     const { t: T } = m.cfg;
-    let ids = [...prompt].map(ch => STOI[ch] ?? 0);
+    let ids = [...encode(prompt)];
     for (let step = 0; step < nChars; step++) {
       const win = ids.slice(-T);
       const X = new Int32Array(T), Y = new Int32Array(T);
@@ -302,10 +369,12 @@
       for (let j = 0; j < m.cfg.vocab; j++) if (logits[row + j] > bv) { bv = logits[row + j]; best = j; }
       ids.push(best);
     }
-    return ids.map(i => CHARS[i]).join("");
+    return decode(ids);
   }
 
-  const api = { init, trainStep, applyUpdate, getFlatParams, setFlatParams, generate, VOCAB, CORPUS };
+  const api = { init, trainStep, applyUpdate, getFlatParams, setFlatParams, generate,
+                streamFineWebEdu, datasetName, loadTokenizer, loadTokenizerData,
+                vocabSize, tokenizerName, encode, decode };
   if (typeof module !== "undefined" && module.exports) { TC = require("./traincore.js"); V = require("./verified_core.js"); module.exports = api; }
   else { TC = root.TrainCore; V = root.Verified; root.Transformer = api; }
 })(typeof self !== "undefined" ? self : this);

web/public/verified_core.js

@@ -28,6 +28,44 @@
     return C;
   }
 
+  // ---- 3xINT8 fast-accurate GEMM --------------------------------------------
+  // The CUTLASS example-27 "3xTF32" scheme, ported to the verified units:
+  // split each float into a coarse int8 part plus an int8-quantized residual,
+  // run three EXACT LUT GEMMs (hi·hi, hi·lo, lo·hi), drop the negligible
+  // lo·lo, and recombine. Same big/small decomposition NVIDIA uses to recover
+  // near-fp32 accuracy from TF32 tensor cores — here it recovers ~14-bit
+  // accuracy from the 8-bit units, at 3× the unit ops. Every product still
+  // goes through the verified mul8 LUT.
+  function quantize2(X) {
+    const hi = quantize(X);
+    const r = new Float32Array(X.length);
+    for (let i = 0; i < X.length; i++) r[i] = X[i] - hi.q[i] * hi.scale;
+    const lo = quantize(r);
+    return { hi, lo };
+  }
+  function combine3(hh, hl, lh, x, w, len) {
+    const out = new Float32Array(len);
+    const shh = x.hi.scale * w.hi.scale, shl = x.hi.scale * w.lo.scale, slh = x.lo.scale * w.hi.scale;
+    for (let i = 0; i < len; i++) out[i] = hh[i] * shh + hl[i] * shl + lh[i] * slh;
+    return out;
+  }
+  function lutMatmul3JS(Xf, Wf, m, k, n, L) {          // sync, CPU LUT path
+    const x = quantize2(Xf), w = quantize2(Wf);
+    return combine3(lutMatmulJS(x.hi.q, w.hi.q, m, k, n, L),
+                    lutMatmulJS(x.hi.q, w.lo.q, m, k, n, L),
+                    lutMatmulJS(x.lo.q, w.hi.q, m, k, n, L), x, w, m * n);
+  }
+  async function lutMatmul3(Xf, Wf, m, k, n, L, matmulInt8) {  // any backend
+    const x = quantize2(Xf), w = quantize2(Wf);
+    const mm = matmulInt8 || lutMatmulJS;
+    const [hh, hl, lh] = await Promise.all([
+      mm(x.hi.q, w.hi.q, m, k, n, L),
+      mm(x.hi.q, w.lo.q, m, k, n, L),
+      mm(x.lo.q, w.hi.q, m, k, n, L),
+    ]);
+    return combine3(hh, hl, lh, x, w, m * n);
+  }
+
   // one verified layer forward; returns float out (+ cache for STE backward).
   // Every product goes through the verified INT8 multiply (mul8 LUT) with exact
   // int32 accumulation — i.e. an emulated INT8 tensor-core GEMM — then dequant.
@@ -78,7 +116,7 @@
     for (let j = 0; j < W2.length; j++) W2[j] -= lr * gAvg[W1.length + j];
   }
 
-  const api = { quantize, lutMatmulJS, linearFwd, forward, backward, splitApply };
+  const api = { quantize, quantize2, lutMatmulJS, lutMatmul3JS, lutMatmul3, 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

@@ -3,10 +3,19 @@
 // without WebGPU (e.g. old PCs via Supermium). initCompute() returns
 //   { backend, label, matmulInt8(Xq, Wq, m, k, n, L) -> Int32Array }
 // matching Verified.lutMatmulJS, so the trainer is device-blind.
+//
+// High-throughput path: when the browser exposes WGSL's
+// packed_4x8_integer_dot_product feature, we use dot4I8Packed — it compiles to
+// the GPU's DP4A/INT8 dot-product hardware (the same units tensor-core INT8
+// paths are built on): 4 exact int8 MACs per instruction, int32 accumulation,
+// and 4× less memory traffic from packing. Because int8×int8→int32 is exact,
+// it is bit-identical to the verified mul8 LUT — and we PROVE that at init by
+// cross-checking random matmuls against the LUT before trusting it. If the
+// hardware ever disagrees with the units, we fall back to the LUT shader.
 (function (root) {
   "use strict";
 
-  const WGSL = `
+  const WGSL_LUT = `
     @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>;   // 65536 signed products
@@ -26,6 +35,25 @@
       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;
+    }`;
+
   async function loadLUTs(base) {
     base = base || "";
     const [mulB, reqB, reluB, meta] = await Promise.all([
@@ -38,6 +66,21 @@
              relu: new Int8Array(reluB), shift: meta.shift };
   }
 
+  // pack a row-major int8 matrix (rows×cols) into u32 words of 4 bytes along
+  // cols, zero-padded to kw words per row (zeros contribute 0 to the dot)
+  function packRows(Q, rows, cols, kw) {
+    const out = new Uint32Array(rows * kw);
+    const bytes = new Uint8Array(out.buffer);
+    for (let r = 0; r < rows; r++)
+      for (let c = 0; c < cols; c++) bytes[(r * kw * 4) + c] = Q[r * cols + c] & 0xFF;
+    return out;
+  }
+  function transposeI8(Q, rows, cols) {
+    const out = new Int8Array(rows * cols);
+    for (let r = 0; r < rows; r++) for (let c = 0; c < cols; c++) out[c * rows + r] = Q[r * cols + c];
+    return out;
+  }
+
   async function initCompute(L) {
     const cpu = { backend: "cpu", label: "CPU (JS)",
                   matmulInt8: (Xq, Wq, m, k, n, LL) => root.Verified.lutMatmulJS(Xq, Wq, m, k, n, LL) };
@@ -46,32 +89,45 @@
       const adapter = await navigator.gpu.requestAdapter();
       if (!adapter) return cpu;
       const device = await adapter.requestDevice();
-      const module = device.createShaderModule({ code: WGSL });
-      const pipeline = device.createComputePipeline({ layout: "auto", compute: { module, entryPoint: "main" } });
-      // upload the multiply LUT once (as i32)
+      const info = adapter.info || {};
+      const gpuName = info.description || info.vendor || "WebGPU";
+
+      // LUT pipeline (always built — the fallback and the verification oracle)
+      const lutModule = device.createShaderModule({ code: WGSL_LUT });
+      const lutPipe = device.createComputePipeline({ layout: "auto", compute: { module: lutModule, entryPoint: "main" } });
       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 info = adapter.info || {};
-      return { backend: "webgpu", label: info.description || info.vendor || "WebGPU",
-               matmulInt8: (Xq, Wq, m, k, n) => gpuMatmul(device, pipeline, lutBuf, Xq, Wq, m, k, n) };
+      const viaLUT = { backend: "webgpu", label: `${gpuName} (LUT shader)`,
+                       matmulInt8: (Xq, Wq, m, k, n) => gpuMatmulLUT(device, lutPipe, lutBuf, Xq, Wq, m, k, n) };
+
+      // DP4A pipeline — only if the WGSL feature exists AND it reproduces the
+      // verified units exactly on random self-tests
+      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; }
+      }
+      return { backend: "webgpu", label: `${gpuName} (DP4A int8 dot — HW verified vs units)`, matmulInt8: dp4mm };
     } catch (e) { console.warn("WebGPU init failed, CPU fallback:", e); return cpu; }
   }
 
-  async function gpuMatmul(device, pipeline, lutBuf, Xq, Wq, m, k, n) {
-    const X32 = Int32Array.from(Xq), W32 = Int32Array.from(Wq);   // byte -> i32
-    const bufX = mk(device, X32.byteLength, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
-    const bufW = mk(device, W32.byteLength, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
+  // shared dispatch/readback plumbing
+  async function runPass(device, pipeline, entries, m, n) {
     const bytesC = m * n * 4;
     const bufC = mk(device, bytesC, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
-    const bufD = mk(device, 16, GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST);
-    device.queue.writeBuffer(bufX, 0, X32);
-    device.queue.writeBuffer(bufW, 0, W32);
-    device.queue.writeBuffer(bufD, 0, new Uint32Array([m, k, n]));
-    const bind = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [
-      { binding: 0, resource: { buffer: bufX } }, { binding: 1, resource: { buffer: bufW } },
-      { binding: 2, resource: { buffer: lutBuf } }, { binding: 3, resource: { buffer: bufC } },
-      { binding: 4, resource: { buffer: bufD } } ] });
+    const bind = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0),
+      entries: entries(bufC) });
     const enc = device.createCommandEncoder();
     const pass = enc.beginComputePass();
     pass.setPipeline(pipeline); pass.setBindGroup(0, bind);
@@ -82,10 +138,42 @@
     await read.mapAsync(GPUMapMode.READ);
     const out = new Int32Array(read.getMappedRange().slice(0));
     read.unmap();
-    [bufX, bufW, bufC, bufD, read].forEach(b => b.destroy());
-    return out;
+    return { out, bufC, read };
+  }
+
+  async function gpuMatmulLUT(device, pipeline, lutBuf, Xq, Wq, m, k, n) {
+    const X32 = Int32Array.from(Xq), W32 = Int32Array.from(Wq);   // byte -> i32
+    const bufX = up(device, X32, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
+    const bufW = up(device, W32, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST);
+    const bufD = up(device, new Uint32Array([m, k, 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: lutBuf } }, { binding: 3, resource: { buffer: bufC } },
+      { binding: 4, resource: { buffer: bufD } } ], m, n);
+    [bufX, bufW, bufD, r.bufC, r.read].forEach(b => b.destroy());
+    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;
   }
+
   function mk(device, size, usage) { return device.createBuffer({ size, usage }); }
+  function up(device, arr, usage) {
+    const b = mk(device, Math.max(16, arr.byteLength), usage);
+    device.queue.writeBuffer(b, 0, arr);
+    return b;
+  }
 
   root.Compute = { initCompute, loadLUTs };
 })(self);

web/test_transformer.js

@@ -20,7 +20,7 @@ const cfg = { c: 32, t: 32, b: 8, layers: 2, heads: 2, steps: 120, lr: 0.02 };
 (async function () {
   const A = X.init(cfg, L, matmulInt8);
   const B = X.init(cfg, L, matmulInt8);
-  console.log(`vocab=${X.VOCAB}, params=${A.nParams}, baseline loss=${Math.log(X.VOCAB).toFixed(3)}`);
+  console.log(`vocab=${X.vocabSize()}, params=${A.nParams}, baseline loss=${Math.log(X.vocabSize()).toFixed(3)}`);
   const oa = T.makeAdam(A.nParams, { lr: cfg.lr });
   const ob = T.makeAdam(B.nParams, { lr: cfg.lr });
   let first = 0, loss = 0;
@@ -40,7 +40,7 @@ const cfg = { c: 32, t: 32, b: 8, layers: 2, heads: 2, steps: 120, lr: 0.02 };
   const sample = await X.generate(A, "the ", 60);
   console.log(`sample: "${sample}"`);
   console.log(`replica max param diff: ${diff.toExponential(3)}`);
-  const ok = loss < first * 0.75 && loss < Math.log(X.VOCAB) && diff === 0;
+  const ok = loss < first * 0.75 && loss < Math.log(X.vocabSize()) && diff === 0;
   console.log(ok ? "TRANSFORMER TEST PASSED" : "TRANSFORMER TEST FAILED");
   process.exit(ok ? 0 : 1);
 })();