Spaces:

borkiss
/

qwen35-2b-webgpu-demo

Running

App Files Files Community

borkiss commited on Jul 6

Commit

51e8f9a

verified ·

1 Parent(s): fcb061c

Upload folder using huggingface_hub

Browse files

Files changed (12) hide show

README.md +8 -4
index.html +79 -17
web/bench-main.js +108 -0
web/bench.html +31 -0
web/config.js +22 -0
web/gpu.js +91 -0
web/index.html +81 -0
web/kernels.js +1361 -0
web/main.js +171 -0
web/model.js +440 -0
web/npy.js +16 -0
web/tokenizer.js +122 -0

README.md CHANGED Viewed

@@ -1,10 +1,14 @@
 ---
-title: Qwen35 2b Webgpu Demo
-emoji: 🏆
 colorFrom: purple
-colorTo: gray
 sdk: static
 pinned: false
 ---
-Check out the configuration reference at https://huggingface.co/docs/hub/spaces-config-reference

 ---
+title: Qwen3.5-2B WebGPU (custom WGSL kernels)
+emoji: ⚡
 colorFrom: purple
+colorTo: blue
 sdk: static
 pinned: false
+license: mit
 ---
+Qwen3.5-2B (hybrid gated DeltaNet + gated attention) running in the browser
+on hand-written WGSL kernels only — the first DeltaNet hybrid on WebGPU.
+Weights stream from [borkiss/qwen35-2b-webgpu](https://huggingface.co/borkiss/qwen35-2b-webgpu).
+Needs WebGPU with subgroups (Chrome/Edge 125+; shader-f16 optional).

index.html CHANGED Viewed

@@ -1,19 +1,81 @@
 <!doctype html>
-<html>
-	<head>
-		<meta charset="utf-8" />
-		<meta name="viewport" content="width=device-width" />
-		<title>My static Space</title>
-		<link rel="stylesheet" href="style.css" />
-	</head>
-	<body>
-		<div class="card">
-			<h1>Welcome to your static Space!</h1>
-			<p>You can modify this app directly by editing <i>index.html</i> in the Files and versions tab.</p>
-			<p>
-				Also don't forget to check the
-				<a href="https://huggingface.co/docs/hub/spaces" target="_blank">Spaces documentation</a>.
-			</p>
-		</div>
-	</body>
 </html>

 <!doctype html>
+<html lang="en">
+<meta charset="utf-8" />
+<meta name="viewport" content="width=device-width, initial-scale=1" />
+<title>Qwen3.5-2B · WebGPU</title>
+<style>
+  :root {
+    --bg: #0b0d10; --panel: #14171c; --line: #232830; --fg: #e8ebef;
+    --dim: #8b93a1; --accent: #7c5cff; --ok: #3ecf8e; --err: #ff6b6b;
+  }
+  * { box-sizing: border-box; }
+  body {
+    margin: 0; background: var(--bg); color: var(--fg);
+    font: 15px/1.55 ui-sans-serif, system-ui, -apple-system, "Segoe UI", sans-serif;
+    height: 100dvh; display: flex; flex-direction: column;
+  }
+  header {
+    display: flex; align-items: baseline; gap: 12px; padding: 12px 18px;
+    border-bottom: 1px solid var(--line); flex: none;
+  }
+  .wordmark { font-weight: 700; letter-spacing: .2px; }
+  .sub { color: var(--dim); font-size: 12.5px; }
+  #gpuname { margin-left: auto; color: var(--dim); font-size: 12.5px; }
+  #toksbox { color: var(--dim); font-size: 12.5px; }
+  #toks { color: var(--ok); font-variant-numeric: tabular-nums; }
+  #chat { flex: 1; overflow-y: auto; padding: 18px; }
+  .msg { max-width: 780px; margin: 0 auto 16px; }
+  .msg .who { font-size: 11px; text-transform: uppercase; letter-spacing: .8px; color: var(--dim); margin-bottom: 4px; }
+  .msg.user .who { color: var(--accent); }
+  .msg .body { white-space: pre-wrap; word-break: break-word; }
+  .stats { font-size: 11.5px; color: var(--dim); margin-top: 6px; }
+  .caret { display: inline-block; width: 7px; height: 15px; background: var(--accent); animation: blink 1s steps(1) infinite; vertical-align: -2px; }
+  @keyframes blink { 50% { opacity: 0; } }
+  form { flex: none; display: flex; gap: 10px; padding: 14px 18px 18px; border-top: 1px solid var(--line); }
+  #input {
+    flex: 1; max-width: 780px; margin-left: auto; resize: none; max-height: 160px;
+    background: var(--panel); color: var(--fg); border: 1px solid var(--line);
+    border-radius: 10px; padding: 10px 12px; font: inherit; outline: none;
+  }
+  #input:focus { border-color: var(--accent); }
+  #send {
+    margin-right: auto; background: var(--accent); border: 0; color: #fff;
+    border-radius: 10px; padding: 0 18px; font: inherit; font-weight: 600; cursor: pointer;
+  }
+  #send:disabled, #input:disabled { opacity: .45; cursor: default; }
+  #boot {
+    position: fixed; inset: 0; background: var(--bg); display: flex; flex-direction: column;
+    align-items: center; justify-content: center; gap: 14px; transition: opacity .4s; z-index: 10;
+  }
+  #boot.done { opacity: 0; pointer-events: none; }
+  #boottitle { font-size: 22px; font-weight: 700; }
+  #bootsub { color: var(--dim); font-size: 13px; max-width: 420px; text-align: center; }
+  #bar { width: 300px; height: 6px; background: var(--line); border-radius: 3px; overflow: hidden; }
+  #barfill { width: 100%; height: 100%; background: var(--accent); transform: scaleX(0); transform-origin: left; transition: transform .2s; }
+  #log { font: 12px/1.7 ui-monospace, monospace; color: var(--dim); width: 420px; }
+  #log .ok { color: var(--ok); }
+  #log .err { color: var(--err); }
+</style>
+<body>
+  <header>
+    <span class="wordmark">Qwen3.5-2B</span>
+    <span class="sub">hybrid DeltaNet + attention · hand-written WGSL · int4</span>
+    <span id="toksbox"><span id="toks">–</span> tok/s</span>
+    <span id="gpuname"></span>
+  </header>
+  <div id="chat"><div id="msgs"></div></div>
+  <form id="form">
+    <textarea id="input" rows="1" placeholder="Message… (Enter to send)" disabled></textarea>
+    <button id="send" type="submit" disabled>Send</button>
+  </form>
+  <div id="boot">
+    <div id="boottitle">Qwen3.5-2B</div>
+    <div id="bootsub">First hybrid gated-DeltaNet LLM in the browser — pure WebGPU, hand-written WGSL kernels.
+      Downloading <span id="dlmb">~1.1 GB</span> of int4 weights <span id="dlpct"></span></div>
+    <div id="bar"><div id="barfill"></div></div>
+    <div id="dlnote" class="sub"></div>
+    <div id="log"></div>
+  </div>
+  <script type="module" src="/web/main.js"></script>
+</body>
 </html>

web/bench-main.js ADDED Viewed

	@@ -0,0 +1,108 @@

+// Public benchmark page: runs the standard suite on the visitor's GPU and
+// produces a shareable .log (human-readable + embedded JSON).
+// Community format (as lfm2): window.__done = { ok, meta, results }.
+import { initDevice } from "/web/gpu.js";
+import { loadModel } from "/web/model.js";
+import { Tokenizer } from "/web/tokenizer.js";
+import { resolvePaths } from "/web/config.js";
+const $ = (id) => document.getElementById(id);
+const median = (a) => a.slice().sort((x, y) => x - y)[a.length >> 1];
+const lines = [];
+function log(text, cls = "") {
+  lines.push(text);
+  const el = document.createElement("div");
+  if (cls) el.className = cls;
+  el.textContent = text;
+  if ($("console").textContent === "idle — press Run bench") $("console").textContent = "";
+  $("console").appendChild(el);
+  $("console").scrollTop = $("console").scrollHeight;
+}
+async function bench() {
+  const R = { meta: {}, results: {} };
+  R.meta.date = new Date().toISOString();
+  R.meta.ua = navigator.userAgent;
+  R.meta.model = "Qwen3.5-2B";
+  R.meta.quant = new URLSearchParams(location.search).get("quant") ?? "q4";
+  const adapter = await navigator.gpu?.requestAdapter();
+  if (!adapter) throw new Error("WebGPU is not available in this browser");
+  R.meta.gpu = `${adapter.info.vendor} ${adapter.info.architecture} ${adapter.info.device ?? ""}`.trim();
+  R.meta.features = ["shader-f16", "subgroups", "timestamp-query"].filter((f) => adapter.features.has(f));
+  log(`gpu: ${R.meta.gpu}`, "hl");
+  log(`features: ${R.meta.features.join(", ")}`);
+  const device = await initDevice();
+  const paths = await resolvePaths();
+  R.meta.weights_source = paths.source;
+  log(`weights source: ${paths.source}`);
+  log(`loading Qwen3.5-2B ${R.meta.quant} weights…`);
+  let t0 = performance.now();
+  const model = await loadModel(device, { base: paths.weights, variant: R.meta.quant });
+  R.results.load_s = +((performance.now() - t0) / 1000).toFixed(1);
+  log(`loaded in ${R.results.load_s}s${model.weightsFromCache ? " (browser cache)" : ""}`, "ok");
+  const tok = await Tokenizer.load(paths.tokenizer);
+  log("warmup (pipeline compilation)…");
+  model.startGenerate(tok.encode("warmup run of the benchmark"));
+  await model.decodeChunk(4);
+  // prefill: 512 tokens x3
+  const P = 512;
+  const words = "the quick brown fox jumps over the lazy dog and keeps running through the quiet forest ".repeat(60);
+  let ids = tok.encode(words).slice(0, P);
+  const pre = [];
+  for (let r = 0; r < 3; r++) {
+    model.reset();
+    t0 = performance.now();
+    model.prefill(ids);
+    await device.queue.onSubmittedWorkDone();
+    pre.push(P / ((performance.now() - t0) / 1000));
+  }
+  R.results.prefill_tok_s = +median(pre).toFixed(1);
+  log(`prefill (${P} tok, median of 3): ${R.results.prefill_tok_s} tok/s`, "ok");
+  // decode: 128 tokens x5, GPU-resident chunks of 16
+  const D = 128;
+  const runs = [];
+  for (let r = 0; r < 5; r++) {
+    model.startGenerate(tok.applyChatTemplate([{ role: "user", content: "Tell me a long story about a lighthouse keeper." }]));
+    await device.queue.onSubmittedWorkDone();
+    t0 = performance.now();
+    for (let done = 1; done < D; done += 16) await model.decodeChunk(Math.min(16, D - done));
+    runs.push((D - 1) / ((performance.now() - t0) / 1000));
+    log(`  decode run ${r + 1}: ${runs[r].toFixed(1)} tok/s`);
+  }
+  R.results.decode_tok_s = +median(runs).toFixed(1);
+  R.results.decode_tok_s_best = +Math.max(...runs).toFixed(1);
+  log(`decode (${D} tok x5, chunk 16): median ${R.results.decode_tok_s}, best ${R.results.decode_tok_s_best} tok/s`, "ok");
+  log("");
+  log(JSON.stringify({ meta: R.meta, results: R.results }), "hl");
+  return R;
+}
+let lastR = null;
+$("run").addEventListener("click", async () => {
+  $("run").disabled = true;
+  try {
+    lastR = await bench();
+    $("save").disabled = false;
+    window.__done = { ok: true, meta: lastR.meta, results: lastR.results };
+  } catch (e) {
+    log(String(e.message ?? e), "err");
+    window.__done = { ok: false, error: String(e.message ?? e) };
+  }
+  $("run").disabled = false;
+});
+$("save").addEventListener("click", () => {
+  const blob = new Blob([lines.join("\n") + "\n"], { type: "text/plain" });
+  const a = document.createElement("a");
+  a.href = URL.createObjectURL(blob);
+  a.download = `qwen35-webgpu-bench-${Date.now()}.log`;
+  a.click();
+});
+if (new URLSearchParams(location.search).has("auto")) $("run").click();

web/bench.html ADDED Viewed

	@@ -0,0 +1,31 @@

