web: B2B MLP chain (CUTLASS ex. 13 two-GEMM fusion + ex. 23 epilogue reduction), test_b2b.js, ten-suite npm test

web/TEST_RESULTS.md

@@ -108,6 +108,40 @@ width 32 (the shared operand is only 32 KB there — the saved quantize work
 scales quadratically with model width). Two-device live run: both replicas
 at step 71/300 with identical loss to the last digit, no sync-guard trips.
 
+## B2B MLP chain (CUTLASS ex. 13 two-GEMM fusion + ex. 23 epilogue reduction)
+
+The MLP's two GEMMs now run back-to-back on the GPU: gemm1 (ReLU fused) and a
+per-row |max| reduction share one command encoder, ~1 KB of absmax comes back
+to JS (scale derivation needs division, which WGSL only guarantees to 2.5 ULP
+— JS f64 division is exactly rounded and device-identical), then h1 is
+quantized ON-DEVICE and fed straight to gemm2. h1 returns to JS only because
+the STE backward needs it; it never goes up again.
+
+This required respeccing the intermediate quantize from `round(x / scale)` to
+`floor(f32(x * invScale) + 0.5)` — WGSL multiply/add are correctly rounded and
+floor/clamp exact, so the GPU kernel and the fround-stepped JS mirror agree
+bit-for-bit, and CPU-fallback devices run the mirror so mixed fleets stay
+bit-identical. The respec is a real (bounded) math change: old and new builds
+cannot co-train, and the per-step divergence guard stops such mixed groups.
+
+`test_b2b.js` (Node):
+
+```
+scales from the fused absmax bit-identical to quantizeRows (3958 rows incl. zero rows)
+respec moves an int8 by at most 1 step (1/112363 = 0.001% of values moved)
+chain gemm1 (hence h1 and the ReLU mask) byte-identical to the un-chained GEMM
+chain output equals the manual composition of its stages; deterministic
+convergence unchanged: old 2.0500 vs new 2.0512 final loss (0.1% apart, 40 steps)
+```
+
+On GPU: both chain variants (LUT shader and DP4A) pass an exact `!==` init
+gate against the mirror chain over ragged shapes including pack-tail padding.
+Discriminating proof that the kernel implements the RESPEC and not the old
+spec: a searched-for boundary input (int8 68→69 under the respec) run through
+the GPU chain matches the new-spec mirror exactly and differs from the
+old-spec composition. Two-device live run: step 141/300 with identical loss
+(6.23958) on both replicas, per-step weight-hash divergence checks silent.
+
 One bug was caught during the rework, by the bit-identity check itself: the
 fused q+k+v sum initially ran in f64 and rounded once, where the old code
 rounded to f32 after each add — a last-ulp fork that would have split

web/package.json

