discrete-GPU validation ?bench=discrete: bandwidth + live + spec-flip in one page

index.html

@@ -220,6 +220,79 @@ async function runPerfBench() {
   log.innerHTML = tbl; st.textContent = "live decode profile · done"; console.log("[perfbench]", runs);
 }
 
+// ── DISCRETE-GPU VALIDATION (?bench=discrete) ── the whole thesis in one page, on whatever GPU opens it:
+// (1) real VRAM bandwidth + decode roofline, (2) live BitNet tok/s warmed, (3) does spec-decode FLIP from the
+// iGPU loss to a win once bandwidth-bound? Confirms the discrete GPU (bandwidth >350 GB/s) is actually in use.
+const BW_FILL = `@group(0) @binding(0) var<storage,read_write> d: array<u32>;
+@group(0) @binding(1) var<uniform> P: vec4<u32>;
+@compute @workgroup_size(256) fn main(@builtin(global_invocation_id) g:vec3<u32>){ let n=P.x; var i=g.x; loop{ if(i>=n){break;} d[i]=(i*2654435761u+1u); i=i+P.y; } }`;
+const BW_READ = `@group(0) @binding(0) var<storage,read> d: array<vec4<u32>>;
+@group(0) @binding(1) var<storage,read_write> sink: array<u32>;
+@group(0) @binding(2) var<uniform> P: vec4<u32>;
+@compute @workgroup_size(256) fn main(@builtin(global_invocation_id) g:vec3<u32>){ let n=P.x; let stride=P.y; var acc=vec4<u32>(0u); var i=g.x; loop{ if(i>=n){break;} acc=acc^d[i]; i=i+stride; } sink[g.x]=acc.x^acc.y^acc.z^acc.w; }`;
+async function measureVramBW(dev) {
+  const L = dev.limits;
+  const bytes = Math.floor(Math.min(L.maxStorageBufferBindingSize, L.maxBufferSize, 512 * 1024 * 1024) / 16) * 16, nVec = bytes / 16;
+  const buf = dev.createBuffer({ size: bytes, usage: GPUBufferUsage.STORAGE });
+  const wg = Math.min(L.maxComputeWorkgroupsPerDimension, 65535), TOTAL = wg * 256;
+  const sink = dev.createBuffer({ size: TOTAL * 4, usage: GPUBufferUsage.STORAGE });
+  const P = dev.createBuffer({ size: 16, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST });
+  dev.queue.writeBuffer(P, 0, new Uint32Array([bytes / 4, TOTAL, 0, 0]));
+  const fp = dev.createComputePipeline({ layout: "auto", compute: { module: dev.createShaderModule({ code: BW_FILL }), entryPoint: "main" } });
+  const fbg = dev.createBindGroup({ layout: fp.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: buf } }, { binding: 1, resource: { buffer: P } }] });
+  { const e = dev.createCommandEncoder(); const p = e.beginComputePass(); p.setPipeline(fp); p.setBindGroup(0, fbg); p.dispatchWorkgroups(wg); p.end(); dev.queue.submit([e.finish()]); await dev.queue.onSubmittedWorkDone(); }
+  dev.queue.writeBuffer(P, 0, new Uint32Array([nVec, TOTAL, 0, 0]));
+  const rp = dev.createComputePipeline({ layout: "auto", compute: { module: dev.createShaderModule({ code: BW_READ }), entryPoint: "main" } });
+  const rbg = dev.createBindGroup({ layout: rp.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: buf } }, { binding: 1, resource: { buffer: sink } }, { binding: 2, resource: { buffer: P } }] });
+  const run = async (passes) => { const e = dev.createCommandEncoder(); for (let k = 0; k < passes; k++) { const p = e.beginComputePass(); p.setPipeline(rp); p.setBindGroup(0, rbg); p.dispatchWorkgroups(wg); p.end(); } const t0 = performance.now(); dev.queue.submit([e.finish()]); await dev.queue.onSubmittedWorkDone(); return performance.now() - t0; };
+  await run(4); let best = 1e9; for (let k = 0; k < 5; k++) best = Math.min(best, await run(32));
+  buf.destroy(); sink.destroy();
+  return (bytes * 32 / 1073741824) / (best / 1000);
+}
+async function runDiscreteBench() {
+  log.innerHTML = ""; input.disabled = send.disabled = true;
+  const rep = m.rep ?? 1.3, gpu = engine._gpu, dev = gpu && gpu._dev && gpu._dev();
+  let adapterStr = "unknown";
+  try { const a = await navigator.gpu.requestAdapter({ powerPreference: "high-performance" }); const inf = (a && a.info) || {}; adapterStr = ((inf.vendor || "") + " " + (inf.architecture || "") + " " + (inf.device || "")).trim() || "unknown"; } catch {}
+  st.textContent = "measuring VRAM bandwidth…";
+  let gbps = 0; try { gbps = await measureVramBW(dev); } catch (e) { console.warn("bw", e); }
+  const roofTok = gbps / 0.69, discrete = gbps > 350;
+  st.textContent = "warming + measuring live decode…";
+  globalThis.__spec = false;
+  engine.reset(); await engine.generate(engine.tokenize(engine.frameTurn("Write one sentence about the sea.", false)), { maxNew: 48, repPenalty: rep });
+  let live = 0;
+  for (let i = 0; i < 3; i++) { engine.reset(); const r = await engine.generate(engine.tokenize(engine.frameTurn("Write a detailed paragraph about ocean currents.", false)), { maxNew: 128, repPenalty: rep }); live = Math.max(live, (r.stats && r.stats.tokps) || 0); }
+  st.textContent = "measuring spec-decode flip…";
+  const sIds = engine.tokenize(engine.frameTurn("Passage: \"The quick brown fox jumps over the lazy dog near the river bank at dawn.\" Repeat that passage back to me word for word.", false));
+  globalThis.__spec = false; engine.reset(); let t0 = performance.now(); const rb = await engine.generate(sIds.slice(), { maxNew: 192, repPenalty: rep }); const bTok = rb.outIds.length / ((performance.now() - t0) / 1000);
+  let sTok = 0, perVerify = 0, accept = 0, same = true, hasSpec = !!engine.specAvailable;
+  if (hasSpec) {
+    globalThis.__spec = true; engine.reset(); t0 = performance.now(); const rs = await engine.generate(sIds.slice(), { maxNew: 192, repPenalty: rep }); sTok = rs.outIds.length / ((performance.now() - t0) / 1000); globalThis.__spec = false;
+    const sp = (rs.stats && rs.stats.spec) || null;
+    perVerify = sp && sp.windows ? 1 + sp.accepted / sp.windows : 0; accept = sp && sp.drafted ? 100 * sp.accepted / sp.drafted : 0;
+    same = rb.outIds.length === rs.outIds.length && rb.outIds.every((x, i) => x === rs.outIds[i]);
+  }
+  const flip = hasSpec && sTok > bTok * 1.1, spRatio = bTok ? sTok / bTok : 0;
+  const row = (k, v) => `<tr style="border-top:1px solid var(--line)"><td style="padding:7px 8px;color:var(--dim)">${k}</td><td style="padding:7px 8px;font-weight:600">${v}</td></tr>`;
+  log.innerHTML = `<div style="font-family:ui-monospace,monospace;font-size:13px;max-width:820px;margin:0 auto;padding:8px">
+    <div style="font-size:18px;font-weight:700;margin-bottom:4px">Discrete-GPU validation</div>
+    <div style="color:var(--dim);margin-bottom:12px">adapter: ${adapterStr} · BitNet-2B</div>
+    <table style="width:100%;border-collapse:collapse">
+      ${row("GPU in use", `<span style="color:${discrete ? "#48c26c" : "#e0a94a"}">${discrete ? "DISCRETE ✓" : "integrated (bandwidth "+gbps.toFixed(0)+" GB/s — not a discrete GPU)"}</span>`)}
+      ${row("VRAM bandwidth", `${gbps.toFixed(0)} GB/s`)}
+      ${row("Decode roofline", `${roofTok.toFixed(0)} tok/s`)}
+      ${row("Live decode (warmed)", `${live.toFixed(0)} tok/s · ${(100*live/roofTok).toFixed(0)}% of roofline`)}
+      ${row("Spec-decode (retrieval)", hasSpec ? `${sTok.toFixed(0)} vs ${bTok.toFixed(0)} baseline · <b style="color:${flip ? "#48c26c" : "#e0a94a"}">${spRatio.toFixed(2)}×</b> · ${perVerify.toFixed(2)} tok/verify · ${accept.toFixed(0)}% accept · ${same ? "byte-exact ✓" : "DIVERGED ✗"}` : "unavailable")}
+    </table>
+    <div style="margin-top:14px;font-size:15px;font-weight:600;color:${flip ? "#48c26c" : "#e0a94a"}">${!discrete
+      ? "⚠ This is still an integrated GPU (bandwidth ≤350 GB/s). Open on a machine with a discrete GPU to validate the >1000 path — the browser may be picking the iGPU for power saving."
+      : flip
+        ? `✓ CONFIRMED: on discrete silicon spec-decode FLIPPED to a ${spRatio.toFixed(2)}× win (byte-exact), and the roofline is ${roofTok.toFixed(0)} tok/s. The code we built delivers the high-throughput path unchanged on this hardware.`
+        : `Spec-decode is ${spRatio.toFixed(2)}× here — not yet a clear win. Bandwidth ${gbps.toFixed(0)} GB/s (roofline ${roofTok.toFixed(0)}); the forward may still be compute-bound on this GPU. Send me the numbers.`}</div>
+  </div>`;
+  st.textContent = "discrete validation · done"; console.log("[discrete]", { adapterStr, gbps, roofTok, live, bTok, sTok, perVerify, accept, same });
+}
+
 // ── PER-PASS GPU TRACE (?bench=trace) ── where does a token's ~40ms go? Runs a profiled forward through the
 // step() path (window.__profile → timestamp-query'd ns per pass) and dumps the breakdown sorted by cost, so the
 // 83% non-weight overhead is named exactly (attention / argmax / lm_head / norms / dispatch count) — no guessing.
@@ -265,6 +338,7 @@ try {
   if (params.get("bench") === "spec") { globalThis.__spec = false; await runSpecBench(); }
   else if (params.get("bench") === "perf") { await runPerfBench(); }
   else if (params.get("bench") === "trace") { await runTraceBench(); }
+  else if (params.get("bench") === "discrete") { globalThis.__spec = false; await runDiscreteBench(); }
   else if (pending) { const w = [...log.querySelectorAll(".a")].reverse().find((x) => x.dataset.pending); if (w) w.remove(); const p = pending; pending = null; generate(p, true); }
   else await proactiveGreeting();
 } catch (e) { st.textContent = "⚠ " + e.message; bubble("a", "Could not start: " + e.message); }