+<!doctype html>
+<html lang="en">
+<meta charset="utf-8" />
+<meta name="viewport" content="width=device-width, initial-scale=1" />
+<title>Qwen3.5-2B WebGPU · bench</title>
+<style>
+  body { margin: 0; background: #0b0d10; color: #e8ebef; font: 14px/1.6 ui-monospace, monospace; }
+  main { max-width: 760px; margin: 40px auto; padding: 0 16px; }
+  h1 { font-size: 18px; }
+  .sub { color: #8b93a1; font-size: 12.5px; }
+  #console { background: #14171c; border: 1px solid #232830; border-radius: 10px; padding: 14px; min-height: 280px; max-height: 60vh; overflow-y: auto; white-space: pre-wrap; margin: 16px 0; }
+  #console .ok { color: #3ecf8e; }
+  #console .hl { color: #7c5cff; }
+  #console .err { color: #ff6b6b; }
+  button { background: #7c5cff; border: 0; color: #fff; border-radius: 8px; padding: 9px 16px; font: inherit; font-weight: 600; cursor: pointer; margin-right: 8px; }
+  button:disabled { opacity: .45; cursor: default; }
+</style>
+<body>
+  <main>
+    <h1>Qwen3.5-2B WebGPU benchmark</h1>
+    <div class="sub">Runs the standard suite on your GPU (downloads ~1.07 GB of int4 weights once, cached).
+      Share the .log — results land in the community table.</div>
+    <p>
+      <button id="run">Run bench</button>
+      <button id="save" disabled>Download .log</button>
+    </p>
+    <div id="console">idle — press Run bench</div>
+  </main>
+  <script type="module" src="/web/bench-main.js"></script>
+</body>
+</html>

web/config.js ADDED Viewed

	@@ -0,0 +1,22 @@

+// Where to fetch weights + tokenizer from. A local checkout (run.mjs dev
+// server, or an HF Space carrying the blobs) serves them same-origin;
+// otherwise fall back to the HF model repo — its resolve/ CDN sends
+// Access-Control-Allow-Origin: *, so any static host can pull cross-origin.
+const HF_REPO = "borkiss/qwen35-2b-webgpu";
+const HF = `https://huggingface.co/${HF_REPO}/resolve/main`;
+let cached = null;
+export async function resolvePaths() {
+  if (cached) return cached;
+  let local = false;
+  for (const m of ["manifest.q4.json", "manifest.f16.json"]) {
+    try {
+      const r = await fetch(`/web/weights/${m}`, { method: "HEAD" });
+      if (r.ok) { local = true; break; }
+    } catch {}
+  }
+  cached = local
+    ? { weights: "/web/weights/", tokenizer: "/web/weights/tokenizer.json", source: "local" }
+    : { weights: `${HF}/weights/`, tokenizer: `${HF}/weights/tokenizer.json`, source: "huggingface" };
+  return cached;
+}

web/gpu.js ADDED Viewed

	@@ -0,0 +1,91 @@

+// Thin WebGPU helpers shared by tests and the model runtime.
+import { setSubgroups, setF16 } from "./kernels.js";
+export async function initDevice() {
+  const adapter = await navigator.gpu?.requestAdapter({ powerPreference: "high-performance" })
+    ?? await navigator.gpu?.requestAdapter();
+  if (!adapter) throw new Error("WebGPU is not available — use Chrome/Edge 113+ or Safari 26+");
+  const params = typeof location !== "undefined" ? new URLSearchParams(location.search) : new URLSearchParams();
+  // shader-f16 is a fast path, not a requirement: without it (Dawn gates it
+  // on NVIDIA/Linux; Qualcomm mobiles lack it) f16 tensors are converted to
+  // f32 at load time and the KV cache lives in f32
+  const hasF16 = !params.has("nof16") && adapter.features.has("shader-f16");
+  setF16(hasF16);
+  const features = hasF16 ? ["shader-f16"] : [];
+  // subgroups are a fast path, not a requirement: without them (iOS Safari)
+  // the matvec kernels fall back to workgroup-shared reductions
+  const hasSubgroups = !params.has("nosg") && adapter.features.has("subgroups");
+  if (hasSubgroups) features.push("subgroups");
+  setSubgroups(hasSubgroups);
+  if (adapter.features.has("timestamp-query")) features.push("timestamp-query");
+  if (hasSubgroups && adapter.features.has("chromium-experimental-subgroup-matrix"))
+    features.push("chromium-experimental-subgroup-matrix");
+  const device = await adapter.requestDevice({
+    requiredFeatures: features,
+    requiredLimits: {
+      maxBufferSize: adapter.limits.maxBufferSize,
+      maxStorageBufferBindingSize: adapter.limits.maxStorageBufferBindingSize,
+      maxStorageBuffersPerShaderStage: adapter.limits.maxStorageBuffersPerShaderStage,
+      maxComputeWorkgroupStorageSize: adapter.limits.maxComputeWorkgroupStorageSize,
+    },
+  });
+  device.hasSubgroups = hasSubgroups;
+  device.hasF16 = hasF16;
+  device.lost.then((info) => console.error("GPU device lost:", info.reason, info.message));
+  return device;
+}
+export function createBuffer(device, sizeBytes, { storage = true, uniform = false } = {}) {
+  let usage = GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC;
+  if (storage) usage |= GPUBufferUsage.STORAGE;
+  if (uniform) usage |= GPUBufferUsage.UNIFORM;
+  return device.createBuffer({ size: Math.max(sizeBytes, 16), usage });
+}
+export function uploadF32(device, data) {
+  const buf = createBuffer(device, data.byteLength);
+  device.queue.writeBuffer(buf, 0, data);
+  return buf;
+}
+export function makePipeline(device, code, label = "") {
+  const module = device.createShaderModule({ code, label });
+  return device.createComputePipeline({ label, layout: "auto", compute: { module } });
+}
+export function bind(device, pipeline, entries) {
+  return device.createBindGroup({
+    layout: pipeline.getBindGroupLayout(0),
+    entries: Object.entries(entries).map(([binding, buffer]) => ({
+      binding: Number(binding),
+      resource: { buffer },
+    })),
+  });
+}
+export async function readback(device, buffer, sizeBytes, offset = 0) {
+  const rb = device.createBuffer({
+    size: sizeBytes,
+    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
+  });
+  const enc = device.createCommandEncoder();
+  enc.copyBufferToBuffer(buffer, offset, rb, 0, sizeBytes);
+  device.queue.submit([enc.finish()]);
+  await rb.mapAsync(GPUMapMode.READ);
+  const data = new Float32Array(rb.getMappedRange().slice(0));
+  rb.destroy();
+  return data;
+}
+// f16 helpers (Float16Array is available in Chrome 135+)
+export function toF16(f32arr) {
+  return new Float16Array(f32arr);
+}
+export function uploadF16(device, f32arr) {
+  const f16 = new Float16Array(f32arr);
+  const buf = createBuffer(device, f16.byteLength);
+  device.queue.writeBuffer(buf, 0, f16.buffer, f16.byteOffset, f16.byteLength);
+  return buf;
+}

web/index.html ADDED Viewed

	@@ -0,0 +1,81 @@

+<!doctype html>
+<html lang="en">
+<meta charset="utf-8" />
+<meta name="viewport" content="width=device-width, initial-scale=1" />
+<title>Qwen3.5-2B · WebGPU</title>
+<style>
+  :root {
+    --bg: #0b0d10; --panel: #14171c; --line: #232830; --fg: #e8ebef;
+    --dim: #8b93a1; --accent: #7c5cff; --ok: #3ecf8e; --err: #ff6b6b;
+  }
+  * { box-sizing: border-box; }
+  body {
+    margin: 0; background: var(--bg); color: var(--fg);
+    font: 15px/1.55 ui-sans-serif, system-ui, -apple-system, "Segoe UI", sans-serif;
+    height: 100dvh; display: flex; flex-direction: column;
+  }
+  header {
+    display: flex; align-items: baseline; gap: 12px; padding: 12px 18px;
+    border-bottom: 1px solid var(--line); flex: none;
+  }
+  .wordmark { font-weight: 700; letter-spacing: .2px; }
+  .sub { color: var(--dim); font-size: 12.5px; }
+  #gpuname { margin-left: auto; color: var(--dim); font-size: 12.5px; }
+  #toksbox { color: var(--dim); font-size: 12.5px; }
+  #toks { color: var(--ok); font-variant-numeric: tabular-nums; }
+  #chat { flex: 1; overflow-y: auto; padding: 18px; }
+  .msg { max-width: 780px; margin: 0 auto 16px; }
+  .msg .who { font-size: 11px; text-transform: uppercase; letter-spacing: .8px; color: var(--dim); margin-bottom: 4px; }
+  .msg.user .who { color: var(--accent); }
+  .msg .body { white-space: pre-wrap; word-break: break-word; }
+  .stats { font-size: 11.5px; color: var(--dim); margin-top: 6px; }
+  .caret { display: inline-block; width: 7px; height: 15px; background: var(--accent); animation: blink 1s steps(1) infinite; vertical-align: -2px; }
+  @keyframes blink { 50% { opacity: 0; } }
+  form { flex: none; display: flex; gap: 10px; padding: 14px 18px 18px; border-top: 1px solid var(--line); }
+  #input {
+    flex: 1; max-width: 780px; margin-left: auto; resize: none; max-height: 160px;
+    background: var(--panel); color: var(--fg); border: 1px solid var(--line);
+    border-radius: 10px; padding: 10px 12px; font: inherit; outline: none;
+  }
+  #input:focus { border-color: var(--accent); }
+  #send {
+    margin-right: auto; background: var(--accent); border: 0; color: #fff;
+    border-radius: 10px; padding: 0 18px; font: inherit; font-weight: 600; cursor: pointer;
+  }
+  #send:disabled, #input:disabled { opacity: .45; cursor: default; }
+  #boot {
+    position: fixed; inset: 0; background: var(--bg); display: flex; flex-direction: column;
+    align-items: center; justify-content: center; gap: 14px; transition: opacity .4s; z-index: 10;
+  }
+  #boot.done { opacity: 0; pointer-events: none; }
+  #boottitle { font-size: 22px; font-weight: 700; }
+  #bootsub { color: var(--dim); font-size: 13px; max-width: 420px; text-align: center; }
+  #bar { width: 300px; height: 6px; background: var(--line); border-radius: 3px; overflow: hidden; }
+  #barfill { width: 100%; height: 100%; background: var(--accent); transform: scaleX(0); transform-origin: left; transition: transform .2s; }
+  #log { font: 12px/1.7 ui-monospace, monospace; color: var(--dim); width: 420px; }
+  #log .ok { color: var(--ok); }
+  #log .err { color: var(--err); }
+</style>
+<body>
+  <header>
+    <span class="wordmark">Qwen3.5-2B</span>
+    <span class="sub">hybrid DeltaNet + attention · hand-written WGSL · int4</span>
+    <span id="toksbox"><span id="toks">–</span> tok/s</span>
+    <span id="gpuname"></span>
+  </header>
+  <div id="chat"><div id="msgs"></div></div>
+  <form id="form">
+    <textarea id="input" rows="1" placeholder="Message… (Enter to send)" disabled></textarea>
+    <button id="send" type="submit" disabled>Send</button>
+  </form>
+  <div id="boot">
+    <div id="boottitle">Qwen3.5-2B</div>
+    <div id="bootsub">First hybrid gated-DeltaNet LLM in the browser — pure WebGPU, hand-written WGSL kernels.
+      Downloading <span id="dlmb">~1.1 GB</span> of int4 weights <span id="dlpct"></span></div>
+    <div id="bar"><div id="barfill"></div></div>
+    <div id="dlnote" class="sub"></div>
+    <div id="log"></div>
+  </div>
+  <script type="module" src="/web/main.js"></script>
+</body>
+</html>

web/kernels.js ADDED Viewed

	@@ -0,0 +1,1361 @@