@@ -6,7 +6,7 @@
   "author": "Dean Byrne (Quazim0t0) / DaisyChainAI",
   "scripts": {
     "start": "node server.js",
-    "test": "node test_core.js && node test_verified.js && node test_ieee.js && node test_gates.js && node test_metamorphic.js && node test_corpus.js && node test_optimizer.js && node test_transformer.js && node test_unit_backward.js"
+    "test": "node test_core.js && node test_verified.js && node test_ieee.js && node test_gates.js && node test_metamorphic.js && node test_corpus.js && node test_b2b.js && node test_optimizer.js && node test_transformer.js && node test_unit_backward.js"
   },
   "dependencies": {
     "hyparquet": "^1.26.2",

web/public/transformer.js

@@ -182,6 +182,7 @@
       cfg: { ...cfg, layers, heads, hidden, vocab: vocabSize() },
       ctx: { L, bgemm: (engine && engine.bgemm) || null,
              att: (engine && engine.att) || null, fgemm: (engine && engine.fgemm) || null,
+             mlp: (engine && engine.mlp) || null,
              audit: audit || null, unitBackward: !!cfg.unitBackward },
       emb: add("emb", mk(vocabSize() * c, 0.08)),
       pos: add("pos", mk(cfg.t * c, 0.02)),
@@ -292,11 +293,16 @@
       for (let i = 0; i < x2.length; i++) x2[i] = x[i] + attnOut[i];
       cb.x2 = x2;
       const l2 = lnFwd(x2, BT, C); cb.ln2 = l2;
-      const h1 = await vmm(l2.y, bl.W1, BT, C, hidden, ctx, true);  // ReLU fused in the epilogue
+      // CUTLASS ex. 13 + 23: both MLP GEMMs run back-to-back on the GPU. The
+      // intermediate h1 is quantized ON-DEVICE (exact-gated respec — see
+      // vmlpBlock in verified_core.js) and only its per-row absmax (~1KB)
+      // visits JS between the GEMMs; h1 itself comes back solely because the
+      // STE backward needs it. CPU devices run the bit-identical mirror chain.
+      const { h1, out: mlpOut } = await V.vmlpBlock(l2.y, bl.W1, bl.W2,
+        { m: BT, k: C, h: hidden, n: C }, ctx.L, ctx.mlp, ctx.audit);
       const mask = new Uint8Array(h1.length);
       for (let i = 0; i < h1.length; i++) if (h1[i] > 0) mask[i] = 1;
       cb.h1 = h1; cb.mask = mask;
-      const mlpOut = await vmm(h1, bl.W2, BT, hidden, C, ctx);
       x = new Float32Array(BT * C);
       for (let i = 0; i < x.length; i++) x[i] = x2[i] + mlpOut[i];
       cache.blocks.push(cb);

web/public/verified_core.js

@@ -100,6 +100,82 @@
     return { q, s };
   }
 
+  // ---- B2B MLP chain (CUTLASS ex. 13 two-GEMM fusion + ex. 23 epilogue
+  // reduction), respecced for cross-device exactness ---------------------------
+  // The MLP is the one back-to-back GEMM pair with no layernorm/softmax between
+  // (ReLU is already fused in the epilogue), so the intermediate h1 can be
+  // quantized ON the GPU and fed straight to the second GEMM. Two rules make
+  // that fleet-safe:
+  //   1. The per-row |max| (ex. 23) uses only comparisons — exact on any
+  //      hardware, order-independent — and comes back to JS as ~1KB.
+  //   2. Scale DERIVATION (two divisions) stays in JS f64, which IEEE requires
+  //      to be exactly rounded and is therefore identical on every device.
+  //      WGSL division is only 2.5 ULP — a fork waiting to happen — but WGSL
+  //      multiply/add are correctly rounded and floor/clamp are exact. So the
+  //      quantize step is respecced from  round(x / scale)  to
+  //      floor(f32(x * invScale) + 0.5)  — floor(x+0.5) IS Math.round's tie
+  //      rule — and the fround-stepped mirror below is bit-identical to the
+  //      GPU kernel, which is exact-gated against it at init.
+  // NOTE this changes which int8 a value on a rounding boundary lands on
+  // (≤1 step) vs quantizeRows, so old and new builds cannot co-train — the
+  // divergence guard stops such mixed groups by design.
+  function rowAbsMax(X, rows, cols) {
+    const mx = new Float32Array(rows);
+    for (let r = 0; r < rows; r++) {
+      let m = 0;
+      for (let c = 0; c < cols; c++) { const a = Math.abs(X[r * cols + c]); if (a > m) m = a; }
+      mx[r] = m;
+    }
+    return mx;
+  }
+  function scalesFromAbsMax(mx) {                // f64 divisions: exactly rounded, device-identical
+    const scale = new Float32Array(mx.length), inv = new Float32Array(mx.length);
+    for (let i = 0; i < mx.length; i++) {
+      scale[i] = Math.max(mx[i] / 127, 1e-8);
+      inv[i] = 1 / scale[i];                     // recip of the STORED f32 scale
+    }
+    return { scale, inv };
+  }
+  function quantizeRowsInv(X, rows, cols, inv) { // bit-exact mirror of the GPU quantize kernel
+    const q = new Int8Array(rows * cols);
+    for (let r = 0; r < rows; r++) {
+      const iv = inv[r];
+      for (let c = 0; c < cols; c++) {
+        const n = Math.floor(f32(f32(X[r * cols + c] * iv) + 0.5));
+        q[r * cols + c] = n < -128 ? -128 : n > 127 ? 127 : n;
+      }
+    }
+    return q;
+  }
+  // the chained MLP: X @ W1 -> ReLU (fused) -> absmax -> quantize -> @ W2.
+  //   d = { m, k, h, n };  gpuMlp (from webgpu.js) runs both GEMMs + the
+  // on-GPU quantize with one tiny absmax readback between; without it the CPU
+  // mirror chain runs — SAME math, so mixed GPU/CPU fleets stay bit-identical.
+  async function vmlpBlock(Xf, W1f, W2f, d, L, gpuMlp, audit) {
+    const x = quantizeRows(Xf, d.m, d.k);
+    const w1 = quantizeCols(W1f, d.k, d.h);
+    const w2 = quantizeCols(W2f, d.h, d.n);
+    if (gpuMlp) {
+      const r = await gpuMlp(x.q, w1.q, w2.q, x.s, w1.s, w2.s, d);
+      if (audit && audit.due()) {
+        // audit BOTH live GEMMs: gemm1 against the units directly; gemm2 by
+        // reconstructing its exact operand through the proven quantize mirror
+        const bad1 = auditTile(x.q, w1.q, x.s, w1.s, { m: d.m, k: d.k, n: d.h, relu: true }, r.h1, L, audit.cells);
+        if (bad1) { audit.fail("mlp gemm1: " + bad1); return r; }
+        const sc = scalesFromAbsMax(rowAbsMax(r.h1, d.m, d.h));
+        const hq = quantizeRowsInv(r.h1, d.m, d.h, sc.inv);
+        const bad2 = auditTile(hq, w2.q, sc.scale, w2.s, { m: d.m, k: d.h, n: d.n }, r.out, L, audit.cells);
+        if (bad2) audit.fail("mlp gemm2: " + bad2);
+      }
+      return r;
+    }
+    const h1 = bgemmJS(x.q, w1.q, x.s, w1.s, { m: d.m, k: d.k, n: d.h, batch: 1, relu: true }, L);
+    const sc = scalesFromAbsMax(rowAbsMax(h1, d.m, d.h));
+    const hq = quantizeRowsInv(h1, d.m, d.h, sc.inv);
+    const out = bgemmJS(hq, w2.q, sc.scale, w2.s, { m: d.m, k: d.h, n: d.n, batch: 1 }, L);
+    return { h1, out };
+  }
+
   // ---- 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