+// WGSL kernel sources for Qwen3.5 (hybrid gated DeltaNet + gated attention)
+// inference. All shapes are baked in as consts at pipeline-creation time;
+// only per-step values (t, pos0) live in a small step buffer.
+//
+// Conventions (inherited from the lfm2 project):
+//   - weights: f16 (f32 when !F16), row-major [out, in], bound as array<vec4<HT>>
+//   - activations: f32; DeltaNet recurrent state: always f32
+//   - step buffer: { t: u32, pos0: u32, sstate: u32, acc: u32 } in storage so
+//     GPU kernels can bump it (GPU-resident decode loop, no CPU roundtrip)
+// Whether the device supports WGSL subgroups (set by initDevice). Without
+// them (iOS Safari) matvec kernels emit workgroup-shared reductions.
+export let SG = true;
+export function setSubgroups(v) {
+  SG = v;
+}
+// Whether the device supports shader-f16 (set by initDevice). Without it
+// (Dawn gates f16 on NVIDIA/Linux; Qualcomm mobiles lack it) the runtime
+// converts every f16 tensor to f32 at load time; kernels just switch the
+// bound array type, since all reads are already wrapped in f32() casts.
+export let F16 = true;
+export function setF16(v) {
+  F16 = v;
+}
+// half-type helpers for kernel templates (do not touch these two lines in
+// bulk replacements — see the lfm2 EF16 regression)
+const HT = () => (F16 ? "f16" : "f32");
+const EF16 = () => (F16 ? "enable f16;" : "");
+export const STEP_UNIFORM = `
+struct Step { t: u32, pos0: u32, sstate: u32, acc: u32 }
+@group(0) @binding(15) var<storage, read> step: Step;
+`;
+// ---------------------------------------------------------------------------
+// Gated DeltaNet: recurrent decode step (one token), all heads in one dispatch.
+//
+// Math (per head h, state S[dk, dv], everything f32):
+//   qn = l2norm(q) / sqrt(DK);  kn = l2norm(k)         (eps 1e-6)
+//   decay = exp(g_h)                                    (g_h < 0)
+//   kv_mem[dv] = sum_dk S[dk,dv]*decay * kn[dk]
+//   delta[dv]  = (v[dv] - kv_mem[dv]) * beta_h
+//   S[dk,dv]   = S[dk,dv]*decay + kn[dk]*delta[dv]
+//   out[dv]    = sum_dk S[dk,dv] * qn[dk]
+//
+// One workgroup per head, WG size = DV (=128), thread li owns column dv=li.
+// Single read pass computes acc_q = S·q_raw, acc_k = S·k_raw per column plus
+// shared reductions of q·q, k·k, q·k; out folds algebraically:
+//   out = decay*acc_q*invq + delta * dot(kn,qn)
+// so S is touched twice (read + read-modify-write), never three times.
+//
+// Bindings:
+//   0 qkv   [T, QKV] f32   (q | k | v concatenated, conv+silu already applied)
+//   1 g     [T, H]   f32   (log-decay, precomputed by deltaGates)
+//   2 beta  [T, H]   f32
+//   3 state [H, DK, DV] f32 read_write
+//   4 out   [T, H*DV] f32  (core_attn_out, pre RMSNormGated)
+//   15 step (t = token index into qkv/g/beta/out)
+export function deltaStep({ H = 16, DK = 128, DV = 128 } = {}) {
+  const QKV = 2 * H * DK + H * DV;
+  return /* wgsl */ `
+@group(0) @binding(0) var<storage, read> qkv: array<f32>;
+@group(0) @binding(1) var<storage, read> g: array<f32>;
+@group(0) @binding(2) var<storage, read> beta: array<f32>;
+@group(0) @binding(3) var<storage, read_write> state: array<f32>;
+@group(0) @binding(4) var<storage, read_write> out: array<f32>;
+${STEP_UNIFORM}
+const H = ${H}u;
+const DK = ${DK}u;
+const DV = ${DV}u;
+const QKV = ${QKV}u;
+const KOFF = ${H * DK}u;
+const VOFF = ${2 * H * DK}u;
+const SCALE = ${(1 / Math.sqrt(DK)).toPrecision(17)};
+var<workgroup> qs: array<f32, DK>;
+var<workgroup> ks: array<f32, DK>;
+var<workgroup> red: array<vec3<f32>, DK>;
+@compute @workgroup_size(${DV})
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  let h = wg.x;
+  let ti = step.t - 1u;   // step.t = token count; this kernel does the last one
+  let base = ti * QKV;
+  // stage raw q/k for this head; each thread also keeps its v component
+  qs[li] = qkv[base + h * DK + li];
+  ks[li] = qkv[base + KOFF + h * DK + li];
+  let v_t = qkv[base + VOFF + h * DV + li];
+  red[li] = vec3<f32>(qs[li] * qs[li], ks[li] * ks[li], qs[li] * ks[li]);
+  workgroupBarrier();
+  // tree-reduce (q·q, k·k, q·k); portable, cost is negligible at this size
+  for (var s = DK / 2u; s > 0u; s >>= 1u) {
+    if (li < s) { red[li] += red[li + s]; }
+    workgroupBarrier();
+  }
+  let sums = red[0];
+  let invq = inverseSqrt(sums.x + 1e-6) * SCALE;
+  let invk = inverseSqrt(sums.y + 1e-6);
+  let dot_kq = sums.z * invq * invk;
+  let decay = exp(g[ti * H + h]);
+  let beta_h = beta[ti * H + h];
+  // pass 1 over S: projections of the old state onto k and q
+  let sbase = h * DK * DV + li;
+  var acc_q = 0.0;
+  var acc_k = 0.0;
+  for (var dk = 0u; dk < DK; dk++) {
+    let s = state[sbase + dk * DV];
+    acc_q += s * qs[dk];
+    acc_k += s * ks[dk];
+  }
+  let kv_mem = decay * acc_k * invk;
+  let delta = (v_t - kv_mem) * beta_h;
+  let out_v = decay * acc_q * invq + dot_kq * delta;
+  // pass 2 over S: decay + rank-1 update
+  for (var dk = 0u; dk < DK; dk++) {
+    let idx = sbase + dk * DV;
+    state[idx] = state[idx] * decay + ks[dk] * invk * delta;
+  }
+  out[ti * H * DV + h * DV + li] = out_v;
+}`;
+}
+// ---------------------------------------------------------------------------
+// Gated DeltaNet: chunked prefill (WY representation), chunk size 64.
+// Reference: torch_chunk_gated_delta_rule (transformers qwen3_next).
+//
+// Stage 1 (deltaChunkLocal) is parallel over (chunk, head): per chunk it
+// l2-normalizes q/k, cumsums g, builds the strictly-lower matrix
+//   A[i][j] = -beta_i * kn_i·kn_j * exp(gc_i - gc_j)
+// runs the forward substitution T = (I - A)^{-1} in workgroup memory, and
+// stores qn, kn, tv = T@(v*beta), kcd = T@(kn*beta*exp(gc)), gc into `loc`.
+//
+// Stage 2 (deltaChunkScan) is the sequential inter-chunk scan, one workgroup
+// per head; thread li owns state column S[:, li] in private registers for the
+// whole segment (state storage is touched once at start and once at end):
+//   v_new = tv - kcd @ S
+//   out   = (qn*e^gc) @ S + tril(qn·kn * e^(gc_r-gc_j)) @ v_new
+//   S     = S*e^glast + (kn*e^(glast-gc))^T @ v_new
+// v_new columns never cross threads, so they live in private memory too.
+//
+// loc layout per (chunk, head) block, floats:
+//   qn 64*128 | kn 64*128 | tv 64*128 | kcd 64*128 | gc 64   (stride 32832)
+// step.t = token count of the segment; chunk count = ceil(t / 64).
+export const DELTA_BLK = 4 * 64 * 128 + 64;
+export function deltaChunkLocal({ H = 16, DK = 128, DV = 128 } = {}) {
+  const QKV = 2 * H * DK + H * DV;
+  return /* wgsl */ `
+@group(0) @binding(0) var<storage, read> qkv: array<f32>;
+@group(0) @binding(1) var<storage, read> g: array<f32>;
+@group(0) @binding(2) var<storage, read> beta: array<f32>;
+@group(0) @binding(3) var<storage, read_write> loc: array<f32>;
+${STEP_UNIFORM}
+const H = ${H}u;
+const DK = ${DK}u;
+const DV = ${DV}u;
+const QKV = ${QKV}u;
+const KOFF = ${H * DK}u;
+const VOFF = ${2 * H * DK}u;
+const BLK = ${DELTA_BLK}u;
+const SCALE = ${(1 / Math.sqrt(DK)).toPrecision(17)};
+var<workgroup> Tm: array<f32, 4096>;
+var<workgroup> rowbuf: array<f32, 64>;
+var<workgroup> gcs: array<f32, 64>;
+var<workgroup> invq: array<f32, 64>;
+var<workgroup> invk: array<f32, 64>;
+var<workgroup> betas: array<f32, 64>;
+var<workgroup> vrow: array<f32, ${DV}>;
+var<workgroup> krow: array<f32, ${DK}>;
+var<workgroup> wgT: u32;
+@compute @workgroup_size(128)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  if (li == 0u) { wgT = step.t; }
+  let nT = workgroupUniformLoad(&wgT);
+  let c = wg.x / H;
+  let h = wg.x % H;
+  let blk = wg.x * BLK;
+  // per-token l2 norms, beta, raw g
+  if (li < 64u) {
+    let tok = c * 64u + li;
+    var sq = 0.0;
+    var sk = 0.0;
+    if (tok < nT) {
+      let qb = tok * QKV + h * DK;
+      let kb = tok * QKV + KOFF + h * DK;
+      for (var d = 0u; d < DK; d++) {
+        let qv = qkv[qb + d];
+        let kv = qkv[kb + d];
+        sq += qv * qv;
+        sk += kv * kv;
+      }
+      invq[li] = inverseSqrt(sq + 1e-6) * SCALE;
+      invk[li] = inverseSqrt(sk + 1e-6);
+      betas[li] = beta[tok * H + h];
+      gcs[li] = g[tok * H + h];
+    } else {
+      invq[li] = 0.0; invk[li] = 0.0; betas[li] = 0.0; gcs[li] = 0.0;
+    }
+  }
+  workgroupBarrier();
+  if (li == 0u) {
+    for (var i = 1u; i < 64u; i++) { gcs[i] += gcs[i - 1u]; }
+  }
+  workgroupBarrier();
+  // A[i][j] (strictly lower) into Tm
+  for (var idx = li; idx < 4096u; idx += 128u) {
+    let i = idx >> 6u;
+    let j = idx & 63u;
+    var val = 0.0;
+    if (j < i) {
+      let toki = c * 64u + i;
+      let tokj = c * 64u + j;
+      if (tokj < nT && toki < nT) {
+        let kbi = toki * QKV + KOFF + h * DK;
+        let kbj = tokj * QKV + KOFF + h * DK;
+        var dt = 0.0;
+        for (var d = 0u; d < DK; d++) { dt += qkv[kbi + d] * qkv[kbj + d]; }
+        val = -betas[i] * dt * invk[i] * invk[j] * exp(gcs[i] - gcs[j]);
+      }
+    }
+    Tm[idx] = val;
+  }
+  workgroupBarrier();
+  // forward substitution: T = (I - A)^{-1}, unit lower-triangular
+  for (var i = 1u; i < 64u; i++) {
+    if (li < 64u) { rowbuf[li] = Tm[i * 64u + li]; }
+    workgroupBarrier();
+    if (li < i) {
+      var acc = rowbuf[li];
+      for (var m = li + 1u; m < i; m++) { acc += rowbuf[m] * Tm[m * 64u + li]; }
+      Tm[i * 64u + li] = acc;
+    }
+    workgroupBarrier();
+  }
+  if (li < 64u) { Tm[li * 64u + li] = 1.0; }
+  workgroupBarrier();
+  // tv = T @ (v*beta), kcd = T @ (kn*beta*e^gc); accumulate row-wise while
+  // staging one (v, k) token row at a time
+  var accV = array<f32, 64>();
+  var accK = array<f32, 64>();
+  for (var j = 0u; j < 64u; j++) {
+    let tok = c * 64u + j;
+    if (tok < nT) {
+      vrow[li] = qkv[tok * QKV + VOFF + h * DV + li] * betas[j];
+      krow[li] = qkv[tok * QKV + KOFF + h * DK + li] * invk[j] * betas[j] * exp(gcs[j]);
+    } else {
+      vrow[li] = 0.0; krow[li] = 0.0;
+    }
+    workgroupBarrier();
+    for (var r = j; r < 64u; r++) {
+      let tj = Tm[r * 64u + j];
+      accV[r] += tj * vrow[li];
+      accK[r] += tj * krow[li];
+    }
+    workgroupBarrier();
+  }
+  // store qn, kn, tv, kcd, gc
+  for (var r = 0u; r < 64u; r++) {
+    let tok = c * 64u + r;
+    var qn = 0.0;
+    var kn = 0.0;
+    if (tok < nT) {
+      qn = qkv[tok * QKV + h * DK + li] * invq[r];
+      kn = qkv[tok * QKV + KOFF + h * DK + li] * invk[r];
+    }
+    loc[blk + r * 128u + li] = qn;
+    loc[blk + 8192u + r * 128u + li] = kn;
+    loc[blk + 16384u + r * 128u + li] = accV[r];
+    loc[blk + 24576u + r * 128u + li] = accK[r];
+  }
+  if (li < 64u) { loc[blk + 32768u + li] = gcs[li]; }
+}`;
+}
+export function deltaChunkScan({ H = 16, DK = 128, DV = 128 } = {}) {
+  return /* wgsl */ `
+@group(0) @binding(0) var<storage, read> loc: array<f32>;
+@group(0) @binding(1) var<storage, read_write> state: array<f32>;
+@group(0) @binding(2) var<storage, read_write> out: array<f32>;
+${STEP_UNIFORM}
+const H = ${H}u;
+const DK = ${DK}u;
+const DV = ${DV}u;
+const BLK = ${DELTA_BLK}u;
+var<workgroup> alocal: array<f32, 4096>;
+var<workgroup> arow: array<f32, ${DK}>;
+var<workgroup> brow: array<f32, ${DK}>;
+var<workgroup> gcs: array<f32, 64>;
+var<workgroup> wgT: u32;
+@compute @workgroup_size(128)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  if (li == 0u) { wgT = step.t; }
+  let nT = workgroupUniformLoad(&wgT);
+  let nC = (nT + 63u) / 64u;
+  let h = wg.x;
+  // thread li owns state column S[:, li] for the whole segment
+  var s = array<f32, ${DK}>();
+  for (var dk = 0u; dk < DK; dk++) { s[dk] = state[h * DK * DV + dk * DV + li]; }
+  var vncol = array<f32, 64>();
+  for (var c = 0u; c < nC; c++) {
+    let blk = (c * H + h) * BLK;
+    if (li < 64u) { gcs[li] = loc[blk + 32768u + li]; }
+    workgroupBarrier();
+    let glast = gcs[63];
+    // intra-chunk attention: alocal[r][j] = qn_r·kn_j * e^(gc_r-gc_j), j<=r
+    for (var idx = li; idx < 4096u; idx += 128u) {
+      let r = idx >> 6u;
+      let j = idx & 63u;
+      var val = 0.0;
+      if (j <= r) {
+        var dt = 0.0;
+        for (var d = 0u; d < DK; d++) {
+          dt += loc[blk + r * 128u + d] * loc[blk + 8192u + j * 128u + d];
+        }
+        val = dt * exp(gcs[r] - gcs[j]);
+      }
+      alocal[idx] = val;
+    }
+    workgroupBarrier();
+    // rows in order: v_new, then out (v_new columns stay in this thread)
+    for (var r = 0u; r < 64u; r++) {
+      arow[li] = loc[blk + 24576u + r * 128u + li];            // kcd row
+      brow[li] = loc[blk + r * 128u + li] * exp(gcs[r]);       // qn*e^gc row
+      workgroupBarrier();
+      var vp = 0.0;
+      var ai = 0.0;
+      for (var dk = 0u; dk < DK; dk++) {
+        vp += arow[dk] * s[dk];
+        ai += brow[dk] * s[dk];
+      }
+      let vn = loc[blk + 16384u + r * 128u + li] - vp;
+      vncol[r] = vn;
+      var o = ai;
+      for (var j = 0u; j <= r; j++) { o += alocal[r * 64u + j] * vncol[j]; }
+      let tok = c * 64u + r;
+      if (tok < nT) { out[tok * H * DV + h * DV + li] = o; }
+      workgroupBarrier();
+    }
+    // S = S*e^glast + sum_r kn_r*e^(glast-gc_r) (outer) v_new_r
+    let dl = exp(glast);
+    for (var dk = 0u; dk < DK; dk++) { s[dk] *= dl; }
+    for (var r = 0u; r < 64u; r++) {
+      arow[li] = loc[blk + 8192u + r * 128u + li] * exp(glast - gcs[r]);
+      workgroupBarrier();
+      let vn = vncol[r];
+      for (var dk = 0u; dk < DK; dk++) { s[dk] += arow[dk] * vn; }
+      workgroupBarrier();
+    }
+  }
+  for (var dk = 0u; dk < DK; dk++) { state[h * DK * DV + dk * DV + li] = s[dk]; }
+}`;
+}
+// ---------------------------------------------------------------------------
+// DeltaNet plumbing: gates, causal conv, gated RMSNorm
+// ---------------------------------------------------------------------------
+// beta = sigmoid(b), g = -exp(A_log) * softplus(a + dt_bias).
+// b|a live in the fused in_proj activation at row stride `row`, offset `off`;
+// params = A_log|dt_bias [2*H] f32.
+export function deltaGates({ H = 16, row = 2 * 16, off = 0 } = {}) {
+  return /* wgsl */ `
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> ba: array<f32>;
+@group(0) @binding(1) var<storage, read> params: array<f32>;
+@group(0) @binding(2) var<storage, read_write> g: array<f32>;
+@group(0) @binding(3) var<storage, read_write> beta: array<f32>;
+const H = ${H}u;
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= step.t * H) { return; }
+  let t = i / H;
+  let h = i % H;
+  let b = ba[t * ${row}u + ${off}u + h];
+  let a = ba[t * ${row}u + ${off}u + H + h];
+  beta[i] = 1.0 / (1.0 + exp(-b));
+  let x = a + params[H + h];
+  let sp = select(log(1.0 + exp(x)), x, x > 20.0);
+  g[i] = -exp(params[h]) * sp;
+}`;
+}
+// Depthwise causal conv (kernel 4, bias-free) + SiLU over the fused qkv
+// activation (first C of each `xrow`-wide row). state[c, 0..3] = raw inputs
+// at t-3..t (persistent). One thread per channel, sequential over t.
+export function convQkv({ C, xrow = null }) {
+  const XROW = xrow ?? C;
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> x: array<f32>;
+@group(0) @binding(1) var<storage, read> w: array<${HT()}>;       // [C, 4]
+@group(0) @binding(2) var<storage, read_write> state: array<f32>; // [C, 4]
+@group(0) @binding(3) var<storage, read_write> y: array<f32>;     // [T, C]
+const C = ${C}u;
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let c = gid.x;
+  if (c >= C) { return; }
+  let w0 = f32(w[c * 4u]);
+  let w1 = f32(w[c * 4u + 1u]);
+  let w2 = f32(w[c * 4u + 2u]);
+  let w3 = f32(w[c * 4u + 3u]);
+  var s0 = state[c * 4u];
+  var s1 = state[c * 4u + 1u];
+  var s2 = state[c * 4u + 2u];
+  var s3 = state[c * 4u + 3u];
+  for (var t = 0u; t < step.t; t++) {
+    let xv = x[t * ${XROW}u + c];
+    let acc = w0 * s1 + w1 * s2 + w2 * s3 + w3 * xv;
+    y[t * C + c] = acc / (1.0 + exp(-acc));
+    s0 = s1; s1 = s2; s2 = s3; s3 = xv;
+  }
+  state[c * 4u] = s0;
+  state[c * 4u + 1u] = s1;
+  state[c * 4u + 2u] = s2;
+  state[c * 4u + 3u] = s3;
+}`;
+}
+// Gated RMSNorm per value head (dim DV): y = rmsnorm(x)*w * silu(z).
+// Weight is per head-dim, shared across heads (plain scale, NOT (1+w)).
+// z lives in the fused in_proj activation at row stride zrow, offset zoff.
+// One workgroup per (t, head); wg.x = t*H + h.
+export function gatedNorm({ H = 16, DV = 128, eps = 1e-6, zrow = null, zoff = 0 } = {}) {
+  const ZROW = zrow ?? H * DV;
+  return /* wgsl */ `
+${EF16()}
+@group(0) @binding(0) var<storage, read> x: array<f32>;
+@group(0) @binding(1) var<storage, read> z: array<f32>;
+@group(0) @binding(2) var<storage, read> w: array<${HT()}>;
+@group(0) @binding(3) var<storage, read_write> y: array<f32>;
+const DV = ${DV}u;
+const H = ${H}u;
+const EPS = ${eps};
+var<workgroup> red: array<f32, ${DV}>;
+@compute @workgroup_size(${DV})
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  let base = wg.x * DV;
+  let t = wg.x / H;
+  let h = wg.x % H;
+  let v = x[base + li];
+  red[li] = v * v;
+  workgroupBarrier();
+  for (var s = ${DV / 2}u; s > 0u; s >>= 1u) {
+    if (li < s) { red[li] += red[li + s]; }
+    workgroupBarrier();
+  }
+  let inv = inverseSqrt(red[0] / f32(DV) + EPS);
+  let zv = z[t * ${ZROW}u + ${zoff}u + h * DV + li];
+  y[base + li] = v * inv * f32(w[li]) * (zv / (1.0 + exp(-zv)));
+}`;
+}
+// ---------------------------------------------------------------------------
+// Gated attention (full_attention layers)
+// ---------------------------------------------------------------------------
+// Post-QKV processing for Qwen3.5 gated attention. Fused projection rows:
+//   [ NH * (q 256 | gate 256) | NKV k 256 | NKV v 256 ]  (ROW = 5120)
+// Per-head RMSNorm over the full HD on q and k (weights pre-normalized to
+// plain scale by the converter), then neox RoPE on the first RD dims only
+// (theta 1e7), remaining dims pass through. Gates stay in the qkv activation
+// (gateMul reads them after attention). wg.x = t*(NH+NKV)+slot; WG = HD.
+export function qkvRope({ NH = 8, NKV = 2, HD = 256, RD = 64, eps = 1e-6, ropeTheta = 1e7, maxSeq = 4096 } = {}) {
+  const ROW = NH * HD * 2 + 2 * NKV * HD;
+  const KO = NH * HD * 2;
+  const VO = NH * HD * 2 + NKV * HD;
+  const RH = RD / 2;
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> qkv: array<f32>;       // [T, ${ROW}]
+@group(0) @binding(3) var<storage, read> qw: array<${HT()}>;    // [HD]
+@group(0) @binding(4) var<storage, read> kw: array<${HT()}>;    // [HD]
+@group(0) @binding(5) var<storage, read_write> qout: array<f32>;  // [T, NH*HD]
+@group(0) @binding(6) var<storage, read_write> kcache: array<${HT()}>; // [maxSeq, NKV*HD]
+@group(0) @binding(7) var<storage, read_write> vcache: array<${HT()}>; // [maxSeq, NKV*HD]
+const NH = ${NH}u;
+const NKV = ${NKV}u;
+const HD = ${HD}u;
+const RD = ${RD}u;
+const RH = ${RH}u;
+const EPS = ${eps};
+const SLOTS = ${NH + NKV}u;
+const ROW = ${ROW}u;
+var<workgroup> red: array<f32, ${HD}>;
+var<workgroup> sn: array<f32, ${HD}>;
+fn ropeMix(d: u32, pos: u32) -> f32 {
+  // normed value at dim d after partial neox rope (reads staged sn[])
+  if (d >= RD) { return sn[d]; }
+  let p = d % RH;
+  let theta = f32(pos) * pow(${ropeTheta}, -f32(p) / f32(RH));
+  let c = cos(theta);
+  let s = sin(theta);
+  if (d < RH) { return sn[d] * c - sn[d + RH] * s; }
+  return sn[d] * c + sn[d - RH] * s;
+}
+@compute @workgroup_size(${HD})
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) d: u32) {
+  let t = wg.x / SLOTS;
+  if (t >= step.t) { return; }
+  let slot = wg.x % SLOTS;
+  let pos = step.pos0 + t;
+  if (slot < NH) {
+    let base = t * ROW + slot * HD * 2u;   // q part of this head
+    let v = qkv[base + d];
+    red[d] = v * v;
+    workgroupBarrier();
+    for (var s = ${HD / 2}u; s > 0u; s >>= 1u) {
+      if (d < s) { red[d] += red[d + s]; }
+      workgroupBarrier();
+    }
+    let inv = inverseSqrt(red[0] / f32(HD) + EPS);
+    sn[d] = v * inv * f32(qw[d]);
+    workgroupBarrier();
+    qout[t * NH * HD + slot * HD + d] = ropeMix(d, pos);
+  } else {
+    let h = slot - NH;
+    let kbase = t * ROW + ${KO}u + h * HD;
+    let vbase = t * ROW + ${VO}u + h * HD;
+    let cbase = pos * NKV * HD + h * HD;
+    let v = qkv[kbase + d];
+    red[d] = v * v;
+    workgroupBarrier();
+    for (var s = ${HD / 2}u; s > 0u; s >>= 1u) {
+      if (d < s) { red[d] += red[d + s]; }
+      workgroupBarrier();
+    }
+    let inv = inverseSqrt(red[0] / f32(HD) + EPS);
+    sn[d] = v * inv * f32(kw[d]);
+    workgroupBarrier();
+    kcache[cbase + d] = ${HT()}(ropeMix(d, pos));
+    vcache[cbase + d] = ${HT()}(qkv[vbase + d]);
+  }
+}`;
+}
+// Causal attention over the KV cache (GQA). One workgroup per (t, head).
+// Same structure as the lfm2 kernel, sized for HD=256.
+export function attention({ NH = 8, NKV = 2, HD = 256, maxSeq = 4096 } = {}) {
+  const group = NH / NKV;
+  const scale = 1 / Math.sqrt(HD);
+  const HD4 = HD / 4;
+  const KVW4 = (NKV * HD) / 4;
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> q: array<vec4<f32>>;        // [T, NH*HD/4]
+@group(0) @binding(1) var<storage, read> kcache: array<vec4<${HT()}>>;   // [maxSeq, NKV*HD/4]
+@group(0) @binding(2) var<storage, read> vcache: array<vec4<${HT()}>>;   // [maxSeq, NKV*HD/4]
+@group(0) @binding(3) var<storage, read_write> out: array<vec4<f32>>; // [T, NH*HD/4]
+const NH = ${NH}u;
+const HD4 = ${HD4}u;
+const KVW4 = ${KVW4}u;
+const SCALE = ${scale};
+var<workgroup> qs4: array<vec4<f32>, ${HD4}>;
+var<workgroup> scores: array<f32, ${maxSeq}>;
+var<workgroup> red: array<f32, ${HD}>;
+var<workgroup> red4: array<vec4<f32>, ${HD}>;
+@compute @workgroup_size(${HD})
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  let t = wg.x / NH;
+  if (t >= step.t) { return; }
+  let h = wg.x % NH;
+  let kvh = h / ${group}u;
+  let S = step.pos0 + t + 1u;   // visible keys
+  if (li < HD4) { qs4[li] = q[(t * NH + h) * HD4 + li]; }
+  workgroupBarrier();
+  var lmax = -3.4e38;
+  for (var j = li; j < S; j += ${HD}u) {
+    var dp = 0.0;
+    let kbase = j * KVW4 + kvh * HD4;
+    for (var d = 0u; d < HD4; d++) {
+      dp += dot(vec4<f32>(kcache[kbase + d]), qs4[d]);
+    }
+    let sc = dp * SCALE;
+    scores[j] = sc;
+    lmax = max(lmax, sc);
+  }
+  red[li] = lmax;
+  workgroupBarrier();
+  var stride = ${HD / 2}u;
+  while (stride > 0u) {
+    if (li < stride) { red[li] = max(red[li], red[li + stride]); }
+    workgroupBarrier();
+    stride /= 2u;
+  }
+  let m = red[0];
+  workgroupBarrier();
+  var lsum = 0.0;
+  for (var j = li; j < S; j += ${HD}u) {
+    let e = exp(scores[j] - m);
+    scores[j] = e;
+    lsum += e;
+  }
+  red[li] = lsum;
+  workgroupBarrier();
+  stride = ${HD / 2}u;
+  while (stride > 0u) {
+    if (li < stride) { red[li] += red[li + stride]; }
+    workgroupBarrier();
+    stride /= 2u;
+  }
+  let z = red[0];
+  let d4 = li % HD4;
+  let blk = li / HD4;
+  var acc = vec4<f32>(0.0);
+  for (var j = blk; j < S; j += ${HD / HD4}u) {
+    acc += scores[j] * vec4<f32>(vcache[j * KVW4 + kvh * HD4 + d4]);
+  }
+  red4[li] = acc;
+  workgroupBarrier();
+  if (li < HD4) {
+    var o = red4[li];
+    for (var b = 1u; b < ${HD / HD4}u; b++) {
+      o += red4[li + b * HD4];
+    }
+    out[(t * NH + h) * HD4 + li] = o / z;
+  }
+}`;
+}
+// attn_out *= sigmoid(gate); gates live in the fused qkv activation at
+// per-head offset h*2*HD + HD. y is [T, NH*HD].
+export function gateMul({ NH = 8, NKV = 2, HD = 256 } = {}) {
+  const ROW = NH * HD * 2 + 2 * NKV * HD;
+  return /* wgsl */ `
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> qkv: array<f32>;
+@group(0) @binding(1) var<storage, read_write> y: array<f32>;
+const NH = ${NH}u;
+const HD = ${HD}u;
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= step.t * NH * HD) { return; }
+  let t = i / (NH * HD);
+  let c = i % (NH * HD);
+  let h = c / HD;
+  let d = c % HD;
+  let gv = qkv[t * ${ROW}u + h * HD * 2u + HD + d];
+  y[i] = y[i] / (1.0 + exp(-gv));
+}`;
+}
+// ---------------------------------------------------------------------------
+// Generic kernels, copied from the lfm2 project (f16 path)
+// ---------------------------------------------------------------------------
+// y[row] = dot(W[row, :], x[:])  -- decode-path matvec, one subgroup per row.
+export function gemv({ N, K }) {
+  const K4 = K / 4;
+  return /* wgsl */ `
+${EF16()}
+enable subgroups;
+@group(0) @binding(0) var<storage, read> w: array<vec4<${HT()}>>;
+@group(0) @binding(1) var<storage, read> x: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+const N = ${N}u;
+const K4 = ${K4}u;
+@compute @workgroup_size(256)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_index) li: u32,
+  @builtin(subgroup_invocation_id) sgi: u32,
+  @builtin(subgroup_size) sgs: u32,
+) {
+  let rows_per_wg = 256u / sgs;
+  let row = wg.x * rows_per_wg + li / sgs;
+  let base = min(row, N - 1u) * K4;
+  var acc = 0.0;
+  for (var i = sgi; i < K4; i += sgs) {
+    acc += dot(vec4<f32>(w[base + i]), x[i]);
+  }
+  acc = subgroupAdd(acc);
+  if (sgi == 0u && row < N) { y[row] = acc; }
+}`;
+}
+// gemv with fused residual add: y[row] += dot(W[row,:], x).
+export function gemvAdd({ N, K }) {
+  const K4 = K / 4;
+  return /* wgsl */ `
+${EF16()}
+enable subgroups;
+@group(0) @binding(0) var<storage, read> w: array<vec4<${HT()}>>;
+@group(0) @binding(1) var<storage, read> x: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+const N = ${N}u;
+const K4 = ${K4}u;
+@compute @workgroup_size(256)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_index) li: u32,
+  @builtin(subgroup_invocation_id) sgi: u32,
+  @builtin(subgroup_size) sgs: u32,
+) {
+  let rows_per_wg = 256u / sgs;
+  let row = wg.x * rows_per_wg + li / sgs;
+  let base = min(row, N - 1u) * K4;
+  var acc = 0.0;
+  for (var i = sgi; i < K4; i += sgs) {
+    acc += dot(vec4<f32>(w[base + i]), x[i]);
+  }
+  acc = subgroupAdd(acc);
+  if (sgi == 0u && row < N) { y[row] = y[row] + acc; }
+}`;
+}
+// fused SwiGLU matvec: y = silu(W1@x) * (W3@x)
+export function gemvSwiglu({ N, K }) {
+  const K4 = K / 4;
+  return /* wgsl */ `
+${EF16()}
+enable subgroups;
+@group(0) @binding(0) var<storage, read> w1: array<vec4<${HT()}>>;
+@group(0) @binding(1) var<storage, read> x: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+@group(0) @binding(3) var<storage, read> w3: array<vec4<${HT()}>>;
+const N = ${N}u;
+const K4 = ${K4}u;
+@compute @workgroup_size(256)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_index) li: u32,
+  @builtin(subgroup_invocation_id) sgi: u32,
+  @builtin(subgroup_size) sgs: u32,
+) {
+  let row = wg.x * (256u / sgs) + li / sgs;
+  let base = min(row, N - 1u) * K4;
+  var a = 0.0;
+  var b = 0.0;
+  for (var i = sgi; i < K4; i += sgs) {
+    let xv = x[i];
+    a += dot(vec4<f32>(w1[base + i]), xv);
+    b += dot(vec4<f32>(w3[base + i]), xv);
+  }
+  a = subgroupAdd(a);
+  b = subgroupAdd(b);
+  if (sgi == 0u && row < N) { y[row] = a / (1.0 + exp(-a)) * b; }
+}`;
+}
+// naive tiled-less GEMM for prefill: y[m, n] = dot(W[n,:], x[m,:])
+export function gemm({ N, K }) {
+  const K4 = K / 4;
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> w: array<vec4<${HT()}>>;
+@group(0) @binding(1) var<storage, read> x: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+const N = ${N}u;
+const K4 = ${K4}u;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let n = gid.x;
+  let m = gid.y;
+  if (n >= N || m >= step.t) { return; }
+  var acc = 0.0;
+  let wbase = n * K4;
+  let xbase = m * K4;
+  for (var i = 0u; i < K4; i++) {
+    acc += dot(vec4<f32>(w[wbase + i]), x[xbase + i]);
+  }
+  y[m * N + n] = acc;
+}`;
+}
+// gemm with fused residual add: y[m, n] += dot(W[n,:], x[m,:])
+export function gemmAdd({ N, K }) {
+  const K4 = K / 4;
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> w: array<vec4<${HT()}>>;
+@group(0) @binding(1) var<storage, read> x: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+const N = ${N}u;
+const K4 = ${K4}u;
+@compute @workgroup_size(8, 8)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let n = gid.x;
+  let m = gid.y;
+  if (n >= N || m >= step.t) { return; }
+  var acc = 0.0;
+  let wbase = n * K4;
+  let xbase = m * K4;
+  for (var i = 0u; i < K4; i++) {
+    acc += dot(vec4<f32>(w[wbase + i]), x[xbase + i]);
+  }
+  y[m * N + n] = y[m * N + n] + acc;
+}`;
+}
+// RMSNorm over rows of length C: y = x / rms(x) * w. One workgroup per row.
+// Weights are pre-normalized by the converter ((1+w) folded in).
+export function rmsnorm({ C, eps }) {
+  return /* wgsl */ `
+${EF16()}
+@group(0) @binding(0) var<storage, read> x: array<f32>;
+@group(0) @binding(1) var<storage, read> w: array<${HT()}>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+const C = ${C}u;
+const EPS = ${eps};
+var<workgroup> partial: array<f32, 256>;
+@compute @workgroup_size(256)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  let base = wg.x * C;
+  var s = 0.0;
+  for (var i = li; i < C; i += 256u) {
+    let v = x[base + i];
+    s += v * v;
+  }
+  partial[li] = s;
+  workgroupBarrier();
+  var stride = 128u;
+  while (stride > 0u) {
+    if (li < stride) { partial[li] += partial[li + stride]; }
+    workgroupBarrier();
+    stride /= 2u;
+  }
+  let inv = inverseSqrt(partial[0] / f32(C) + EPS);
+  for (var i = li; i < C; i += 256u) {
+    y[base + i] = x[base + i] * inv * f32(w[i]);
+  }
+}`;
+}
+// y = silu(a) * b, elementwise over T*C
+export function swiglu({ C }) {
+  return /* wgsl */ `
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read> b: array<f32>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= step.t * ${C}u) { return; }
+  let v = a[i];
+  y[i] = v / (1.0 + exp(-v)) * b[i];
+}`;
+}
+// y += x, elementwise over T*C
+export function addInPlace({ C }) {
+  return /* wgsl */ `
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> x: array<f32>;
+@group(0) @binding(1) var<storage, read_write> y: array<f32>;
+@compute @workgroup_size(256)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+  let i = gid.x;
+  if (i >= step.t * ${C}u) { return; }
+  y[i] = y[i] + x[i];
+}`;
+}
+// out[t, :] = f32(emb[hist[pos0 + t], :]); C % 4 == 0.
+export function embedGather({ C }) {
+  const C4 = C / 4;
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> emb: array<vec4<${HT()}>>;
+@group(0) @binding(1) var<storage, read> hist: array<u32>;
+@group(0) @binding(2) var<storage, read_write> y: array<vec4<f32>>;
+const C4 = ${C4}u;
+@compute @workgroup_size(256)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  let t = wg.x;
+  if (t >= step.t) { return; }
+  let base = hist[step.pos0 + t] * C4;
+  for (var i = li; i < C4; i += 256u) {
+    y[t * C4 + i] = vec4<f32>(emb[base + i]);
+  }
+}`;
+}
+// ---------------------------------------------------------------------------
+// int4 wire format (from the lfm2 project): symmetric per-group along K,
+// byte-sliced nibbles (u32 word = weights q0..q7: low nibbles of bytes 0..3 =
+// q0..q3, high nibbles = q4..q7), biased by +8. f16 scales [N, K/group].
+// One vec4<u32> = 32 weights.
+// ---------------------------------------------------------------------------
+const Q8_UNPACK = /* wgsl */ `
+fn unpack_q8(w: u32) -> vec4<f32> {
+  return vec4<f32>(
+    f32(i32(w << 24u) >> 24u),
+    f32(i32(w << 16u) >> 24u),
+    f32(i32(w << 8u) >> 24u),
+    f32(i32(w) >> 24u));
+}`;
+const Q4_UNPACK = /* wgsl */ `
+fn q4_lo(w: u32) -> vec4<f32> {
+  return vec4<f32>(f32(w & 0xFu), f32((w >> 8u) & 0xFu), f32((w >> 16u) & 0xFu), f32((w >> 24u) & 0xFu)) - vec4<f32>(8.0);
+}
+fn q4_hi(w: u32) -> vec4<f32> {
+  return vec4<f32>(f32((w >> 4u) & 0xFu), f32((w >> 12u) & 0xFu), f32((w >> 20u) & 0xFu), f32((w >> 28u) & 0xFu)) - vec4<f32>(8.0);
+}
+fn q4_dot(w: u32, x0: vec4<f32>, x1: vec4<f32>) -> f32 {
+  return dot(q4_lo(w), x0) + dot(q4_hi(w), x1);
+}`;
+// f16-bitfield unpack: (nibble | 0x6400) bitcast to f16 is exactly 1024+q —
+// bit-exact with q4_lo/q4_hi at ~half the unpack ALU (lfm2 measured win on
+// hot layer matvecs; the cache-cold logits matvec prefers the plain unpack).
+const Q4_UNPACK_FT = /* wgsl */ `
+fn q4_dot(w: u32, x0: vec4<f32>, x1: vec4<f32>) -> f32 {
+  let p0 = bitcast<vec2<f16>>((w & 0x000F000Fu) | 0x64006400u);
+  let p1 = bitcast<vec2<f16>>(((w >> 8u) & 0x000F000Fu) | 0x64006400u);
+  let p2 = bitcast<vec2<f16>>(((w >> 4u) & 0x000F000Fu) | 0x64006400u);
+  let p3 = bitcast<vec2<f16>>(((w >> 12u) & 0x000F000Fu) | 0x64006400u);
+  let lo = vec4<f32>(f32(p0.x), f32(p1.x), f32(p0.y), f32(p1.y)) - vec4<f32>(1032.0);
+  let hi = vec4<f32>(f32(p2.x), f32(p3.x), f32(p2.y), f32(p3.y)) - vec4<f32>(1032.0);
+  return dot(lo, x0) + dot(hi, x1);
+}`;
+// q4 matvec, vec4<u32> loads, 2 rows per subgroup; epilogue "plain"/"add".
+// group 32: one scale per chunk. Subgroup-free fallback: 32-lane virtual warps.
+export function gemvQ4({ N, K, epilogue = "plain", group = 32, ft = true }) {
+  const KC = K / 32;
+  const SIDX = group === 32 ? "c" : "c / 2u";
+  const SGRP = K / group;
+  const store = (rowExpr, accExpr) =>
+    epilogue === "add" ? `y[${rowExpr}] = y[${rowExpr}] + ${accExpr};` : `y[${rowExpr}] = ${accExpr};`;
+  const head = SG
+    ? `  let lane = sgi;
+  let warp = li / sgs;
+  let wsz = sgs;
+  let row0 = (wg.x * (256u / sgs) + warp) * 2u;`
+    : `  let lane = li % 32u;
+  let warp = li / 32u;
+  let wsz = 32u;
+  let row0 = (wg.x * 8u + warp) * 2u;`;
+  const reduce = SG
+    ? `  a0 = subgroupAdd(a0);
+  a1 = subgroupAdd(a1);`
+    : `  red0[li] = a0;
+  red1[li] = a1;
+  workgroupBarrier();
+  var stride = 16u;
+  while (stride > 0u) {
+    if (lane < stride) {
+      red0[li] += red0[li + stride];
+      red1[li] += red1[li + stride];
+    }
+    workgroupBarrier();
+    stride /= 2u;
+  }
+  a0 = red0[li];
+  a1 = red1[li];`;
+  return /* wgsl */ `
+${EF16()}
+${SG ? "enable subgroups;" : ""}
+@group(0) @binding(0) var<storage, read> wq: array<vec4<u32>>;
+@group(0) @binding(1) var<storage, read> x: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+@group(0) @binding(4) var<storage, read> scales: array<${HT()}>;
+const N = ${N}u;
+const KC = ${KC}u;
+${ft && F16 ? Q4_UNPACK_FT : Q4_UNPACK}
+${SG ? "" : "var<workgroup> red0: array<f32, 256>;\nvar<workgroup> red1: array<f32, 256>;"}
+fn q4_chunk(w: vec4<u32>, c: u32) -> f32 {
+  return q4_dot(w.x, x[c * 8u], x[c * 8u + 1u]) +
+         q4_dot(w.y, x[c * 8u + 2u], x[c * 8u + 3u]) +
+         q4_dot(w.z, x[c * 8u + 4u], x[c * 8u + 5u]) +
+         q4_dot(w.w, x[c * 8u + 6u], x[c * 8u + 7u]);
+}
+@compute @workgroup_size(256)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_index) li: u32,
+${SG ? "  @builtin(subgroup_invocation_id) sgi: u32,\n  @builtin(subgroup_size) sgs: u32," : ""}
+) {
+${head}
+  let r = min(row0, N - 2u);
+  let wb0 = r * KC;
+  let wb1 = wb0 + KC;
+  let sb0 = r * ${SGRP}u;
+  let sb1 = sb0 + ${SGRP}u;
+  var a0 = 0.0;
+  var a1 = 0.0;
+  for (var c = lane; c < KC; c += wsz) {
+    a0 += f32(scales[sb0 + ${SIDX}]) * q4_chunk(wq[wb0 + c], c);
+    a1 += f32(scales[sb1 + ${SIDX}]) * q4_chunk(wq[wb1 + c], c);
+  }
+${reduce}
+  if (lane == 0u && row0 + 2u <= N) {
+    ${store("row0", "a0")}
+    ${store("row0 + 1u", "a1")}
+  }
+}`;
+}
+// q4 fused SwiGLU matvec
+export function gemvSwigluQ4({ N, K, group = 32 }) {
+  const KC = K / 32;
+  const SIDX = group === 32 ? "c" : "c / 2u";
+  const SGRP = K / group;
+  const head = SG
+    ? `  let lane = sgi;
+  let wsz = sgs;
+  let row = wg.x * (256u / sgs) + li / sgs;`
+    : `  let lane = li % 32u;
+  let wsz = 32u;
+  let row = wg.x * 8u + li / 32u;`;
+  const reduce = SG
+    ? `  a = subgroupAdd(a);
+  b = subgroupAdd(b);`
+    : `  red0[li] = a;
+  red1[li] = b;
+  workgroupBarrier();
+  var stride = 16u;
+  while (stride > 0u) {
+    if (lane < stride) {
+      red0[li] += red0[li + stride];
+      red1[li] += red1[li + stride];
+    }
+    workgroupBarrier();
+    stride /= 2u;
+  }
+  a = red0[li];
+  b = red1[li];`;
+  return /* wgsl */ `
+${EF16()}
+${SG ? "enable subgroups;" : ""}
+@group(0) @binding(0) var<storage, read> w1q: array<vec4<u32>>;
+@group(0) @binding(1) var<storage, read> x: array<vec4<f32>>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+@group(0) @binding(3) var<storage, read> w3q: array<vec4<u32>>;
+@group(0) @binding(4) var<storage, read> s1: array<${HT()}>;
+@group(0) @binding(5) var<storage, read> s3: array<${HT()}>;
+const N = ${N}u;
+const KC = ${KC}u;
+${F16 ? Q4_UNPACK_FT : Q4_UNPACK}
+${SG ? "" : "var<workgroup> red0: array<f32, 256>;\nvar<workgroup> red1: array<f32, 256>;"}
+fn q4_chunk(w: vec4<u32>, c: u32) -> f32 {
+  return q4_dot(w.x, x[c * 8u], x[c * 8u + 1u]) +
+         q4_dot(w.y, x[c * 8u + 2u], x[c * 8u + 3u]) +
+         q4_dot(w.z, x[c * 8u + 4u], x[c * 8u + 5u]) +
+         q4_dot(w.w, x[c * 8u + 6u], x[c * 8u + 7u]);
+}
+@compute @workgroup_size(256)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_index) li: u32,
+${SG ? "  @builtin(subgroup_invocation_id) sgi: u32,\n  @builtin(subgroup_size) sgs: u32," : ""}
+) {
+${head}
+  let r = min(row, N - 1u);
+  let wbase = r * KC;
+  let sbase = r * ${SGRP}u;
+  var a = 0.0;
+  var b = 0.0;
+  for (var c = lane; c < KC; c += wsz) {
+    let g = sbase + ${SIDX};
+    a += f32(s1[g]) * q4_chunk(w1q[wbase + c], c);
+    b += f32(s3[g]) * q4_chunk(w3q[wbase + c], c);
+  }
+${reduce}
+  if (lane == 0u && row < N) { y[row] = a / (1.0 + exp(-a)) * b; }
+}`;
+}
+// tiled prefill GEMM over quantized weights (q4/q8), 64x64 tiles, BK=32.
+// Unlike the lfm2 original: guards for N not divisible by 64 (our fused
+// delta in_proj is 8224 rows) and an "add" epilogue for fused residuals.
+export function gemmTiledQ({ N, K, dtype = "q4", group = 32, xstride = K, epilogue = "plain" }) {
+  const q4 = dtype === "q4";
+  const wordsPerCol = q4 ? 4 : 8;
+  const weightsPerWord = q4 ? 8 : 4;
+  const KWROW = K / weightsPerWord;
+  const SGRP = K / group;
+  const store = (dst, acc) => (epilogue === "add" ? `y[${dst}] = y[${dst}] + ${acc};` : `y[${dst}] = ${acc};`);
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> wq: array<u32>;
+@group(0) @binding(1) var<storage, read> x: array<f32>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+@group(0) @binding(4) var<storage, read> scales: array<${HT()}>;
+const N = ${N}u;
+const K = ${K}u;
+${Q8_UNPACK}
+${Q4_UNPACK}
+var<workgroup> As: array<array<f32, 32>, 64>;      // [m][k]
+var<workgroup> Bs: array<array<f32, 65>, 32>;      // [k][n] (+1 pad col)
+@compute @workgroup_size(16, 16)
+fn main(
+  @builtin(workgroup_id) wg: vec3<u32>,
+  @builtin(local_invocation_id) lid: vec3<u32>,
+  @builtin(local_invocation_index) li: u32,
+) {
+  let n0 = wg.x * 64u;
+  let m0 = wg.y * 64u;
+  var acc: array<array<f32, 4>, 4>;
+  for (var k0 = 0u; k0 < K; k0 += 32u) {
+    for (var e = li; e < 512u; e += 256u) {
+      let m = e / 8u;
+      let k = (e % 8u) * 4u;
+      let mm = m0 + m;
+      if (mm < step.t) {
+        let base = mm * ${xstride}u + k0 + k;
+        As[m][k] = x[base];
+        As[m][k + 1u] = x[base + 1u];
+        As[m][k + 2u] = x[base + 2u];
+        As[m][k + 3u] = x[base + 3u];
+      } else {
+        As[m][k] = 0.0; As[m][k + 1u] = 0.0; As[m][k + 2u] = 0.0; As[m][k + 3u] = 0.0;
+      }
+    }
+    for (var e = li; e < ${64 * wordsPerCol}u; e += 256u) {
+      let n = e / ${wordsPerCol}u;
+      let kw = e % ${wordsPerCol}u;
+      let col = min(n0 + n, N - 1u);
+      let kabs = k0 + kw * ${weightsPerWord}u;
+      let s = f32(scales[col * ${SGRP}u + kabs / ${group}u]);
+      let w = wq[col * ${KWROW}u + kabs / ${weightsPerWord}u];
+      let kk = kw * ${weightsPerWord}u;
+${
+  q4
+    ? `      let lo = q4_lo(w) * s;
+      let hi = q4_hi(w) * s;
+      Bs[kk][n] = lo.x; Bs[kk + 1u][n] = lo.y; Bs[kk + 2u][n] = lo.z; Bs[kk + 3u][n] = lo.w;
+      Bs[kk + 4u][n] = hi.x; Bs[kk + 5u][n] = hi.y; Bs[kk + 6u][n] = hi.z; Bs[kk + 7u][n] = hi.w;`
+    : `      let v = unpack_q8(w) * s;
+      Bs[kk][n] = v.x; Bs[kk + 1u][n] = v.y; Bs[kk + 2u][n] = v.z; Bs[kk + 3u][n] = v.w;`
+}
+    }
+    workgroupBarrier();
+    for (var k = 0u; k < 32u; k++) {
+      let a0 = As[lid.y * 4u][k];
+      let a1 = As[lid.y * 4u + 1u][k];
+      let a2 = As[lid.y * 4u + 2u][k];
+      let a3 = As[lid.y * 4u + 3u][k];
+      let b0 = Bs[k][lid.x * 4u];
+      let b1 = Bs[k][lid.x * 4u + 1u];
+      let b2 = Bs[k][lid.x * 4u + 2u];
+      let b3 = Bs[k][lid.x * 4u + 3u];
+      acc[0][0] += a0 * b0; acc[0][1] += a0 * b1; acc[0][2] += a0 * b2; acc[0][3] += a0 * b3;
+      acc[1][0] += a1 * b0; acc[1][1] += a1 * b1; acc[1][2] += a1 * b2; acc[1][3] += a1 * b3;
+      acc[2][0] += a2 * b0; acc[2][1] += a2 * b1; acc[2][2] += a2 * b2; acc[2][3] += a2 * b3;
+      acc[3][0] += a3 * b0; acc[3][1] += a3 * b1; acc[3][2] += a3 * b2; acc[3][3] += a3 * b3;
+    }
+    workgroupBarrier();
+  }
+  for (var i = 0u; i < 4u; i++) {
+    let m = m0 + lid.y * 4u + i;
+    if (m >= step.t) { continue; }
+    for (var j = 0u; j < 4u; j++) {
+      let n = n0 + lid.x * 4u + j;
+      if (n < N) { ${store("m * N + n", "acc[i][j]")} }
+    }
+  }
+}`;
+}
+// q4 embedding gather: out[t, :] = dequant(emb[hist[pos0+t], :]).
+// One workgroup per token; each thread unpacks one u32 word (8 weights).
+export function embedGatherQ4({ C, group = 32 }) {
+  const CW = C / 8;      // u32 words per row
+  const SGRP = C / group;
+  return /* wgsl */ `
+${EF16()}
+${STEP_UNIFORM}
+@group(0) @binding(0) var<storage, read> emb: array<u32>;
+@group(0) @binding(1) var<storage, read> hist: array<u32>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+@group(0) @binding(4) var<storage, read> scales: array<${HT()}>;
+const CW = ${CW}u;
+${Q4_UNPACK}
+@compute @workgroup_size(256)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  let t = wg.x;
+  if (t >= step.t) { return; }
+  let row = hist[step.pos0 + t];
+  for (var wi = li; wi < CW; wi += 256u) {
+    let s = f32(scales[row * ${SGRP}u + (wi * 8u) / ${group}u]);
+    let w = emb[row * CW + wi];
+    let lo = q4_lo(w) * s;
+    let hi = q4_hi(w) * s;
+    let base = t * ${C}u + wi * 8u;
+    y[base] = lo.x; y[base + 1u] = lo.y; y[base + 2u] = lo.z; y[base + 3u] = lo.w;
+    y[base + 4u] = hi.x; y[base + 5u] = hi.y; y[base + 6u] = hi.z; y[base + 7u] = hi.w;
+  }
+}`;
+}
+// ---------------------------------------------------------------------------
+// Cascaded greedy argmax (lfm2 design, full-width "direct" mode only —
+// partial coarse is provably unsafe without tail-norm bounds). The exact
+// logits are computed by the normal logits matvec; argmaxSliceTop picks the
+// top-2 of every 64-token slice (the global winner provably survives), and
+// argmaxCandAdvance reduces candidates straight off the exact logits, writes
+// hist[pos] and bumps the step buffer — GPU-resident decode, no CPU readback.
+// ---------------------------------------------------------------------------
+// top-2 of each 64-wide slice of the logits -> 2 candidates per slice.
+// dispatch: V/64 workgroups.
+export function argmaxSliceTop({ V }) {
+  return /* wgsl */ `
+@group(0) @binding(0) var<storage, read> logits: array<f32>;
+@group(0) @binding(2) var<storage, read_write> cidxs: array<u32>;
+var<workgroup> vals: array<f32, 64>;
+var<workgroup> idxs: array<u32, 64>;
+var<workgroup> win: u32;
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wg: vec3<u32>, @builtin(local_invocation_index) li: u32) {
+  let idx = wg.x * 64u + li;
+  let v = logits[idx];
+  for (var round = 0u; round < 2u; round++) {
+    vals[li] = v;
+    idxs[li] = idx;
+    if (round == 1u && idx == win) { vals[li] = -3.4e38; }
+    workgroupBarrier();
+    var stride = 32u;
+    while (stride > 0u) {
+      if (li < stride) {
+        let ov = vals[li + stride];
+        let oi = idxs[li + stride];
+        if (ov > vals[li] || (ov == vals[li] && oi < idxs[li])) {
+          vals[li] = ov;
+          idxs[li] = oi;
+        }
+      }
+      workgroupBarrier();
+      stride /= 2u;
+    }
+    if (li == 0u) {
+      cidxs[wg.x * 2u + round] = idxs[0];
+      win = idxs[0];
+    }
+    workgroupBarrier();
+  }
+}`;
+}
+// final reduce over NC candidates read straight from the exact logits
+// (ties -> lowest token id, matching torch.argmax) + GPU-side advance:
+// hist[pos0 + t] = winner, step becomes {1, pos0 + t}. dispatch: 1
+export function argmaxCandAdvance({ NC }) {
+  return /* wgsl */ `
+struct Step { t: u32, pos0: u32, sstate: u32, acc: u32 }
+@group(0) @binding(0) var<storage, read> logits: array<f32>;
+@group(0) @binding(1) var<storage, read_write> step: Step;
+@group(0) @binding(2) var<storage, read_write> hist: array<u32>;
+@group(0) @binding(3) var<storage, read> cidxs: array<u32>;
+var<workgroup> vals: array<f32, 256>;
+var<workgroup> idxs: array<u32, 256>;
+@compute @workgroup_size(256)
+fn main(@builtin(local_invocation_index) li: u32) {
+  var bv = -3.4e38;
+  var bi = 0xffffffffu;
+  for (var c = li; c < ${NC}u; c += 256u) {
+    let ti = cidxs[c];
+    let v = logits[ti];
+    if (v > bv || (v == bv && ti < bi)) { bv = v; bi = ti; }
+  }
+  vals[li] = bv;
+  idxs[li] = bi;
+  workgroupBarrier();
+  var stride = 128u;
+  while (stride > 0u) {
+    if (li < stride) {
+      let ov = vals[li + stride];
+      let oi = idxs[li + stride];
+      if (ov > vals[li] || (ov == vals[li] && oi < idxs[li])) {
+        vals[li] = ov;
+        idxs[li] = oi;
+      }
+    }
+    workgroupBarrier();
+    stride /= 2u;
+  }
+  if (li == 0u) {
+    let next = step.pos0 + step.t;
+    hist[next] = idxs[0];
+    step.t = 1u;
+    step.pos0 = next;
+  }
+}`;
+}

web/main.js ADDED Viewed

	@@ -0,0 +1,171 @@

+import { initDevice } from "/web/gpu.js";
+import { loadModel } from "/web/model.js";
+import { Tokenizer } from "/web/tokenizer.js";
+import { resolvePaths } from "/web/config.js";
+const $ = (id) => document.getElementById(id);
+const CHUNK = 8;
+const MAX_NEW = 768;
+const log = (text, cls = "") => {
+  const el = document.createElement("div");
+  el.textContent = text;
+  if (cls) el.className = cls;
+  $("log").appendChild(el);
+  while ($("log").children.length > 8) $("log").firstChild.remove();
+};
+let model, tok, EOS, EOT;
+const history = [];
+const QUANT = new URLSearchParams(location.search).get("quant") ?? "q4";
+const BLOB_GB = { f16: 3.76, q4: 1.07 };
+async function boot() {
+  try {
+    log("requesting a WebGPU adapter…");
+    const adapter = await navigator.gpu?.requestAdapter();
+    if (!adapter) throw new Error("WebGPU is not available in this browser");
+    const info = adapter.info;
+    $("gpuname").textContent = `${info.vendor} ${info.architecture} · ${QUANT}`.trim();
+    const device = await initDevice();
+    const paths = await resolvePaths();
+    log(`weights source: ${paths.source}`);
+    log("loading the tokenizer (248k BPE)…");
+    const tokP = Tokenizer.load(paths.tokenizer);
+    const GB = BLOB_GB[QUANT] ?? 1.07;
+    $("dlmb").textContent = `${GB} GB`;
+    log(`loading weights: ${GB} GB ${QUANT}…`);
+    const t0 = performance.now();
+    model = await loadModel(device, {
+      base: paths.weights,
+      variant: QUANT,
+      onProgress: (got, total) => {
+        const f = got / total;
+        $("barfill").style.transform = `scaleX(${f * 0.9})`;
+        $("dlpct").textContent = `· ${Math.round(f * 100)}%`;
+      },
+    });
+    $("dlnote").textContent = model.weightsFromCache
+      ? "model loaded from the browser cache ✓"
+      : "model downloaded ✓ — cached for next time";
+    tok = await tokP;
+    EOS = tok.eos;
+    EOT = tok.eot;
+    $("barfill").style.transform = "scaleX(0.95)";
+    log(`weights in GPU buffers in ${((performance.now() - t0) / 1000).toFixed(1)} s`, "ok");
+    log("warmup: compiling pipelines…");
+    model.startGenerate(tok.encode("warmup"));
+    await model.decodeChunk(1);
+    model.reset();
+    $("barfill").style.transform = "scaleX(1)";
+    log("ready.", "ok");
+    setTimeout(() => {
+      $("boot").classList.add("done");
+      $("input").disabled = false;
+      $("send").disabled = false;
+      $("input").focus();
+      const demo = new URLSearchParams(location.search).get("demo");
+      if (demo !== null) {
+        $("input").value = demo || "What is the capital of France?";
+        $("form").requestSubmit();
+      } else {
+        window.__done = { ok: true, stage: "loaded" };
+      }
+    }, 400);
+  } catch (e) {
+    log(String(e.message ?? e), "err");
+    window.__done = { ok: false, error: String(e.message ?? e) };
+  }
+}
+function addMsg(role) {
+  const el = document.createElement("div");
+  el.className = `msg ${role}`;
+  el.innerHTML = `<div class="who">${role === "user" ? "you" : "qwen3.5"}</div><div class="body"></div>`;
+  $("msgs").appendChild(el);
+  return el.querySelector(".body");
+}
+let busy = false;
+async function generate(userText) {
+  busy = true;
+  $("send").disabled = true;
+  $("input").disabled = true;
+  addMsg("user").textContent = userText;
+  history.push({ role: "user", content: userText });
+  const body = addMsg("assistant");
+  const caret = document.createElement("span");
+  caret.className = "caret";
+  body.appendChild(caret);
+  // each turn re-prefills the whole conversation: keeps DeltaNet/KV state
+  // exactly right with no post-EOS pollution
+  const ids = tok.applyChatTemplate(history);
+  const t0 = performance.now();
+  model.startGenerate(ids);
+  await model.device.queue.onSubmittedWorkDone();
+  const prefillMs = performance.now() - t0;
+  let outIds = [];
+  let text = "";
+  const t1 = performance.now();
+  const budget = Math.min(MAX_NEW, model.maxSeq - ids.length - CHUNK - 4);
+  let done = 0;
+  while (done < budget) {
+    const n = Math.min(CHUNK, budget - done);
+    const gen = await model.decodeChunk(n);
+    done += n;
+    const stop = gen.findIndex((id) => id === EOS || id === EOT);
+    outIds = stop >= 0 ? gen.slice(0, stop) : gen;
+    text = tok.decode(outIds);
+    body.textContent = text;
+    body.appendChild(caret);
+    const dt = (performance.now() - t1) / 1000;
+    $("toks").textContent = (done / dt).toFixed(1);
+    $("chat").scrollTop = $("chat").scrollHeight;
+    if (stop >= 0) break;
+  }
+  const decodeS = (performance.now() - t1) / 1000;
+  caret.remove();
+  history.push({ role: "assistant", content: text });
+  const stats = document.createElement("div");
+  stats.className = "stats";
+  stats.textContent =
+    `prefill ${ids.length} tok · ${prefillMs.toFixed(0)} ms — ` +
+    `decode ${done} tok · ${(done / decodeS).toFixed(1)} tok/s`;
+  body.parentElement.appendChild(stats);
+  $("chat").scrollTop = $("chat").scrollHeight;
+  busy = false;
+  $("send").disabled = false;
+  $("input").disabled = false;
+  $("input").focus();
+  window.__done = { ok: true, stage: "generated", text, tok_s: +(done / decodeS).toFixed(1) };
+}
+$("form").addEventListener("submit", (e) => {
+  e.preventDefault();
+  const text = $("input").value.trim();
+  if (!text || busy) return;
+  $("input").value = "";
+  $("input").style.height = "auto";
+  generate(text);
+});
+$("input").addEventListener("keydown", (e) => {
+  if (e.key === "Enter" && !e.shiftKey) {
+    e.preventDefault();
+    $("form").requestSubmit();
+  }
+});
+$("input").addEventListener("input", (e) => {
+  e.target.style.height = "auto";
+  e.target.style.height = e.target.scrollHeight + "px";
+});
+boot();

web/model.js ADDED Viewed

	@@ -0,0 +1,440 @@

+// Qwen3.5-2B WebGPU runtime: weight loader + hybrid forward
+// (18x gated DeltaNet + 6x gated attention), f16 weights / f32 activations.
+import * as K from "./kernels.js";
+import { createBuffer, makePipeline } from "./gpu.js";
+export const CFG = {
+  L: 24, C: 2048, V: 248320, FF: 6144, EPS: 1e-6,
+  H: 16, DK: 128, DV: 128,              // DeltaNet
+  NH: 8, NKV: 2, HD: 256, RD: 64,       // attention
+  ROPE_THETA: 1e7,
+  DELTA_IN: 8224,                       // qkv 6144 | z 2048 | b 16 | a 16
+  CONV_C: 6144,
+  ATTN_IN: 5120,                        // per-head (q|gate)*8 | k*2 | v*2
+};
+const LAYER_TYPES = Array.from({ length: 24 }, (_, i) => (i % 4 === 3 ? "attn" : "delta"));
+// Stream the blob straight into per-GPU-buffer segments: the f16 blob (3.76
+// GB) does not fit in a single ArrayBuffer, so JS never holds more than one
+// segment (max ~1 GB, the tied embedding) at a time. q4 tensors contribute
+// two segments: packed weights and f16 scales (key "<name>::scales").
+export async function loadModel(device, { base = "/web/weights/", variant = "q4", onProgress } = {}) {
+  const manifest = await (await fetch(`${base}manifest.${variant}.json`)).json();
+  const segments = [];
+  for (const [name, t] of Object.entries(manifest.tensors)) {
+    segments.push({ key: name, offset: t.offset, bytes: t.bytes });
+    if (t.scales_offset !== undefined) {
+      segments.push({ key: name + "::scales", offset: t.scales_offset, bytes: t.scales_bytes });
+    }
+  }
+  segments.sort((a, b) => a.offset - b.offset);
+  const url = `${base}qwen35-2b.${variant}.bin`;
+  // Cache API: q4 (~1 GB) is cached for instant reloads; the f16 blob is too
+  // big for most quotas and is fetched fresh every time.
+  let resp = null;
+  let fromCache = false;
+  let cache = null;
+  const cacheable = variant !== "f16" && typeof caches !== "undefined";
+  if (cacheable) {
+    try {
+      cache = await caches.open("qwen35-weights");
+      const hit = await cache.match(url);
+      if (hit) { resp = hit; fromCache = true; }
+    } catch {}
+  }
+  if (!resp) {
+    resp = await fetch(url);
+    if (resp.ok && cache) {
+      try { await cache.put(url, resp.clone()); } catch {}
+    }
+  }
+  if (!resp.ok) throw new Error(`fetch blob: ${resp.status}`);
+  const total = Number(resp.headers.get("content-length") ?? manifest.total_bytes);
+  const reader = resp.body.getReader();
+  const tensors = {};
+  let idx = 0;
+  let cur = null;      // Uint8Array being filled
+  let curFill = 0;
+  let pos = 0;         // absolute blob position
+  const finish = () => {
+    const seg = segments[idx];
+    const buf = createBuffer(device, seg.bytes);
+    device.queue.writeBuffer(buf, 0, cur.buffer, 0, seg.bytes);
+    tensors[seg.key] = buf;
+    cur = null;
+    curFill = 0;
+    idx++;
+  };
+  for (;;) {
+    const { done, value } = await reader.read();
+    if (done) break;
+    let off = 0;
+    while (off < value.byteLength && idx < segments.length) {
+      const seg = segments[idx];
+      if (pos < seg.offset) {                     // alignment gap
+        const skip = Math.min(seg.offset - pos, value.byteLength - off);
+        pos += skip; off += skip;
+        continue;
+      }
+      if (!cur) cur = new Uint8Array(seg.bytes);
+      const take = Math.min(seg.bytes - curFill, value.byteLength - off);
+      cur.set(value.subarray(off, off + take), curFill);
+      curFill += take; pos += take; off += take;
+      if (curFill === seg.bytes) finish();
+    }
+    pos += value.byteLength - off;
+    onProgress?.(pos, total);
+  }
+  if (idx < segments.length) throw new Error(`blob truncated: ${idx}/${segments.length} segments`);
+  const model = new Qwen35(device, manifest, tensors);
+  model.weightsFromCache = fromCache;
+  return model;
+}
+export class Qwen35 {
+  constructor(device, manifest, w) {
+    this.device = device;
+    this.manifest = manifest;
+    this.w = w;
+    this.q4 = manifest.quant === "q4";
+    this.maxSeq = 4096;
+    this.maxT = 512;      // prefill segment length
+    this.pos = 0;
+    this.hist = [];       // token ids (CPU mirror of bufHist)
+    this.#alloc();
+    this.#pipelines();
+    this.#bindGroups();
+  }
+  #alloc() {
+    const d = this.device;
+    const { C, V, FF, H, DK, DV, NH, NKV, HD, DELTA_IN, CONV_C, ATTN_IN, L } = CFG;
+    const T = this.maxT;
+    const hb = 2; // f16 KV cache
+    this.bufHist = createBuffer(d, this.maxSeq * 4);
+    this.bufStep = createBuffer(d, 16);
+    this.bufX = createBuffer(d, T * C * 4);
+    this.bufXn = createBuffer(d, T * C * 4);
+    this.bufAct = createBuffer(d, T * DELTA_IN * 4);
+    this.bufQkv = createBuffer(d, T * CONV_C * 4);
+    this.bufG = createBuffer(d, T * H * 4);
+    this.bufBeta = createBuffer(d, T * H * 4);
+    this.bufCore = createBuffer(d, T * H * DV * 4);
+    this.bufGn = createBuffer(d, T * H * DV * 4);
+    this.bufQout = createBuffer(d, T * NH * HD * 4);
+    this.bufAttnOut = createBuffer(d, T * NH * HD * 4);
+    this.bufMlpA = createBuffer(d, T * FF * 4);
+    this.bufMlpB = createBuffer(d, T * FF * 4);
+    this.bufMlpC = createBuffer(d, T * FF * 4);
+    this.bufLoc = createBuffer(d, Math.ceil(T / 64) * H * K.DELTA_BLK * 4);
+    this.bufLogits = createBuffer(d, V * 4);
+    this.bufCidxs = createBuffer(d, (V / 64) * 2 * 4);
+    this.convState = [];
+    this.recState = [];
+    this.kCache = [];
+    this.vCache = [];
+    for (let i = 0; i < L; i++) {
+      if (LAYER_TYPES[i] === "delta") {
+        this.convState[i] = createBuffer(d, CONV_C * 4 * 4);
+        this.recState[i] = createBuffer(d, H * DK * DV * 4);
+      } else {
+        this.kCache[i] = createBuffer(d, this.maxSeq * NKV * HD * hb);
+        this.vCache[i] = createBuffer(d, this.maxSeq * NKV * HD * hb);
+      }
+    }
+  }
+  #pipelines() {
+    const d = this.device;
+    const { C, V, FF, H, DK, DV, NH, NKV, HD, RD, EPS, ROPE_THETA, DELTA_IN, CONV_C, ATTN_IN } = CFG;
+    const mk = (code, label) => makePipeline(d, code, label);
+    if (this.q4) {
+      this.p = {
+        embed: mk(K.embedGatherQ4({ C }), "embedQ4"),
+        norm: mk(K.rmsnorm({ C, eps: EPS }), "rmsnorm"),
+        conv: mk(K.convQkv({ C: CONV_C, xrow: DELTA_IN }), "convQkv"),
+        gates: mk(K.deltaGates({ H, row: DELTA_IN, off: DELTA_IN - 2 * H }), "deltaGates"),
+        dstep: mk(K.deltaStep({ H, DK, DV }), "deltaStep"),
+        dlocal: mk(K.deltaChunkLocal({ H, DK, DV }), "deltaChunkLocal"),
+        dscan: mk(K.deltaChunkScan({ H, DK, DV }), "deltaChunkScan"),
+        gnorm: mk(K.gatedNorm({ H, DV, eps: EPS, zrow: DELTA_IN, zoff: CONV_C }), "gatedNorm"),
+        rope: mk(K.qkvRope({ NH, NKV, HD, RD, eps: EPS, ropeTheta: ROPE_THETA, maxSeq: this.maxSeq }), "qkvRope"),
+        attn: mk(K.attention({ NH, NKV, HD, maxSeq: this.maxSeq }), "attention"),
+        gateMul: mk(K.gateMul({ NH, NKV, HD }), "gateMul"),
+        swiglu: mk(K.swiglu({ C: FF }), "swiglu"),
+        gvDeltaIn: mk(K.gemvQ4({ N: DELTA_IN, K: C }), "gemvQ4.deltaIn"),
+        gvAttnIn: mk(K.gemvQ4({ N: ATTN_IN, K: C }), "gemvQ4.attnIn"),
+        gvOut: mk(K.gemvQ4({ N: C, K: C, epilogue: "add" }), "gemvQ4add.out"),
+        gvSwiglu: mk(K.gemvSwigluQ4({ N: FF, K: C }), "gemvSwigluQ4"),
+        gvDown: mk(K.gemvQ4({ N: C, K: FF, epilogue: "add" }), "gemvQ4add.down"),
+        gvLogits: mk(K.gemvQ4({ N: V, K: C, ft: false }), "gemvQ4.logits"),
+        gmDeltaIn: mk(K.gemmTiledQ({ N: DELTA_IN, K: C }), "gemmQ4.deltaIn"),
+        gmAttnIn: mk(K.gemmTiledQ({ N: ATTN_IN, K: C }), "gemmQ4.attnIn"),
+        gmOut: mk(K.gemmTiledQ({ N: C, K: C, epilogue: "add" }), "gemmQ4add.out"),
+        gmFF: mk(K.gemmTiledQ({ N: FF, K: C }), "gemmQ4.ff"),
+        gmDown: mk(K.gemmTiledQ({ N: C, K: FF, epilogue: "add" }), "gemmQ4add.down"),
+        sliceTop: mk(K.argmaxSliceTop({ V }), "argmaxSliceTop"),
+        candAdv: mk(K.argmaxCandAdvance({ NC: (V / 64) * 2 }), "argmaxCandAdvance"),
+      };
+      return;
+    }
+    this.p = {
+      embed: mk(K.embedGather({ C }), "embed"),
+      norm: mk(K.rmsnorm({ C, eps: EPS }), "rmsnorm"),
+      conv: mk(K.convQkv({ C: CONV_C, xrow: DELTA_IN }), "convQkv"),
+      gates: mk(K.deltaGates({ H, row: DELTA_IN, off: DELTA_IN - 2 * H }), "deltaGates"),
+      dstep: mk(K.deltaStep({ H, DK, DV }), "deltaStep"),
+      dlocal: mk(K.deltaChunkLocal({ H, DK, DV }), "deltaChunkLocal"),
+      dscan: mk(K.deltaChunkScan({ H, DK, DV }), "deltaChunkScan"),
+      gnorm: mk(K.gatedNorm({ H, DV, eps: EPS, zrow: DELTA_IN, zoff: CONV_C }), "gatedNorm"),
+      rope: mk(K.qkvRope({ NH, NKV, HD, RD, eps: EPS, ropeTheta: ROPE_THETA, maxSeq: this.maxSeq }), "qkvRope"),
+      attn: mk(K.attention({ NH, NKV, HD, maxSeq: this.maxSeq }), "attention"),
+      gateMul: mk(K.gateMul({ NH, NKV, HD }), "gateMul"),
+      swiglu: mk(K.swiglu({ C: FF }), "swiglu"),
+      // decode matvecs
+      gvDeltaIn: mk(K.gemv({ N: DELTA_IN, K: C }), "gemv.deltaIn"),
+      gvAttnIn: mk(K.gemv({ N: ATTN_IN, K: C }), "gemv.attnIn"),
+      gvOut: mk(K.gemvAdd({ N: C, K: C }), "gemvAdd.out"),
+      gvSwiglu: mk(K.gemvSwiglu({ N: FF, K: C }), "gemvSwiglu"),
+      gvDown: mk(K.gemvAdd({ N: C, K: FF }), "gemvAdd.down"),
+      gvLogits: mk(K.gemv({ N: V, K: C }), "gemv.logits"),
+      // prefill matmuls
+      gmDeltaIn: mk(K.gemm({ N: DELTA_IN, K: C }), "gemm.deltaIn"),
+      gmAttnIn: mk(K.gemm({ N: ATTN_IN, K: C }), "gemm.attnIn"),
+      gmOut: mk(K.gemmAdd({ N: C, K: C }), "gemmAdd.out"),
+      gmFF: mk(K.gemm({ N: FF, K: C }), "gemm.ff"),
+      gmDown: mk(K.gemmAdd({ N: C, K: FF }), "gemmAdd.down"),
+      sliceTop: mk(K.argmaxSliceTop({ V }), "argmaxSliceTop"),
+      candAdv: mk(K.argmaxCandAdvance({ NC: (V / 64) * 2 }), "argmaxCandAdvance"),
+    };
+  }
+  #bind(pipeline, entries) {
+    return this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: Object.entries(entries).map(([binding, r]) => ({
+        binding: Number(binding),
+        resource: r instanceof GPUBuffer ? { buffer: r } : r,
+      })),
+    });
+  }
+  #bindGroups() {
+    const { C } = CFG;
+    const w = this.w;
+    const p = this.p;
+    const step = this.bufStep;
+    // q4 linears carry scales at binding 4 (5 for the swiglu up-matrix)
+    const sc = (name, b = 4) => (this.q4 ? { [b]: w[name + "::scales"] } : {});
+    this.bgEmbed = this.#bind(p.embed, { 0: w["embed_tokens.weight"], 1: this.bufHist, 2: this.bufX, 15: step, ...sc("embed_tokens.weight") });
+    this.bgFinalNorm = this.#bind(p.norm, { 0: this.bufX, 1: w["norm.weight"], 2: this.bufXn });
+    this.bgLogits = (tOff) => this.#bind(p.gvLogits, {
+      0: w["embed_tokens.weight"],
+      1: { buffer: this.bufXn, offset: tOff * C * 4, size: C * 4 },
+      2: this.bufLogits,
+      ...sc("embed_tokens.weight"),
+    });
+    this.bgSliceTop = this.#bind(p.sliceTop, { 0: this.bufLogits, 2: this.bufCidxs });
+    this.bgCandAdv = this.#bind(p.candAdv, { 0: this.bufLogits, 1: step, 2: this.bufHist, 3: this.bufCidxs });
+    this.layers = [];
+    for (let i = 0; i < CFG.L; i++) {
+      const pre = `layers.${i}.`;
+      const lg = {
+        type: LAYER_TYPES[i],
+        bgNorm1: this.#bind(p.norm, { 0: this.bufX, 1: w[pre + "input_layernorm.weight"], 2: this.bufXn }),
+        bgNorm2: this.#bind(p.norm, { 0: this.bufX, 1: w[pre + "post_attention_layernorm.weight"], 2: this.bufXn }),
+      };
+      if (lg.type === "delta") {
+        const q = pre + "linear_attn.";
+        lg.bgIn = this.#bind(p.gvDeltaIn, { 0: w[q + "in_proj.weight"], 1: this.bufXn, 2: this.bufAct, ...sc(q + "in_proj.weight") });
+        lg.bgInP = this.#bind(p.gmDeltaIn, { 0: w[q + "in_proj.weight"], 1: this.bufXn, 2: this.bufAct, 15: step, ...sc(q + "in_proj.weight") });
+        lg.bgConv = this.#bind(p.conv, { 0: this.bufAct, 1: w[q + "conv1d.weight"], 2: this.convState[i], 3: this.bufQkv, 15: step });
+        lg.bgGates = this.#bind(p.gates, { 0: this.bufAct, 1: w[q + "gate_params"], 2: this.bufG, 3: this.bufBeta, 15: step });
+        lg.bgStep = this.#bind(p.dstep, { 0: this.bufQkv, 1: this.bufG, 2: this.bufBeta, 3: this.recState[i], 4: this.bufCore, 15: step });
+        lg.bgLocal = this.#bind(p.dlocal, { 0: this.bufQkv, 1: this.bufG, 2: this.bufBeta, 3: this.bufLoc, 15: step });
+        lg.bgScan = this.#bind(p.dscan, { 0: this.bufLoc, 1: this.recState[i], 2: this.bufCore, 15: step });
+        lg.bgGnorm = this.#bind(p.gnorm, { 0: this.bufCore, 1: this.bufAct, 2: w[q + "norm.weight"], 3: this.bufGn });
+        lg.bgOut = this.#bind(p.gvOut, { 0: w[q + "out_proj.weight"], 1: this.bufGn, 2: this.bufX, ...sc(q + "out_proj.weight") });
+        lg.bgOutP = this.#bind(p.gmOut, { 0: w[q + "out_proj.weight"], 1: this.bufGn, 2: this.bufX, 15: step, ...sc(q + "out_proj.weight") });
+      } else {
+        const q = pre + "self_attn.";
+        lg.bgIn = this.#bind(p.gvAttnIn, { 0: w[q + "in_proj.weight"], 1: this.bufXn, 2: this.bufAct, ...sc(q + "in_proj.weight") });
+        lg.bgInP = this.#bind(p.gmAttnIn, { 0: w[q + "in_proj.weight"], 1: this.bufXn, 2: this.bufAct, 15: step, ...sc(q + "in_proj.weight") });
+        lg.bgRope = this.#bind(p.rope, { 0: this.bufAct, 3: w[q + "q_norm.weight"], 4: w[q + "k_norm.weight"], 5: this.bufQout, 6: this.kCache[i], 7: this.vCache[i], 15: step });
+        lg.bgAttn = this.#bind(p.attn, { 0: this.bufQout, 1: this.kCache[i], 2: this.vCache[i], 3: this.bufAttnOut, 15: step });
+        lg.bgGateMul = this.#bind(p.gateMul, { 0: this.bufAct, 1: this.bufAttnOut, 15: step });
+        lg.bgOut = this.#bind(p.gvOut, { 0: w[q + "o_proj.weight"], 1: this.bufAttnOut, 2: this.bufX, ...sc(q + "o_proj.weight") });
+        lg.bgOutP = this.#bind(p.gmOut, { 0: w[q + "o_proj.weight"], 1: this.bufAttnOut, 2: this.bufX, 15: step, ...sc(q + "o_proj.weight") });
+      }
+      const m = pre + "mlp.";
+      lg.bgSwiglu = this.#bind(p.gvSwiglu, { 0: w[m + "gate_proj.weight"], 1: this.bufXn, 2: this.bufMlpA, 3: w[m + "up_proj.weight"], ...sc(m + "gate_proj.weight"), ...sc(m + "up_proj.weight", 5) });
+      lg.bgDown = this.#bind(p.gvDown, { 0: w[m + "down_proj.weight"], 1: this.bufMlpA, 2: this.bufX, ...sc(m + "down_proj.weight") });
+      lg.bgGateP = this.#bind(p.gmFF, { 0: w[m + "gate_proj.weight"], 1: this.bufXn, 2: this.bufMlpA, 15: step, ...sc(m + "gate_proj.weight") });
+      lg.bgUpP = this.#bind(p.gmFF, { 0: w[m + "up_proj.weight"], 1: this.bufXn, 2: this.bufMlpB, 15: step, ...sc(m + "up_proj.weight") });
+      lg.bgSwP = this.#bind(p.swiglu, { 0: this.bufMlpA, 1: this.bufMlpB, 2: this.bufMlpC, 15: step });
+      lg.bgDownP = this.#bind(p.gmDown, { 0: w[m + "down_proj.weight"], 1: this.bufMlpC, 2: this.bufX, 15: step, ...sc(m + "down_proj.weight") });
+      this.layers.push(lg);
+    }
+  }
+  reset() {
+    const { CONV_C, H, DK, DV, NKV, HD } = CFG;
+    this.pos = 0;
+    this.hist = [];
+    for (let i = 0; i < CFG.L; i++) {
+      if (LAYER_TYPES[i] === "delta") {
+        this.device.queue.writeBuffer(this.convState[i], 0, new Float32Array(CONV_C * 4));
+        this.device.queue.writeBuffer(this.recState[i], 0, new Float32Array(H * DK * DV));
+      }
+    }
+  }
+  #writeStep(t, pos0) {
+    this.device.queue.writeBuffer(this.bufStep, 0, new Uint32Array([t, pos0, 0, 0]));
+  }
+  // one full forward over T tokens already in bufHist at [pos0, pos0+T).
+  // debug: GPUBuffer to receive per-stage copies of the hidden state
+  // (slot 0 = embed, 1+i = after layer i, L+1 = final norm), T*C floats each.
+  #encodeForward(T, prefill, { debug = null, cascade = false } = {}) {
+    const { C, H, FF, NH, NKV, HD, DELTA_IN, CONV_C, ATTN_IN } = CFG;
+    const enc = this.device.createCommandEncoder();
+    let pass = enc.beginComputePass();
+    const disp = (pipe, bg, x, y = 1) => {
+      pass.setPipeline(pipe);
+      pass.setBindGroup(0, bg);
+      pass.dispatchWorkgroups(x, y);
+    };
+    const snap = (buf, slot) => {
+      if (!debug) return;
+      pass.end();
+      enc.copyBufferToBuffer(buf, 0, debug, slot * T * C * 4, T * C * 4);
+      pass = enc.beginComputePass();
+    };
+    const p = this.p;
+    // f16: gemv 8 rows/wg (sg=32), gemm 8x8 grid; q4: gemv 16 rows/wg
+    // (2 rows/subgroup), swiglu 8, tiled gemm 64x64 tiles
+    const rows = (n) => Math.ceil(n / (this.q4 ? 16 : 8));
+    const rowsSw = (n) => Math.ceil(n / 8);
+    const gmX = (n) => Math.ceil(n / (this.q4 ? 64 : 8));
+    const cols = Math.ceil(T / (this.q4 ? 64 : 8));
+    disp(p.embed, this.bgEmbed, T);
+    snap(this.bufX, 0);
+    for (let i = 0; i < CFG.L; i++) {
+      const lg = this.layers[i];
+      disp(p.norm, lg.bgNorm1, T);
+      if (lg.type === "delta") {
+        if (prefill) disp(p.gmDeltaIn, lg.bgInP, gmX(DELTA_IN), cols);
+        else disp(p.gvDeltaIn, lg.bgIn, rows(DELTA_IN));
+        disp(p.conv, lg.bgConv, Math.ceil(CONV_C / 256));
+        disp(p.gates, lg.bgGates, Math.ceil((T * H) / 64));
+        if (prefill) {
+          disp(p.dlocal, lg.bgLocal, Math.ceil(T / 64) * H);
+          disp(p.dscan, lg.bgScan, H);
+        } else {
+          disp(p.dstep, lg.bgStep, H);
+        }
+        disp(p.gnorm, lg.bgGnorm, T * H);
+        if (prefill) disp(p.gmOut, lg.bgOutP, gmX(C), cols);
+        else disp(p.gvOut, lg.bgOut, rows(C));
+      } else {
+        if (prefill) disp(p.gmAttnIn, lg.bgInP, gmX(ATTN_IN), cols);
+        else disp(p.gvAttnIn, lg.bgIn, rows(ATTN_IN));
+        disp(p.rope, lg.bgRope, T * (NH + NKV));
+        disp(p.attn, lg.bgAttn, T * NH);
+        disp(p.gateMul, lg.bgGateMul, Math.ceil((T * NH * HD) / 256));
+        if (prefill) disp(p.gmOut, lg.bgOutP, gmX(C), cols);
+        else disp(p.gvOut, lg.bgOut, rows(C));
+      }
+      disp(p.norm, lg.bgNorm2, T);
+      if (prefill) {
+        disp(p.gmFF, lg.bgGateP, gmX(FF), cols);
+        disp(p.gmFF, lg.bgUpP, gmX(FF), cols);
+        disp(p.swiglu, lg.bgSwP, Math.ceil((T * FF) / 256));
+        disp(p.gmDown, lg.bgDownP, gmX(C), cols);
+      } else {
+        disp(p.gvSwiglu, lg.bgSwiglu, rowsSw(FF));
+        disp(p.gvDown, lg.bgDown, rows(C));
+      }
+      snap(this.bufX, 1 + i);
+    }
+    disp(p.norm, this.bgFinalNorm, T);
+    snap(this.bufXn, CFG.L + 1);
+    disp(p.gvLogits, this.bgLogits(T - 1), rows(CFG.V));
+    if (cascade) {
+      disp(p.sliceTop, this.bgSliceTop, CFG.V / 64);
+      disp(p.candAdv, this.bgCandAdv, 1);
+    }
+    pass.end();
+    return enc;
+  }
+  pushTokens(ids) {
+    this.device.queue.writeBuffer(this.bufHist, this.pos * 4, new Uint32Array(ids));
+    this.hist.push(...ids);
+  }
+  // prefill; logits of the last token land in bufLogits
+  prefill(ids, opts = {}) {
+    let off = 0;
+    while (off < ids.length) {
+      const seg = ids.slice(off, off + this.maxT);
+      this.pushTokens(seg);
+      this.#writeStep(seg.length, this.pos);
+      this.device.queue.submit([this.#encodeForward(seg.length, true, opts).finish()]);
+      this.pos += seg.length;
+      off += seg.length;
+    }
+  }
+  // prefill + GPU-side argmax of the first generated token: after this,
+  // hist[P] holds the first token and step = {1, P} (P = prompt length).
+  startGenerate(ids) {
+    this.reset();
+    this.promptLen = ids.length;
+    let off = 0;
+    while (off < ids.length) {
+      const seg = ids.slice(off, off + this.maxT);
+      this.pushTokens(seg);
+      this.#writeStep(seg.length, this.pos);
+      const last = off + seg.length >= ids.length;
+      this.device.queue.submit([this.#encodeForward(seg.length, true, { cascade: last }).finish()]);
+      this.pos += seg.length;
+      off += seg.length;
+    }
+    this.gen = 1;
+  }
+  // n GPU-resident decode steps in ONE submit (no CPU roundtrips inside);
+  // returns all generated token ids so far (incl. the pending one).
+  async decodeChunk(n) {
+    const bufs = [];
+    for (let i = 0; i < n; i++) {
+      bufs.push(this.#encodeForward(1, false, { cascade: true }).finish());
+    }
+    this.device.queue.submit(bufs);
+    this.pos += n;
+    this.gen += n;
+    const total = this.promptLen + this.gen;
+    const rb = this.device.createBuffer({ size: total * 4, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
+    const enc = this.device.createCommandEncoder();
+    enc.copyBufferToBuffer(this.bufHist, 0, rb, 0, total * 4);
+    this.device.queue.submit([enc.finish()]);
+    await rb.mapAsync(GPUMapMode.READ);
+    const ids = Array.from(new Uint32Array(rb.getMappedRange().slice(0)));
+    rb.destroy();
+    this.hist = ids;
+    return ids.slice(this.promptLen);
+  }
+  decodeStep(tokenId, opts = {}) {
+    this.pushTokens([tokenId]);
+    this.#writeStep(1, this.pos);
+    this.device.queue.submit([this.#encodeForward(1, false, opts).finish()]);
+    this.pos += 1;
+  }
+}
+// GEMM prefill matvec grid note: gemm uses workgroup_size(8,8) over (N, T);
+// rows() above sizes the X dimension in units of 8 columns.

web/npy.js ADDED Viewed

	@@ -0,0 +1,16 @@

+// Minimal .npy loader (little-endian f32 only).
+export async function loadNpy(url) {
+  const buf = await (await fetch(url)).arrayBuffer();
+  const view = new DataView(buf);
+  if (view.getUint32(0, true) !== 0x4d554e93) throw new Error("not npy: " + url);
+  const major = view.getUint8(6);
+  const headerLen = major >= 2 ? view.getUint32(8, true) : view.getUint16(8, true);
+  const headerStart = major >= 2 ? 12 : 10;
+  const header = new TextDecoder().decode(new Uint8Array(buf, headerStart, headerLen));
+  const descr = header.match(/'descr':\s*'([^']+)'/)[1];
+  if (descr !== "<f4") throw new Error(`unsupported dtype ${descr}`);
+  if (/'fortran_order':\s*True/.test(header)) throw new Error(`fortran order not supported: ${url}`);
+  const shape = (header.match(/'shape':\s*\(([^)]*)\)/)[1].match(/\d+/g) ?? []).map(Number);
+  const data = new Float32Array(buf, headerStart + headerLen);
+  return { shape, data };
+}

web/tokenizer.js ADDED Viewed

	@@ -0,0 +1,122 @@

+// Byte-level BPE tokenizer for Qwen3.5 (HF tokenizer.json format).
+// Differences vs the lfm2 port: NFC normalizer, Qwen pretokenizer regex
+// (\p{M} marks glued to letters, single digits), 26 special tokens, and the
+// Qwen chat template (auto-inserts an empty think block, non-thinking mode).
+// GPT-2 byte <-> unicode table
+function byteMaps() {
+  const bs = [];
+  for (let i = 33; i <= 126; i++) bs.push(i);
+  for (let i = 161; i <= 172; i++) bs.push(i);
+  for (let i = 174; i <= 255; i++) bs.push(i);
+  const cs = bs.slice();
+  let n = 0;
+  for (let b = 0; b < 256; b++) {
+    if (!bs.includes(b)) {
+      bs.push(b);
+      cs.push(256 + n);
+      n++;
+    }
+  }
+  const b2c = new Array(256);
+  const c2b = new Map();
+  bs.forEach((b, i) => {
+    b2c[b] = String.fromCharCode(cs[i]);
+    c2b.set(cs[i], b);
+  });
+  return { b2c, c2b };
+}
+const { b2c, c2b } = byteMaps();
+// JS port of the Qwen pretokenizer regex; (?i:'s|'t|'re|'ve|'m|'ll|'d)
+// rewritten with explicit cases. Source pattern:
+// (?i:'s|'t|'re|'ve|'m|'ll|'d)|[^\r\n\p{L}\p{N}]?[\p{L}\p{M}]+|\p{N}| ?[^\s\p{L}\p{M}\p{N}]+[\r\n]*|\s*[\r\n]+|\s+(?!\S)|\s+
+const PRETOK =
+  /'(?:[sS]|[tT]|[rR][eE]|[vV][eE]|[mM]|[lL][lL]|[dD])|[^\r\n\p{L}\p{N}]?[\p{L}\p{M}]+|\p{N}| ?[^\s\p{L}\p{M}\p{N}]+[\r\n]*|\s*[\r\n]+|\s+(?!\S)|\s+/gu;
+export class Tokenizer {
+  static async load(url) {
+    const t = new Tokenizer();
+    const j = await (await fetch(url)).json();
+    t.vocab = new Map(Object.entries(j.model.vocab));
+    t.inv = new Map([...t.vocab].map(([s, i]) => [i, s]));
+    t.ranks = new Map(j.model.merges.map((m, i) => [Array.isArray(m) ? m.join(" ") : m, i]));
+    t.special = new Map();
+    for (const a of j.added_tokens) {
+      t.special.set(a.content, a.id);
+      t.inv.set(a.id, a.content);
+    }
+    t.specialRe = new RegExp(
+      "(" + [...t.special.keys()].map((s) => s.replace(/[|<>]/g, "\\$&")).join("|") + ")"
+    );
+    t.eos = t.special.get("<|im_end|>");
+    t.eot = t.special.get("<|endoftext|>");
+    return t;
+  }
+  #bpe(word) {
+    let parts = [...word];
+    while (parts.length > 1) {
+      let best = null, bestRank = Infinity;
+      for (let i = 0; i < parts.length - 1; i++) {
+        const r = this.ranks.get(parts[i] + " " + parts[i + 1]);
+        if (r !== undefined && r < bestRank) {
+          bestRank = r;
+          best = i;
+        }
+      }
+      if (best === null) break;
+      parts = [...parts.slice(0, best), parts[best] + parts[best + 1], ...parts.slice(best + 2)];
+    }
+    return parts;
+  }
+  #encodeText(text) {
+    const ids = [];
+    const enc = new TextEncoder();
+    for (const m of text.matchAll(PRETOK)) {
+      const chars = Array.from(enc.encode(m[0]), (b) => b2c[b]).join("");
+      for (const tok of this.#bpe(chars)) {
+        const id = this.vocab.get(tok);
+        if (id === undefined) throw new Error(`token not in vocab: ${JSON.stringify(tok)}`);
+        ids.push(id);
+      }
+    }
+    return ids;
+  }
+  encode(text) {
+    const ids = [];
+    for (const piece of text.normalize("NFC").split(this.specialRe)) {
+      if (!piece) continue;
+      if (this.special.has(piece)) ids.push(this.special.get(piece));
+      else ids.push(...this.#encodeText(piece));
+    }
+    return ids;
+  }
+  decode(ids, { skipSpecial = true } = {}) {
+    const bytes = [];
+    for (const id of ids) {
+      const tok = this.inv.get(id);
+      if (tok === undefined) continue;
+      if (this.special.has(tok)) {
+        if (!skipSpecial) bytes.push(...new TextEncoder().encode(tok));
+        continue;
+      }
+      for (const ch of tok) bytes.push(c2b.get(ch.charCodeAt(0)));
+    }
+    return new TextDecoder().decode(new Uint8Array(bytes));
+  }
+  // chat_template.jinja equivalent for text-only chats (non-thinking mode:
+  // the template pre-fills an empty think block in the generation prompt)
+  applyChatTemplate(messages, { addGenerationPrompt = true } = {}) {
+    let s = "";
+    for (const m of messages) {
+      s += `<|im_start|>${m.role}\n${m.content}<|im_end|>\n`;
+    }
+    if (addGenerationPrompt) s += "<|im_start|>assistant\n<think>\n\n</think>\n\n";
+    return this.encode(s);
+  }
+}