@@ -307,7 +383,8 @@
   }
 
   const api = { quantize, quantize2, quantizeRows, quantizeCols, quantizeHeadCols, lutMatmulJS, lutMatmul3JS, lutMatmul3,
-                bgemmJS, vgemmBlock, auditTile, epi, attScoresJS, attCtxJS, linearFwd, forward, backward, splitApply };
+                bgemmJS, vgemmBlock, auditTile, epi, attScoresJS, attCtxJS, linearFwd, forward, backward, splitApply,
+                rowAbsMax, scalesFromAbsMax, quantizeRowsInv, vmlpBlock };
   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,6 +35,57 @@
       C[row * n + col] = s;
     }`;
 
+  // ---- B2B MLP chain kernels (CUTLASS ex. 13 + 23) ---------------------------
+  // ROWMAX: per-row |max| of a GEMM's f32 output, fused into the same command
+  // encoder (ex. 23 epilogue reduction). Non-negative f32 bit patterns order
+  // like u32, so atomicMax on bitcast(abs(v)) computes an EXACT max, in any
+  // execution order, on any hardware — nothing here can round.
+  const WGSL_ROWMAX = `
+    @group(0) @binding(0) var<storage, read> O : array<f32>;
+    @group(0) @binding(1) var<storage, read_write> MX : array<atomic<u32>>;
+    @group(0) @binding(2) var<uniform> dims : vec4<u32>;         // m, n, _, _
+    @compute @workgroup_size(8, 8, 1)
+    fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
+      let m = dims.x; let n = dims.y;
+      let row = gid.x; let col = gid.y;
+      if (row >= m || col >= n) { return; }
+      atomicMax(&MX[row], bitcast<u32>(abs(O[row * n + col])));
+    }`;
+  // QUANT: h1 (f32, still on the GPU) -> int8 by MULTIPLY with a JS-computed
+  // inverse scale. floor(f32(x*inv)+0.5) uses only ops WGSL guarantees exact
+  // or correctly rounded (mul, add, floor, clamp) — division is 2.5 ULP and
+  // never runs on the GPU. Bit-identical to Verified.quantizeRowsInv, and
+  // exact-gated against it at init. pack=true emits 4 bytes per u32 for the
+  // DP4A kernel; pack=false emits one i32 per element for the LUT kernel.
+  const WGSL_QUANT = (pack) => `
+    @group(0) @binding(0) var<storage, read> H : array<f32>;
+    @group(0) @binding(1) var<storage, read> inv : array<f32>;   // per row
+    @group(0) @binding(2) var<storage, read_write> Q : array<${pack ? "u32" : "i32"}>;
+    @group(0) @binding(3) var<uniform> dims : vec4<u32>;         // m, k, kw, _
+    @compute @workgroup_size(64)
+    fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
+      let m = dims.x; let k = dims.y; let kw = dims.z;
+      let idx = gid.x;
+      ${pack ? `
+      if (idx >= m * kw) { return; }
+      let row = idx / kw;
+      var acc : u32 = 0u;
+      for (var b = 0u; b < 4u; b = b + 1u) {
+        let c = (idx % kw) * 4u + b;
+        var q : i32 = 0;
+        if (c < k) {
+          let v = clamp(floor(H[row * k + c] * inv[row] + 0.5), -128.0, 127.0);
+          q = i32(v);
+        }
+        acc = acc | ((u32(q) & 255u) << (8u * b));
+      }
+      Q[idx] = acc;` : `
+      if (idx >= m * k) { return; }
+      let row = idx / k;
+      let v = clamp(floor(H[idx] * inv[row] + 0.5), -128.0, 127.0);
+      Q[idx] = i32(v);`}
+    }`;
+
   // 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
@@ -252,6 +303,11 @@
       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"]);
+      const rowmaxLayout = mkLayout(["r", "rw", "u"]);
+      const quantLayout = mkLayout(["r", "r", "rw", "u"]);
+      const rowmaxPipe = mkPipeL(WGSL_ROWMAX, rowmaxLayout);
+      const quantI32Pipe = mkPipeL(WGSL_QUANT(false), quantLayout);
+      const quantPackPipe = mkPipeL(WGSL_QUANT(true), quantLayout);
       // 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);
@@ -342,9 +398,34 @@
       // 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; }
+
+      // B2B MLP chain gate (CUTLASS ex. 13+23): run the WHOLE chain — gemm1 +
+      // ReLU + on-GPU rowmax + on-GPU quantize + gemm2 — against the pure-JS
+      // mirror chain, exact `!==` on both h1 and the final output. Sweeps
+      // ragged shapes; h not a multiple of 4 exercises the pack-tail padding.
+      async function gateMlp(mlpFn) {
+        for (const d0 of [{ m: 6, k: 8, h: 12, n: 5 }, { m: 5, k: 16, h: 6, n: 3 },
+                          { m: 17, k: 33, h: 10, n: 9 }, { m: 32, k: 64, h: 128, n: 32 }]) {
+          const rnd = (len) => Float32Array.from({ length: len }, () => Math.random() * 2 - 1);
+          const Xf = rnd(d0.m * d0.k), W1 = rnd(d0.k * d0.h), W2 = rnd(d0.h * d0.n);
+          const hw = await root.Verified.vmlpBlock(Xf, W1, W2, d0, L, mlpFn, null);
+          const ref = await root.Verified.vmlpBlock(Xf, W1, W2, d0, L, null, null);
+          const shape = `${d0.m}x${d0.k}x${d0.h}x${d0.n}`;
+          for (let i = 0; i < ref.h1.length; i++)
+            if (hw.h1[i] !== ref.h1[i]) return `h1 mismatch @${i} (${shape}): ${hw.h1[i]} vs ${ref.h1[i]}`;
+          for (let i = 0; i < ref.out.length; i++)
+            if (hw.out[i] !== ref.out[i]) return `out mismatch @${i} (${shape}): ${hw.out[i]} vs ${ref.out[i]}`;
+        }
+        return null;
+      }
+      const lutMlpEnv = { dp4: false, gemm: bgLutPipe, rowmax: rowmaxPipe, quant: quantI32Pipe, lutBuf };
+      let mlpLut = (xq, w1q, w2q, xs, w1s, w2s, d) => gpuMlpChain(device, lutMlpEnv, xq, w1q, w2q, xs, w1s, w2s, d);
+      const mlpLutBad = await gateMlp(mlpLut);
+      if (mlpLutBad) { console.warn("B2B MLP chain (LUT) failed verification — MLP stays on the CPU mirror chain:", mlpLutBad); mlpLut = null; }
+
       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: bgLut, att, fgemm };
+                       bgemm: bgLut, att, fgemm, mlp: mlpLut };
 
       // DP4A pipeline — only if the WGSL feature exists AND its batched kernel
       // reproduces the verified units exactly across the shape sweep
@@ -357,8 +438,12 @@
         console.warn("batched DP4A disagreed with the verified units — using LUT bgemm:", dp4Bad);
         return viaLUT;
       }
+      const dp4MlpEnv = { dp4: true, gemm: bgDp4Pipe, rowmax: rowmaxPipe, quant: quantPackPipe };
+      let mlpDp4 = (xq, w1q, w2q, xs, w1s, w2s, d) => gpuMlpChain(device, dp4MlpEnv, xq, w1q, w2q, xs, w1s, w2s, d);
+      const mlpDp4Bad = await gateMlp(mlpDp4);
+      if (mlpDp4Bad) { console.warn("B2B MLP chain (DP4A) failed verification — using the LUT chain:", mlpDp4Bad); mlpDp4 = mlpLut; }
       return { backend: "webgpu", label: `${gpuName} (DP4A int8 dot · exact-gated vs units)`,
-               bgemm: bg, att, fgemm };
+               bgemm: bg, att, fgemm, mlp: mlpDp4 };
     } catch (e) { console.warn("WebGPU init failed, CPU fallback:", e); return cpu; }
   }
 
@@ -515,6 +600,85 @@
     return r.out;
   }
 
+  // B2B MLP chain: gemm1 (ReLU fused) + rowmax in one encoder, a 4·m-byte
+  // absmax readback, then quantize + gemm2 in a second encoder. h1 comes back
+  // because the STE backward needs it, but it never goes UP again — gemm2's
+  // left operand is produced and consumed entirely on the GPU.
+  async function gpuMlpChain(device, env, xq, w1q, w2q, xs, w1s, w2s, d) {
+    const { m, k, h, n } = d, dp4 = !!env.dp4;
+    const SU = GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST;
+    let bufX, bufW1, bufW2, kw1, hw;
+    if (dp4) {
+      kw1 = Math.ceil(k / 4); hw = Math.ceil(h / 4);
+      bufX = up(device, packRows(xq, m, k, kw1), SU);
+      bufW1 = up(device, packRows(transposeI8(w1q, k, h), h, k, kw1), SU);
+      bufW2 = up(device, packRows(transposeI8(w2q, h, n), n, h, hw), SU);
+    } else {
+      kw1 = k; hw = h;
+      bufX = up(device, Int32Array.from(xq), SU);
+      bufW1 = up(device, Int32Array.from(w1q), SU);
+      bufW2 = up(device, Int32Array.from(w2q), SU);
+    }
+    const bufRs = up(device, xs, SU), bufCs1 = up(device, w1s, SU), bufCs2 = up(device, w2s, SU);
+    const bufH = mk(device, m * h * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
+    const bufMX = mk(device, m * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);  // zero-initialized
+    const UU = GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST;
+    const bufD1 = up(device, new Uint32Array([m, kw1, h, 1]), UU);       // flags=1: fused ReLU
+    const bufDM = up(device, new Uint32Array([m, h, 0, 0]), UU);
+    const gemmBind = (bufA, bufB, bufR, bufC, bufO, bufD) => device.createBindGroup({
+      layout: env.gemm.getBindGroupLayout(0),
+      entries: (dp4 ? [bufA, bufB, bufR, bufC, bufO, bufD]
+                    : [bufA, bufB, env.lutBuf, bufR, bufC, bufO, bufD])
+        .map((b, i) => ({ binding: i, resource: { buffer: b } })) });
+    const enc1 = device.createCommandEncoder();
+    let pass = enc1.beginComputePass();
+    pass.setPipeline(env.gemm);
+    pass.setBindGroup(0, gemmBind(bufX, bufW1, bufRs, bufCs1, bufH, bufD1));
+    pass.dispatchWorkgroups(Math.ceil(m / 8), Math.ceil(h / 8), 1); pass.end();
+    pass = enc1.beginComputePass();
+    pass.setPipeline(env.rowmax);
+    pass.setBindGroup(0, device.createBindGroup({ layout: env.rowmax.getBindGroupLayout(0), entries: [
+      { binding: 0, resource: { buffer: bufH } }, { binding: 1, resource: { buffer: bufMX } },
+      { binding: 2, resource: { buffer: bufDM } } ] }));
+    pass.dispatchWorkgroups(Math.ceil(m / 8), Math.ceil(h / 8), 1); pass.end();
+    const readH = mk(device, m * h * 4, GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ);
+    const readM = mk(device, m * 4, GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ);
+    enc1.copyBufferToBuffer(bufH, 0, readH, 0, m * h * 4);
+    enc1.copyBufferToBuffer(bufMX, 0, readM, 0, m * 4);
+    device.queue.submit([enc1.finish()]);
+    await Promise.all([readH.mapAsync(GPUMapMode.READ), readM.mapAsync(GPUMapMode.READ)]);
+    const h1 = new Float32Array(readH.getMappedRange().slice(0)); readH.unmap();
+    // the atomicMax'ed u32 bit patterns ARE the f32 |max| values
+    const mx = new Float32Array(readM.getMappedRange().slice(0)); readM.unmap();
+    // scale derivation in JS f64 — exactly rounded, identical on every device
+    // (WGSL division is 2.5 ULP, which is why it never runs on the GPU)
+    const sc = root.Verified.scalesFromAbsMax(mx);
+    const bufInv = up(device, sc.inv, SU), bufHs = up(device, sc.scale, SU);
+    const bufQ = mk(device, m * hw * 4, GPUBufferUsage.STORAGE);
+    const bufDQ = up(device, new Uint32Array([m, h, hw, 0]), UU);
+    const bufD2 = up(device, new Uint32Array([m, hw, n, 0]), UU);
+    const bufO = mk(device, m * n * 4, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
+    const enc2 = device.createCommandEncoder();
+    pass = enc2.beginComputePass();
+    pass.setPipeline(env.quant);
+    pass.setBindGroup(0, device.createBindGroup({ layout: env.quant.getBindGroupLayout(0), entries: [
+      { binding: 0, resource: { buffer: bufH } }, { binding: 1, resource: { buffer: bufInv } },
+      { binding: 2, resource: { buffer: bufQ } }, { binding: 3, resource: { buffer: bufDQ } } ] }));
+    pass.dispatchWorkgroups(Math.ceil((m * (dp4 ? hw : h)) / 64)); pass.end();
+    pass = enc2.beginComputePass();
+    pass.setPipeline(env.gemm);
+    pass.setBindGroup(0, gemmBind(bufQ, bufW2, bufHs, bufCs2, bufO, bufD2));
+    pass.dispatchWorkgroups(Math.ceil(m / 8), Math.ceil(n / 8), 1); pass.end();
+    const readO = mk(device, m * n * 4, GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ);
+    enc2.copyBufferToBuffer(bufO, 0, readO, 0, m * n * 4);
+    device.queue.submit([enc2.finish()]);
+    await readO.mapAsync(GPUMapMode.READ);
+    const out = new Float32Array(readO.getMappedRange().slice(0)); readO.unmap();
+    [bufX, bufW1, bufW2, bufRs, bufCs1, bufCs2, bufH, bufMX, bufD1, bufDM,
+     bufInv, bufHs, bufQ, bufDQ, bufD2, bufO, readH, readM, readO].forEach(b => b.destroy());
+    return { h1, 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);

web/test_b2b.js

@@ -0,0 +1,107 @@
+// B2B MLP chain (CUTLASS ex. 13 + 23): tests for the respecced quantize and
+// the chained forward.
+//
+// What must hold EXACTLY:
+//   - the per-row absmax + scale derivation matches quantizeRows' internals
+//     bit for bit (ex. 23's reduction changes where max runs, not what it is)
+//   - the chain is deterministic and structurally equal to its parts
+//   - gemm1 (hence h1, hence the ReLU mask) is byte-identical to the old path
+// What changes BY DESIGN (and must stay bounded):
+//   - quantizeRowsInv rounds via floor(f32(x*inv)+0.5) instead of
+//     round(x_f64/scale): every int8 moves by AT MOST one step, only on
+//     values that sit within float noise of a rounding boundary
+// The GPU side of the same claims is enforced at init by gateMlp in webgpu.js
+// (exact !== against the mirror chain, ragged shapes, pack-tail padding).
+const fs = require("fs");
+const path = require("path");
+const p = (f) => path.join(__dirname, "public", f);
+const V = require("./public/verified_core.js");
+const T = require("./public/traincore.js");
+const OLD = require("./public/_transformer_prefusion.js");
+const NEW = require("./public/transformer.js");
+const L = { mul: new Int16Array(fs.readFileSync(p("mul_lut.bin")).buffer.slice(0)) };
+
+const rnd = (len) => Float32Array.from({ length: len }, () => Math.random() * 4 - 2);
+let pass = true;
+const ok = (c, msg) => { console.log(`${c ? "  ok  " : "  FAIL"}  ${msg}`); if (!c) pass = false; };
+
+(async () => {
+  // 1) ex. 23 reduction: absmax + scales identical to quantizeRows' own
+  {
+    let bad = 0, n = 0;
+    for (let t = 0; t < 200; t++) {
+      const rows = 1 + (Math.random() * 40 | 0), cols = 1 + (Math.random() * 90 | 0);
+      const X = rnd(rows * cols);
+      if (Math.random() < 0.1) for (let c = 0; c < cols; c++) X[c] = 0;   // zero row -> 1e-8 floor
+      const sOld = V.quantizeRows(X, rows, cols).s;
+      const sNew = V.scalesFromAbsMax(V.rowAbsMax(X, rows, cols)).scale;
+      for (let i = 0; i < rows; i++) { n++; if (sOld[i] !== sNew[i]) bad++; }
+    }
+    ok(bad === 0, `scales from the fused absmax are bit-identical to quantizeRows (${n} rows incl. zero rows)`);
+  }
+  // 2) the respec moves an int8 by at most ONE step
+  {
+    let maxd = 0, diff = 0, n = 0;
+    for (let t = 0; t < 200; t++) {
+      const rows = 1 + (Math.random() * 30 | 0), cols = 1 + (Math.random() * 70 | 0);
+      const X = rnd(rows * cols);
+      const old = V.quantizeRows(X, rows, cols);
+      const inv = V.scalesFromAbsMax(V.rowAbsMax(X, rows, cols)).inv;
+      const q2 = V.quantizeRowsInv(X, rows, cols, inv);
+      for (let i = 0; i < q2.length; i++) {
+        n++;
+        const d = Math.abs(q2[i] - old.q[i]);
+        if (d) diff++;
+        if (d > maxd) maxd = d;
+      }
+    }
+    ok(maxd <= 1, `respecced quantize differs by at most 1 step (max ${maxd}; ${diff}/${n} = ${(100 * diff / n).toFixed(3)}% of values moved)`);
+  }
+  // 3) chain == its parts, and gemm1 is byte-identical to the un-chained GEMM
+  {
+    const d = { m: 17, k: 32, h: 20, n: 9 };
+    const X = rnd(d.m * d.k), W1 = rnd(d.k * d.h), W2 = rnd(d.h * d.n);
+    const r = await V.vmlpBlock(X, W1, W2, d, L, null, null);
+    const x = V.quantizeRows(X, d.m, d.k), w1 = V.quantizeCols(W1, d.k, d.h), w2 = V.quantizeCols(W2, d.h, d.n);
+    const h1ref = V.bgemmJS(x.q, w1.q, x.s, w1.s, { m: d.m, k: d.k, n: d.h, batch: 1, relu: true }, L);
+    let same = true;
+    for (let i = 0; i < h1ref.length; i++) if (r.h1[i] !== h1ref[i]) { same = false; break; }
+    ok(same, "chain gemm1 (hence h1 and the ReLU mask) is byte-identical to the un-chained GEMM");
+    const sc = V.scalesFromAbsMax(V.rowAbsMax(r.h1, d.m, d.h));
+    const outref = V.bgemmJS(V.quantizeRowsInv(r.h1, d.m, d.h, sc.inv), w2.q, sc.scale, w2.s, { m: d.m, k: d.h, n: d.n, batch: 1 }, L);
+    let same2 = true;
+    for (let i = 0; i < outref.length; i++) if (r.out[i] !== outref[i]) { same2 = false; break; }
+    ok(same2, "chain output equals the manual composition of its stages");
+    const r2 = await V.vmlpBlock(X, W1, W2, d, L, null, null);
+    ok(r.out.every((v, i) => v === r2.out[i]), "chain is deterministic (two runs bitwise equal)");
+  }
+  // 4) the transformer trains through the chain, and the spec change does not
+  //    hurt convergence: same seeds, old forward vs chained forward
+  {
+    const cfg = { c: 32, t: 32, b: 8, layers: 2, heads: 2, steps: 40, lr: 0.01 };
+    const run = async (X) => {
+      const m = X.init(cfg, L, null);
+      const opt = T.makeAdam(m.nParams, { lr: cfg.lr });
+      let seed = 31337;
+      const orig = Math.random;
+      Math.random = () => { seed = (Math.imul(seed, 1103515245) + 12345) & 0x7fffffff; return seed / 0x7fffffff; };
+      let first = 0, last = 0;
+      for (let s = 0; s < cfg.steps; s++) {
+        const r = await X.trainStep(m);
+        if (s === 0) first = r.loss;
+        last = r.loss;
+        X.applyUpdate(m, opt.step(r.grad));
+      }
+      Math.random = orig;
+      return { first, last };
+    };
+    const a = await run(OLD), b = await run(NEW);
+    ok(a.first === b.first, `step-0 loss identical before any quantize divergence compounds (${a.first.toFixed(6)})`);
+    ok(b.last < b.first * 0.75, `chained transformer converges (${b.first.toFixed(4)} -> ${b.last.toFixed(4)})`);
+    const rel = Math.abs(a.last - b.last) / a.last;
+    ok(rel < 0.10, `convergence unchanged by the respec: old ${a.last.toFixed(4)} vs new ${b.last.toFixed(4)} (${(100 * rel).toFixed(1)}% apart)`);
+    console.log("  note   the runs are NOT bit-identical past step 0 — the quantize respec is a real (bounded) spec change, which is why old and new builds must not co-train");
+  }
+  console.log(pass ? "\nB2B MLP CHAIN TEST PASSED" : "\nB2B MLP CHAIN TEST FAILED");
+  process.exit(pass ? 0 : 1);
+})();