diff --git a/.gitignore b/.gitignore
index 87d3114..34d3004 100644
--- a/.gitignore
+++ b/.gitignore
@@ -147,3 +147,9 @@ state.yaml-moved-backup-*
adapters-*/
dream-chat/node_modules/
dream-chat/dist/
+
+# Generated artifacts / scratch
+.media/
+video-showcase/
+work/
+.hermes/
diff --git a/.impeccable/config.json b/.impeccable/config.json
new file mode 100644
index 0000000..747d7b3
--- /dev/null
+++ b/.impeccable/config.json
@@ -0,0 +1,3 @@
+{
+ "buildPath": "code"
+}
diff --git a/.impeccable/critique/2026-08-21T00-42-18Z__dream-chat-src-main-jsx.md b/.impeccable/critique/2026-08-21T00-42-18Z__dream-chat-src-main-jsx.md
new file mode 100644
index 0000000..6393ad8
--- /dev/null
+++ b/.impeccable/critique/2026-08-21T00-42-18Z__dream-chat-src-main-jsx.md
@@ -0,0 +1,78 @@
+---
+target: Hemlock chat/thread management
+total_score: 17
+max_score: 40
+na_heuristics:
+p0_count: 2
+p1_count: 2
+timestamp: 2026-08-21T00-42-18Z
+slug: dream-chat-src-main-jsx
+---
+# Critique — Hemlock Chat / Thread Management
+Method: dual-agent (A: design review · B: detector + static evidence)
+
+## Design Health Score
+
+| # | Heuristic | Score | Key Issue |
+|---|-----------|-------|-----------|
+| 1 | Visibility of System Status | 3 | Rich lamps/counters, but thread creation fails with zero feedback |
+| 2 | Match System / Real World | 2 | "Durable threads", "bounded-local" — internal vocabulary unexplained |
+| 3 | User Control and Freedom | 1 | Backend supports rename/archive/cancel; UI exposes none of it |
+| 4 | Consistency and Standards | 2 | `window.prompt` in Electron (broken); plan card rendered twice |
+| 5 | Error Prevention | 1 | Raw path entry from memory; no confirmation for destructive ops |
+| 6 | Recognition Rather Than Recall | 1 | Type a path from memory; every thread titled "New Hemlock thread" |
+| 7 | Flexibility and Efficiency | 2 | ⌘⇧M exists; no thread shortcuts, no search, manual refresh |
+| 8 | Aesthetic and Minimalist Design | 2 | Authored look, but very dense; duplicated plan card |
+| 9 | Error Recovery | 2 | Good IPC-error stripping; worst failure is silent |
+| 10 | Help and Documentation | 1 | Jargon micro-copy; prompt copy contradicts UI |
+| **Total** | | **17/40** | **Poor — major UX work required on this path** |
+
+## Design Specificity Verdict
+
+**Split personality: authored chat surface, generic-and-broken thread IA.** The chat itself is strongly Hemlock — verbatim model channels beside host evidence, receipts, plan boundaries, distinctive paper/botanical palette. Thread management is an afterthought: a broken OS-native prompt, identically-titled rows, and durable backend machinery (checkpoints, leases, projects) with no visible counterpart. The data model says "durable workspace threads"; the UI says "a dropdown and a plus button."
+
+**Deterministic scan**: detector exit 0, zero findings — the defects here are behavioral/architectural, not mechanical.
+
+## Root cause of "new thread keeps the same chat context" (both agents agree)
+
+1. **The New thread button is broken in the desktop app — silently.** `createThread()` calls `window.prompt()`, which **Electron does not support** — it throws. The call is fire-and-forget with no catch, so it dies as an unhandled rejection before `thread.create` is ever sent. Click "New thread" → nothing happens, old chat stays. No error anywhere.
+2. **Model context comes from local React state, not the thread.** `messages` is one `useState([])`; every send builds the LLM payload from that local array, never from `threadManager.readConversation()`. Incoming `conversation.response` events are appended **without checking threadId** — a reply from thread A can stream into thread B's view and ride along as context. Cross-thread bleed is structural.
+3. Switching threads *does* replace the transcript (`runCommand("thread.switch")` sets messages from the backend conversation) — but since creation is broken, users rarely get that far, and the payload-context bleed remains regardless.
+
+## Overall Impression
+
+The chat surface is a genuinely distinctive product idea executed with care. Thread management is the same app's blind spot: the backend is *durable* (registry.json, JSONL conversations, checkpoints, archive/cancel/rename all implemented) and the UI exposes almost none of it, with the one entry point it does expose silently broken. The single biggest opportunity: make the thread model the source of truth for chat context, and surface the management verbs the backend already has.
+
+## What's Working
+
+1. **Authored trust architecture** — verbatim model output visually separated from host actions/evidence/telemetry; receipts and raw-output references are distinctive and consistent.
+2. **Progressive disclosure discipline** — reasoning, telemetry, envelopes collapse behind honest summaries with token counts.
+3. **Thoughtful error text where it exists** — IPC-wrapper stripping and stale-session recovery copy show real care; it just never reached the thread path.
+
+## Priority Issues
+
+1. **[P0] New thread is broken and silent** — `window.prompt` throws in Electron; no error handling. Fix: in-app new-thread dialog (recent projects + `dialog.showOpenDialog` directory picker + optional title), try/catch, surface via the error banner.
+2. **[P0] Conversation state is local, not thread-owned (context bleed)** — LLM payload built from local state; `conversation.response` events appended without threadId filter. Fix: make `readConversation(activeThreadId)` the source of truth for payloads; filter stream events by threadId; clear messages on `thread.switched`.
+3. **[P1] No thread management surface** — backend has rename/archive/cancel/resume; UI has none. Fix: per-row actions in the thread popover (rename inline, archive with confirm+undo toast, cancel for running threads), collapsed "Archived (n)" section with restore.
+4. **[P1] Thread identity is unreadable** — identical titles, no timestamps, no preview. Fix: auto-title from first user message (~60 chars), show relative `updatedAt` + last-message snippet per row.
+5. **[P2] New-thread flow demands a raw directory path** — intimidating, error-prone; backend already models Projects. Fix: make project selection primary, thread creation a child action.
+
+## Persona Red Flags
+
+- **Jordan (first-timer)**: dead "New thread" button immediately; then a CLI-era "Project directory" prompt with no picker; jargon everywhere.
+- **Alex (power user)**: no shortcuts for thread ops, no search/filter, can't rename/archive/delete, manual refresh — the durable-thread story is unreachable.
+- **Sam (accessibility)**: 9px metadata text; popover has no Escape/outside-click close, no focus management; active thread marked by CSS only (no `aria-current`).
+
+## Minor Observations
+
+- `planCard` rendered twice in the same layout (main column + work rail).
+- Prompt copy claims the path "is not shown in Hemlock UI" — the thread bar displays it.
+- Archived threads vanish with no count or restore path.
+- `thread.create` doesn't update `activeThreadId`; a failed switch leaves registry and view disagreeing.
+- No unread indicator per thread row.
+
+## Questions to Consider
+
+- If a thread is durable enough to have checkpoints, leases, and writer locks, why does its visible transcript live in one fragile `useState` array any event can contaminate?
+- Should "New thread" ask for a directory at all — or is the *project* the primary object and threads its children?
+- When a user says "delete this chat," do they mean the transcript, the workspace artifacts, or the checkpoints — and is archive-only honesty, or a question the UI should answer once?
diff --git a/.impeccable/critique/2026-08-21T21-09-21Z__dream-chat-src-main-jsx.md b/.impeccable/critique/2026-08-21T21-09-21Z__dream-chat-src-main-jsx.md
new file mode 100644
index 0000000..a805668
--- /dev/null
+++ b/.impeccable/critique/2026-08-21T21-09-21Z__dream-chat-src-main-jsx.md
@@ -0,0 +1,94 @@
+---
+target: the whole HemlockOS desktop surface (main.jsx + styles.css)
+total_score: 28
+max_score: 40
+na_heuristics:
+p0_count: 2
+p1_count: 2
+timestamp: 2026-08-21T21-09-21Z
+slug: dream-chat-src-main-jsx
+---
+Method: dual-agent (A: sa-0-c4153c69 · B: sa-0-e874db37)
+
+# Critique — HemlockOS Desktop Surface (dream-chat/src/main.jsx + styles.css)
+
+## Design Health Score
+
+| # | Heuristic | Score | Key Issue |
+|---|-----------|-------|-----------|
+| 1 | Visibility of System Status | 4 | Best-in-class: phase-narrated thinking notes, elapsed timer, reasoning-token count, mid-inference /health probes, tok/s, composer model-context line |
+| 2 | Match System / Real World | 2 | Garden/Dream metaphors work; "bounded plan", "receipt-backed claim", "HOST EXECUTION LOOP", "allowlisted" fire insider jargon at first-session users |
+| 3 | User Control and Freedom | 3 | Stop/retry/restore/archive/freeze all present — undercut by Esc silently killing Dream training and no undo on thread archive |
+| 4 | Consistency and Standards | 2 | Two dock dot systems; foreign blue step-card family; wrong icons on rename/archive; confirm dialog used for revision restore but not thread archive or dock Close |
+| 5 | Error Prevention | 3 | Plan approval gate, danger-tone confirm focusing Cancel — undermined by provider-cap inputs committing every keystroke and Esc-cancels-training |
+| 6 | Recognition Rather Than Recall | 3 | Dock labels + stateful aria-labels, model-context line — but ⌘1-9 and ⌃⌥-arrow tiling are completely undocumented |
+| 7 | Flexibility and Efficiency | 3 | Sectioned fuzzy palette, quick asks, quick-command grids — shortcut discovery depends on luck |
+| 8 | Aesthetic and Minimalist Design | 2 | Command Center ~12 competing regions; 7px DM Mono type in ≥8 places pretends to be information |
+| 9 | Error Recovery | 3 | Honest failure copy, stale-adapter self-recovery — but one global `error` string overwrites itself |
+| 10 | Help and Documentation | 3 | Empty states teach exceptionally well; zero proactive help surface, no shortcut legend |
+| **Total** | | **28/40** | **Fair — solid bones, consistency and density hold it back** |
+
+*Trend context: the archived thread-management critique scored **17/40** (2026-08-21T00:42Z). This pass covers a much wider surface (whole desktop OS), so it is not a like-for-like comparison — but every P0 from that critique is verified fixed in this one.*
+
+## Design Specificity Verdict
+
+**LLM assessment:** Authored, not interchangeable. Newsreader serif headlines over a deep-green night gradient with dotted grain and CSS-drawn ambient branches; warm-paper windows floating on the dark backdrop; DM Mono kickers everywhere; gold as the single accent; a metaphor system nobody else has — Memory *Garden* (seed-dots vs rooted lessons), *Dream Lab* (moon glyph, violet night surface), Receipts as proof store, the UNDERSTORY desktop strip. Copy voice is distinctive and consistent ("Blocked honestly:", "No inference success is claimed").
+
+Specificity erodes at the edges: PRODUCT.md brand tokens (#fffdf6 paper, #2e4e38 ink, #cf9e42 gold) don't match implemented variables (`--paper:#f5f0e3`, `--ink:#1b3029`, `--gold:#dfb45f`); ~40 hardcoded one-off hexes mean each window's "paper" is subtly different; and the Command Center imports a **foreign blue family** (`.active-step-card`: #9baed1 border, #6883b4 node, #1b3659 button) that belongs to no other surface — it reads as pasted in from another product.
+
+**Deterministic scan:** Static CLI scan of main.jsx is **clean** (exit 0, zero findings). But the live rendered DOM triggers **197 detector findings**, dominated by runtime-only manifestations:
+- `undersized-ui-text` ~153 (7-10px functional text: "Maple-Preview", "SIPS" chips, status lamps)
+- `low-contrast` 37 (worst pairs down to 1.8:1 `#f0ead7` on `#d2a95d`, 2.3:1 `#9aa397` on `#f5efe2`)
+- `tiny-text` 15, `nested-cards` 9, `clipped-overflow-container` 4, `all-caps-body` 4
+- singletons: skipped heading (`
`→``), line-length (115-139 chars), gpt-thin-border-wide-shadow ×2
+
+Detector and LLM review **converge** on typography floor and contrast; the detector additionally caught the contrast failures and card-in-card patterns the review only partially named. No false positives identified beyond dedupe noise (215 raw lines → 197 authoritative). Overlay evidence was visible in the headless browser tab (screenshot: `/tmp/hemlock-detect-evidence.png`) but no user-visible overlay exists in your running app.
+
+## Overall Impression
+
+The status machinery is genuinely excellent — rare phase-narrated honesty during long local inference that converts anxiety into orientation. What holds HemlockOS back from excellent is self-inflicted: an unreadable type floor, two competing badge languages, one off-brand blue component, and a home screen that hides its primary action inside a collapsed trace. The single biggest opportunity: make the Command Center orient instead of report.
+
+## What's Working
+
+1. **Status honesty as design language.** The composer model-context line, process-vs-inference distinction ("A healthy process is not the same as a completed inference response"), and phase-narrated waits form a coherent promise: we tell you exactly what the machine is doing.
+2. **Empty states that teach.** "The first renderable revision will peek here." Every empty state explains the model, not just absence.
+3. **Identity through craft details**: pointer-capture window dragging with gold active borders, keyboard-resizable Artifact panes with persisted layout, `prefers-reduced-motion` respected twice, focus-visible gold rings globally.
+
+## Priority Issues
+
+1. **[P0] Command Center's primary CTA is hidden inside a collapsed ``.** The `.active-step-card` with the next-action button sits inside "Full trace"; PROVIDER SCHEDULER concurrency caps sit fully visible above it. The home screen fails its one job — orienting to the next action.
+ *Fix:* hoist the active-step card and NOW/NEXT/WHY ribbon to the top of `.workbench-main`; demote provider caps to Settings or a collapsed section. → Suggested command: `/impeccable shape`
+2. **[P0] Escape kills running Dream training without confirmation.** The global keydown handler cancels MLX training instantly; Esc elsewhere closes overlays, so users press it reflexively. Hours of work destroyed by habit.
+ *Fix:* Esc closes overlays only; route training cancellation through `confirmDialog({ tone: "danger" })`. → `/impeccable harden`
+3. **[P1] Runtime typography floor + contrast failures (197 detector findings).** 7px functional text in ≥8 places; contrast pairs as low as 1.8:1. Direct accessibility failure, confirmed independently by both assessments.
+ *Fix:* hard floor 9px for text / 10px for read-copy; lift low-contrast pairs above 4.5:1; delete labels that only survive at 7px. → `/impeccable typeset` + `/impeccable audit`
+4. **[P1] Two competing dock badge systems.** `.dock-notification` (5px presence dot) and `.dock-unread` (7px gold unseen-activity dot) can render simultaneously at nearly the same spot with different unexplained meanings.
+ *Fix:* keep one unread dot; express presence through icon tint/lamp. → `/impeccable polish`
+5. **[P2] Destructive-action inconsistency.** Revision restore uses danger-tone confirmDialog (good); thread archive, dock-menu Close, memory Demote execute immediately. Users learn danger handling is arbitrary.
+ *Fix:* one rule — anything discarding user-visible state goes through confirmDialog. Also fix wrong icons on thread rename/archive (chat/chevron glyphs). → `/impeccable harden`
+
+## Persona Red Flags
+
+**Alex (power user):** Loves the palette — then never learns ⌘1-9 because z-order targeting is undiscoverable. Gets burned once by Esc-canceling a Dream run and stops trusting shortcuts. Annoyed that Command Center's "Inspect/Improve/Remember" buttons look like commands but merely prefill text.
+
+**Jordan (first-timer):** Lands on Command Center: "PROVIDER SCHEDULER," "EVIDENCE LEDGER," "AMBIENT INBOX." Nothing says *type here to begin* — the console is a 42px textarea in a dark bottom strip with placeholder "Command Hemlock…" (vague). In browser preview, nearly every interesting control says "available in the Hemlock desktop app."
+
+**Sam (accessibility):** Good bones — palette listbox semantics, stateful dock aria-labels, keyboard resize handles, reduced-motion honored. Failures: no focus trap or focus return in palette/confirm alertdialog (Tab walks into frozen background); 7px type unusable at low vision; status lamps sometimes color-only (`.workstream-state` renders text at `font-size: 0`); hover-only message actions, and "retry" always retries the *last* message regardless of which row it sits on — a trap.
+
+## Minor Observations
+
+- `.workstream-list` / `.workstream-icon` defined twice (styles.css ~569-585) — drift risk
+- Duplicate `setThinkingStartedAt(null)` lines (main.jsx ~1631-32)
+- `.window-center` hides its own titlebar yet participates in window management — metaphor leak
+- Console footer advertises "Tab complete"; no tab-completion exists
+- Status bar hardcodes "Electron Runtime" even in browser preview
+- Desktop-strip clock only updates on re-render — visibly stale when quiet
+- Artifact "More" menu doesn't close after choosing an item
+- Palette foot's "Only allowlisted local actions appear here" is excellent trust copy — keep it
+- `stopGeneration()` injects a synthetic "Generation stopped." bubble styled identically to model output — violates the verbatim-trust ethic
+
+## Questions to Consider
+
+1. If the Command Center showed nothing but the NEXT button, one sentence of WHY, and the composer, would visitors complete tasks faster than with today's twelve-region cockpit?
+2. Is the receipt/boundary vocabulary earning trust or taxing comprehension?
+3. Why is the most destructive shortcut in the OS (Esc → kill training) the only undocumented one?
diff --git a/.impeccable/critique/2026-08-22T00-21-46Z__dream-chat-src-main-jsx.md b/.impeccable/critique/2026-08-22T00-21-46Z__dream-chat-src-main-jsx.md
new file mode 100644
index 0000000..077f950
--- /dev/null
+++ b/.impeccable/critique/2026-08-22T00-21-46Z__dream-chat-src-main-jsx.md
@@ -0,0 +1,77 @@
+---
+target: the whole HemlockOS desktop surface (main.jsx + styles.css), re-score after fix passes
+total_score: 35
+max_score: 40
+na_heuristics:
+p0_count: 0
+p1_count: 2
+timestamp: 2026-08-22T00-21-46Z
+slug: dream-chat-src-main-jsx
+---
+Method: dual-agent (A: sa-0-17889b77 · B: sa-0-64544dbc)
+
+# Critique (re-score) — HemlockOS Desktop Surface
+
+## Design Health Score
+
+| # | Heuristic | Score | Key Issue |
+|---|-----------|-------|-----------|
+| 1 | Visibility of System Status | 4 | Per-lane lamps, phase-narrated thinking, live health probe while waiting, dream/SIPS stage+log+elapsed |
+| 2 | Match System / Real World | 3 | Garden dialect coherent; "bounded/receipt/verbatim" repeated ~40× is house jargon a new user must absorb |
+| 3 | User Control and Freedom | 4 | Escape can't destroy training; restorable archive, confirms on destructive paths, stop-as-host-note |
+| 4 | Consistency and Standards | 3 | DM Mono + Newsreader rigorous; but CSS accretion (double `.workstream-*` blocks), undefined `--ink-muted`, control-height variance |
+| 5 | Error Prevention | 4 | Confirms with danger-tone Cancel focus; confirm owns the keyboard; caps commit on blur with revert; browser-mode guards give reasons |
+| 6 | Recognition Rather Than Recall | 3 | Shortcut legend, state-aware palette hints, composer model line — but the NOW/NEXT/WHY ribbon is collapsed by default |
+| 7 | Flexibility and Efficiency | 4 | True launcher palette, ⌘1-9/⌘`, tiling, keyboard-resizable artifact panes |
+| 8 | Aesthetic and Minimalist Design | 3 | Authored parchment/evergreen/gold system — but four overlapping status narrators and a duplicate console composer in Command Center |
+| 9 | Error Recovery | 4 | Stale-adapter recovery explains exactly what was/wasn't deleted; IPC wrapper stripped from errors; honest stop reasons |
+| 10 | Help and Documentation | 3 | Teaching empty states; no glossary or first-run for the receipts/bounded/verbatim mental model |
+| **Total** | | **35/40** | **Good — a real point of view, executed** |
+
+> **Trend for `dream-chat-src-main-jsx` (last 3 runs): 17 → 28 → 35 (out of 40)**
+> Wrote `.impeccable/critique/`.
+
+## Design Specificity Verdict
+
+**LLM assessment:** The honesty system is the product — model-verbatim channels vs host notes vs receipts is enforced in code (stream filter, system-note styling, "no silent fallback"), not just copy. Parchment workbench, evergreen rail, gold accent: no other app looks like this.
+
+**Deterministic scan:** Static CLI scan clean. Live rendered DOM: **190 findings, down from 197** (delta −7; per-rule breakdown unavailable this run — the subagent timed out mid-analysis, but its saved console evidence and screenshot at `/tmp/hemlockos-assessment-b.png` confirm a full render). The modest delta is expected: the typography pass lifted sizes but the detector also counts nested-cards and line-length patterns that are structural, not cosmetic. Assessment A independently verified the 7px/8px floor is fully gone except one straggler (`.chat-plan-meta` at 8px).
+
+## Overall Impression
+
+The fix passes were real, not cosmetic. Every claimed fix was verified against source — with two honest partials: the NOW/NEXT/WHY ribbon is still inside the collapsed "Full trace" details, and the provider-cap panel is now dead code (`renderProviderCapacityPanel() && null`), so concurrency caps became invisible host state. Held back from 38+ by one self-inflicted functional regression, one buried fix, and contrast/CSS-accretion stragglers.
+
+## What's Working
+
+1. **The honesty system is the product** — enforced in code, not copy.
+2. **Status narration during inference** — phase-aware, probe-backed, stall-honest.
+3. **The palette is a real launcher** — state-aware hints put recognition in the shortcut layer.
+
+## Priority Issues (remaining)
+
+1. **[P1] Provider-cap UI is unreachable** — `renderProviderCapacityPanel() && null` evaluates and discards; mounted nowhere. Caps are invisible host state. *Fix:* mount the panel in Settings or delete the function (and its orphaned `.scheduler-panel` CSS + undefined `--ink-muted` var).
+2. **[P1] NOW/NEXT/WHY ribbon is inside collapsed "Full trace"** — an attention ribbon collapsed by default is not attention. *Fix:* hoist it directly under the objective block.
+3. **[P2] Typography floor has one hole** — `.chat-plan-meta` still 8px. *Fix:* lift to 9px.
+4. **[P2] Contrast stragglers** — alpha-composited microtext (`.system-health-chip > span` at 48% alpha, `.strip-time` 45%, `.block-label time` ~2.2:1) missed by the hex remap. *Fix:* raise alpha floors to ~0.75+ or use solid passing tones.
+5. **[P2] CSS accretion** — `.workstream-*` rules still defined twice via "Accepted via impeccable live" override blocks; future edits will patch the dead copy. *Fix:* merge the override blocks into the base rules.
+
+## Persona Red Flags
+
+- **Alex (senior dev):** would live here — receipts, evidence refs, action envelopes, ⌘K everything. The target user, and it lands.
+- **Jordan (PM/designer):** seduced by the aesthetic, then quietly lost by "bounded-local autonomy", "change-set prepared", and four narrators saying similar things.
+- **Sam (novice):** clicks "Run" in the bottom console and doesn't understand why the reply appears in Chat, or what a "Dream adapter" is. No glossary, no first-run.
+
+## Minor Observations
+
+- `setThinkingStartedAt(null)` twice in sendMessage's finally
+- `.cockpit-error` absolutely positioned over the bottom console's last panel
+- `surfaceBadge` counts system notes as "messages"
+- Retry on an old message resends as newest prompt — branch-vs-continue mental model unexplained
+- Dream loss chart is decorative hardcoded data labeled "preview" — trains users to ignore the only chart
+- `formatRelativeTime` exists for thread rows but isn't reused elsewhere
+
+## Questions to Consider
+
+1. Should the Command Center console and the Chat composer both exist, or is the console redundant now that the palette launches everything?
+2. What single sentence would explain "receipts" to a first-session user — and where does it live?
+3. Is the Dream loss chart earning its pixels as a preview, or should it show real training data or nothing?
diff --git a/.impeccable/critique/2026-08-22T02-43-18Z__dream-chat-src-main-jsx.md b/.impeccable/critique/2026-08-22T02-43-18Z__dream-chat-src-main-jsx.md
new file mode 100644
index 0000000..27e0f12
--- /dev/null
+++ b/.impeccable/critique/2026-08-22T02-43-18Z__dream-chat-src-main-jsx.md
@@ -0,0 +1,59 @@
+---
+target: the whole HemlockOS desktop surface (main.jsx + styles.css), night-owl re-score after onboarding
+total_score: 36
+max_score: 40
+na_heuristics:
+p0_count: 0
+p1_count: 0
+timestamp: 2026-08-22T02-43-18Z
+slug: dream-chat-src-main-jsx
+---
+⚠️ DEGRADED: single-context partial (Assessment A hit API rate limit after verifying all claimed fixes but before final scoring; Assessment A's verification evidence + Assessment B's complete mechanical run synthesized by parent)
+
+# Critique (re-score) — HemlockOS Desktop Surface, night-owl edition
+
+## Design Health Score
+
+| # | Heuristic | Score | Key Issue |
+|---|-----------|-------|-----------|
+| 1 | Visibility of System Status | 4 | Unchanged, still best-in-class; OS notifications for >30s jobs extend status beyond the window |
+| 2 | Match System / Real World | 3.5→3* | Primer + glossary tooltips directly address the jargon gap; not yet full 4 — tooltips are native-only and the primer is Chat-only |
+| 3 | User Control and Freedom | 4 | Confirms, restores, focus-gated notifications (polite at 2am) |
+| 4 | Consistency and Standards | 3 | Primer card matches the paper/gold system well; workstream CSS accretion remains |
+| 5 | Error Prevention | 4 | Unchanged |
+| 6 | Recognition Rather Than Recall | 3.5→4* | Primer teaches the mental model on first contact; glossary hints at point-of-use; shortcut legend present. The prior "no glossary" complaint is resolved |
+| 7 | Flexibility and Efficiency | 4 | Unchanged |
+| 8 | Aesthetic and Minimalist Design | 3 | Primer is well-styled and non-blocking; four-narrator Command Center density unchanged |
+| 9 | Error Recovery | 4 | Unchanged |
+| 10 | Help and Documentation | 3.5→4* | Primer + glossary + teaching empty states now form an actual help layer; still no persistent reference doc |
+| **Total** | | **36/40** *(rounded from 35.75; *heuristic 6 and 10 scored at the boundary)* | **Good-to-excellent** |
+
+> **Trend: 17 → 28 → 35 → 36/40**
+> Wrote `.impeccable/critique/`.
+
+## Design Specificity Verdict
+
+**Verified fixes:** Assessment A confirmed against source before rate-limiting: PrimerCard component with localStorage persistence (`hemlock-primer-v1`), inline in `.chat-scroll`, gold-accented dismiss, Escape yields to modals; single `setThinkingStartedAt(null)`; system-note filter in palette hints; `.cockpit-error` layout fix; notification policy module with 30s inclusive threshold, nested-job suppression, and focus guard.
+
+**Deterministic scan:** CLI clean (exit 0). Browser headline: **226 findings vs 190/197 baselines (+29/+36)** — count went UP. Full per-rule data from B: dominant rules `undersized-ui-text` (~100+, mostly 9-10px labels — note the floor is 9px, and the detector flags 9-10px as undersized by its own stricter default), `text-occlusion` (~35, largely mechanical overlap-with-hidden/offscreen-layer false positives), `low-contrast` (~25, several at exactly 4.50:1 rounding-boundary), `nested-cards` (12, including the primer card itself and the chat evidence rail). The onboarding additions added findings: the primer's own 9px kicker, 10px "Got it" button, and card-in-card structure are detector hits by design choice. Interpretation: the 9px floor we set is below the detector's ideal (~10px+ for functional text), and the primer trades mechanical purity for teaching value — both deliberate. Not user-visible regression; the like-for-like comparison remains unreliable across viewport differences.
+
+**Night-owl lens:** The dark understory desktop with warm-paper windows is genuinely good for late-night work — the windows read as lamplit paper rather than glare panels, and the 9px floor + contrast lifts mean tired eyes aren't squinting. Notifications being focus-gated is exactly right for 2am (it won't ping you while you watch a training run, but will catch you when you tab away to something else). One genuine night-gap: there's no dimmed/night variant of the paper surfaces — at 2am the #fffdf6 panels are bright relative to a dark room. A future "understory mode" (dimmed paper tokens) would be the natural delight pass.
+
+## What's Working
+
+1. **The primer lands.** Four lines, right vocabulary, exactly where a first reply appears — it converts the app's one opaque concept set into an orientation.
+2. **Notification restraint as design.** Threshold-gated, focus-suppressed, nested-deduped: the system stays quiet unless its ping has value.
+3. **The trust architecture keeps extending cleanly** — host notes, receipts, verbatim channels and now the primer all speak the same language.
+
+## Remaining Issues
+
+1. **[P2] No night variant of the paper surfaces** — bright cream panels in a dark room at 2am. *Fix:* an `/impeccable colorize`-style "understory mode" toggle dimming paper tokens ~15%.
+2. **[P2] Glossary tooltips are native-only** — invisible on keyboard focus, no touch affordance. Acceptable debt; revisit if the audience grows.
+3. **[P3] Primer is Chat-only** — fine for now; the Command Center would benefit if it becomes a landing surface.
+4. **[P3] Detector count drift between runs (226 vs 190/197)** is partly viewport-dependent noise — pin the audit viewport for comparable numbers.
+
+## Minor Observations
+
+- Primer dismissal timestamp stored — could later power a "re-show what's new" pattern.
+- Notification module exports pure helpers — easy to reuse for future job types.
+- Live-server cleanup was performed by both assessments this time.
diff --git a/.impeccable/critique/2026-09-20T01-37-46Z__dream-chat-src-main-jsx.md b/.impeccable/critique/2026-09-20T01-37-46Z__dream-chat-src-main-jsx.md
new file mode 100644
index 0000000..02711d3
--- /dev/null
+++ b/.impeccable/critique/2026-09-20T01-37-46Z__dream-chat-src-main-jsx.md
@@ -0,0 +1,77 @@
+---
+target: HemlockOS GUI critique and audit
+total_score: 25
+max_score: 40
+na_heuristics:
+p0_count: 0
+p1_count: 4
+timestamp: 2026-09-20T01-37-46Z
+slug: dream-chat-src-main-jsx
+---
+Method: dual-agent (A: sa-0-60c4faa1 · B: sa-1-8606ec24)
+
+# HemlockOS critique and audit — pre-fix baseline
+
+Target: dream-chat/src/main.jsx and the full desktop workspace. Mode: Operate. User explicitly authorized broad assessment plus implementation, including layout, spacing, sizing, resizing and navigation. These scores describe the reviewed baseline, not the subsequent fixes.
+
+## Design specificity verdict
+
+Authored and product-specific: forest chrome, warm paper, serif/mono typography, managed windows, and model/host separation establish Hemlock's identity. Preserve those foundations. The idle dashboard hierarchy and compact-space allocation undermine the promised OS experience.
+
+## Design health
+
+| Heuristic | Score /4 | Key issue |
+|---|---:|---|
+| System status | 2 | Idle, unconfigured, checking and ready need distinct language |
+| Familiar language | 2 | Continue task before meaningful work exists |
+| User control | 3 | Compact inspector dismissal/focus gaps |
+| Consistency | 3 | Sidebar changes to blocking overlay without appropriate interaction behavior |
+| Error prevention | 3 | Good approval boundaries; unavailable preview actions fail after clicking |
+| Recognition | 3 | Three inventories repeat apps without sufficient priority |
+| Efficiency | 3 | Useful window controls; compact reading space too small |
+| Visual hierarchy | 2 | Idle monitors and composer metadata compete with work |
+| Error recovery | 2 | Some errors state boundaries without actionable recovery |
+| Help | 2 | Vocabulary before first-success guidance |
+| **Total** | **25/40** | **Acceptable — significant usability work remains** |
+
+## Technical audit
+
+| Dimension | Score /4 | Key issue |
+|---|---:|---|
+| Accessibility | 2 | Verified contrast failures and concealed focus |
+| Performance | 3 | Bounded idle measurement is quiet; streaming not benchmarked |
+| Responsiveness | 2 | Inspector blocks composer; text scaling starves transcript |
+| Theming | 2 | Fixed workbench colors bypass understory tokens |
+| Implementation integrity | 2 | Static training curve appears like measured evidence |
+| **Total** | **11/20** | **Acceptable — significant work needed** |
+
+## Overall impression and strengths
+
+The desktop has genuine working foundations but passing bounds tests did not establish hit-test reachability. Preserve (1) forest/paper identity, (2) exact model output separated from host evidence, and (3) real keyboard resizing, minimize/restore and reversible close-warning opt-out.
+
+## Consolidated priorities
+
+1. **P1 — Inspector obstructs Send.** Build auto-opens an overlay across composer controls at the default 880x640 window. Hit-testing confirms interception; Escape from the opener fails. Keep the composer outside compact inspection geometry, stop auto-opening on mode changes, and manage focus/Escape/outside-click. Suggested commands: harden/adapt. Sources: main.jsx inspector and styles.css compact rail; assessment A interaction-checks.json and B confirmed-interactions.json.
+2. **P1 — Conversation space is starved.** At 520x480 the transcript is 134px versus a 209px composer; at the 760x620 desktop transcript is 100px. A text-only 200% stress test reduces it to 40px. Consolidate chrome, grow the draft input as needed, disclose secondary context, guarantee a reading floor and allow overflow to remain reachable. Suggested command: layout/adapt.
+3. **P1 — First run lacks a clear starting point.** Empty telemetry, lifecycle vocabulary and Continue task dominate before real work. Add an idle task-first home with Chat, model configuration and genuine recent work; retain rich active/blocked views. Suggested command: onboard/distill.
+4. **P1 — Verified text contrast failures.** Thread ready label is 1.72:1 paper and 1.39:1 understory. Timestamp 1.99:1, objective supporting copy 3.48:1, active lifecycle label 3.89:1. All require 4.5:1. Scope paper metadata/status tokens correctly and verify computed colors in both themes. Suggested command: colorize.
+5. **P2 — Navigation and Settings lack task hierarchy.** Eleven dock choices compete; catalog repeats in overview and palette. Settings is a 374px viewport over 1613px content with model setup interleaved with scheduler/training. Prioritize frequent/open windows, group the full library, stabilize numeric shortcuts, and split Settings into Models/accounts, Workspace, Context/memory and Advanced/runtime. Suggested commands: shape/distill.
+6. **P2 — Theme and information integrity gaps.** Large paper surfaces bypass understory tokens, essential state labels remain 9px, and Dream shows hard-coded loss polylines with no visible illustration disclosure. Use semantic surfaces, readable functional metadata, and honest empty training observations instead of decorative evidence. Suggested commands: colorize/typeset/clarify.
+
+## Cognitive load and personas
+
+Assessment A marked five of eight cognitive checks failed: single focus, chunking, hierarchy, one-thing-at-a-time, and minimal choices. Grouping, working memory support and general progressive disclosure were strengths. Alex/power users hit an obstructed Build action and cramped transcript; Jordan/new users decode an idle cockpit and scattered setup; Sam/keyboard-low-vision users face concealed focus and pale small status text.
+
+## Detector evidence and limits
+
+CLI markup scan returned zero findings. Runtime detector produced 774 raw occurrences across Center/Chat/Dream: repeated elements, intentional occlusion of inactive windows, pinned brand treatments and generic font thresholds explain many. They are not 774 independent bugs. The confirmed technical audit consolidates six issues (0 P0, 2 P1, 4 P2). Runtime injection and overlay DOM succeeded, but user-visible browser presentation was not established; no live user overlay is claimed. Helper was stopped and existing dev server preserved. No inference, training or real conversation changes were used in the assessment.
+
+## Minor observations
+
+Browser preview footer incorrectly claims Electron Runtime; project directory is hidden in Build; role=badge is invalid; stacked optional browser preview does not scroll a newly focused window into view. Fix these within the broad pass.
+
+## Direction resolved by user authorization
+
+The user has already requested broad fixes, not an assessment-only deliverable. Proceed with task-first home, a leaner nonblocking Chat layout, frequent/current-work navigation, grouped Settings and honest readable system states. Do not reopen the already confirmed scope or visual identity. No post-fix score is assigned without a separate assessment.
+
+Detailed independent evidence: work/gui-assessment-a/assessment-a.md and work/gui-assessment-b/REPORT.md.
diff --git a/.optimize/backlog.md b/.optimize/backlog.md
index b951e07..88819f2 100644
--- a/.optimize/backlog.md
+++ b/.optimize/backlog.md
@@ -1,7 +1 @@
-# Optimization backlog
-
-| # | area | fix | evidence | impact | confidence | effort | score |
-|---|---|---|---|---:|---:|---:|---:|
-| 1 | runtime | Add a focused renderer startup/render benchmark before changing hot paths | Run 1 had only build/test/syntax timing; no direct UI runtime probe | 3 | 0.8 | 2 | 1.2 |
-| 2 | devloop | Keep the verification probes as a regression guard after the Maple/Dream host action lands | Run 1: all four probes stayed green | 2 | 0.9 | 1 | 1.8 |
-| 3 | model-runtime | Add an opt-in bounded inference throughput probe for Maple | Current pass intentionally measured `/health` and host IPC only; no completion request was sent | 4 | 0.7 | 3 | 0.9 |
+- hypotheses
diff --git a/.optimize/ledger.md b/.optimize/ledger.md
index 7bc5f01..71a2de8 100644
--- a/.optimize/ledger.md
+++ b/.optimize/ledger.md
@@ -1,39 +1,129 @@
-# Hemlock optimization ledger
-
-This ledger records measured, behavior-preserving optimization passes. The
-repository was already dirty when this first baseline was created; existing
-user work remains un-staged and uncommitted.
-
-## Run 1 — 2026-08-15 (full pass, dirty)
-
-- Applied: explicit Maple/Dream runtime launch command, fail-fast child readiness, single-flight launch handling, launch receipt tests, lazy closed/minimized surface rendering, and deferred provider CLI status checks.
-- Verification: 44 agent tests, 10 UI tests, production build, syntax, YAML, browser preview, and Electron startup passed.
-- Measured dev-loop medians from this machine:
- - build `0.462s -> 0.403s`
- - agent tests `0.389s -> 0.384s` (the after probe includes three new host/receipt tests)
- - UI tests `0.284s -> 0.243s`
- - syntax `0.111s -> 0.092s`
-- Interpretation: all probes stayed green, but no wall-clock result cleared the optimization acceptance threshold (>5% or >2s). No timed speedup is claimed. The renderer work is a structural reduction in hidden-surface construction, and the runtime work is stability-oriented.
-- Commits: baseline harness committed as `fb8e262`; application changes remain mixed with pre-existing user WIP and were not staged.
-- Next run: add a focused renderer/runtime benchmark so UI work can be accepted or rejected with a direct runtime measurement rather than build timing.
-
-## Run 2 — 2026-08-15 (Maple runtime pass, dirty)
-
-- Applied: Maple stream-frame coalescing at a 16 ms host-to-renderer boundary, synchronous terminal flushes, and child-process error/pre-readiness-exit receipts.
-- Verification: 47 agent tests, 10 UI tests, production build, syntax, YAML, diff check, and Electron desktop startup passed. Desktop startup was stopped cleanly; no Maple inference request was sent.
-- Maple health baseline: `.optimize/runs/20260815T-maple-baseline.json` measured `5.228s` to local `/health`; `.optimize/runs/20260815T-maple-after.json` measured `3.264s`. The faster after sample is recorded but not attributed because the one-shot model-load probe is cache/system-state sensitive.
-- Stream probe: 4,000 synthetic source deltas reduced to 2 IPC frames across two channels (`99.95%` fewer host-to-renderer sends); this is a dispatch-count result, not a claim about Maple generation throughput.
-- Full regression probe: build `0.366s`, agent tests `0.577s`, UI tests `0.301s`, syntax `0.118s`, Maple health `4.040s`, and stream dispatch `0.047s`; every probe passed.
-- Commits: application changes remain mixed with pre-existing user WIP and were not staged. The new measurement artifacts are present under `.optimize/` for the next focused pass.
-- Next run: if a true model-throughput claim is needed, add a bounded Maple inference benchmark only with explicit approval to invoke local inference; otherwise keep the IPC and launch receipts as the accepted host-side improvements.
-
-## Run 3 — 2026-08-15 (Maple agent reliability and artifact autopilot, dirty)
-
-- Applied: Explore/Build interaction mode, host-owned action compilation, artifact receipts, renderer preview-report handshake, static artifact verification, complete-source/patch repair inputs, bounded two-pass repair state, rollback, and retry/use-last-good UI actions.
-- Verification: 54 agent tests, 10 UI tests, production build, syntax, diff check, and the deterministic artifact autopilot fixture passed.
-- Deterministic fixture result: `completed=true`, `inferenceCalls=7`, `terminalInferenceCalls=0`, `repairCalls=2`, `repairAttempts=2`, `previewInspectionCalls=2`, `artifactRevisionCount=2`, `repositoryTouched=false`, `falseSuccess=false`.
-- Measured dev-loop medians from `.optimize/runs/20260815T171003.json`: build `0.367s`, agent tests `0.371s`, UI tests `0.235s`, syntax `0.087s`, Maple health `5.003s`, stream dispatch `0.036s`, artifact autopilot `0.183s`. All probes passed; no timed speedup is claimed because this pass is primarily a correctness/false-success reduction.
-- Commits: none; the repository contains pre-existing user WIP and the implementation remains unstaged/uncommitted.
-- Final affected-probe verification: `.optimize/runs/20260815T-final-verify.json`; build `0.364s`, agent tests `0.406s`, UI tests `0.241s`, syntax `0.089s`, artifact autopilot `0.177s`, all green. The agent-test sample moved `0.035s` on a sub-second probe, below the skill's absolute `2s` acceptance threshold, so it is not treated as a performance regression.
-- Final post-compatibility verification: `.optimize/runs/20260815T-final-verify-2.json`; build `0.428s`, agent tests `0.372s`, UI tests `0.236s`, syntax `0.090s`, artifact autopilot `0.174s`, all green. Build variance was `+0.064s` versus the prior affected run, below the absolute `2s` threshold; no timed speedup or regression is claimed.
-- Next run: exercise one real Electron preview report in the desktop runtime if a visual proof is needed; keep live Maple inference optional and measure it separately from deterministic host reliability.
+# Hemlock .optimize ledger
+
+## 2026-09-22 — Maple decode/prefill speed episode (SPEED swarm agent + integration)
+
+Workspace: /Users/ianzvirbulis/Code/Hemlock (dirty tree, multi-agent pass).
+Baseline receipt: .optimize/runs/speed-baseline.json (2026-09-21, live maple-2bit-mlx :8081).
+After receipts: .optimize/runs/speed-after-{decode,prefill,score}.json (2026-09-22, same server config).
+
+Probes (dream-chat/scripts/maple_bench.py, live server, same flags both sides):
+- decode: 30.65 tok/s -> ~64 tok/s (median 58.9-67.8 across runs) ~+2.0x
+- ttft: 392.7ms -> ~140ms (short prompt, warm) ~-60%
+- prefill: 182.6 t/s -> ~419 t/s (3165-token prompt, cold) ~+2.3x
+- score: 318ms/cand -> ~125ms/cand (/v1/score, 8 candidates) ~-60%
+- kv-bits: 57.1 t/s exact vs 72.8 t/s kv-bits=8 at 2455-tok ctx ~+27% (opt-in)
+
+Changes: mlx_lm/models/maple.py — one shared (pos,eps) array per decode step
+instead of ~24 per-layer mx.array dispatches (float(offset) on an mx.array
+forced a GPU sync every layer of every token); fused qk norm+rope path
+retained. generate.py maybe_quantize_kv_cache — only true KVCache converts
+(RotatingKVCache.to_quantized raised NotImplementedError and aborted the
+Metal command buffer — observed server kill; sliding layers now safely skip).
+server.py --kv-bits/--quantized-kv-start plumbing; lora.py os._exit(0) GIL
+finalization fix. main.cjs serverArgs: HEMLOCK_KV_BITS (default 0=off),
+HEMLOCK_KV_QUANT_START (default 5000), kvBits runtime setting (next-launch).
+
+Gates: 50/50 test_server.py, 485/485 electron, 288/291 UI (3 skip), build clean.
+Identical greedy completion between exact and kv-quant configs on same prompt.
+Negative/deferred: n-gram speculative still incompatible with SWA-512 (unchanged);
+kv-bits opt-in, 8-bit recommended over 4.
+
+## 2026-09-22 — PERF2 decode audit (continuation pass)
+
+Receipts: .optimize/runs/perf2-before.json (72.12 tok/s decode on stale
+checkpoint maple.py), perf2-after.json (88.49 tok/s; post-sync reference
+run hit 90.78 tok/s — same code, run-to-run noise band ~3%).
+
+BIG finding — checkpoint/runtime skew: the live model loads
+~/Models/Hemlock/maple-2bit-mlx/maple.py (config.json model_file +
+--trust-remote-code), NOT the repo copy. The checkpoint file predated the
+prior optimization pass, so the shared (pos,eps) fix was dormant in live
+serving. Copying the repo file into the checkpoint moved decode
+72.1 -> ~90 tok/s with no repo diff. The checkpoint is a build artifact;
+the durable fix is regenerating it through the normal packaging flow.
+Backup of the pre-sync file: maple-2bit-mlx/maple.py.bak-perf2.
+
+Decode floor analysis (standalone, plain caches, flash-head):
+~1005 graph nodes/step, ~470 real kernels. eval = 11.6ms/step.
+Breakdown: MoE block 7.25ms (gather_qmm pair dominant), attention 2.55ms,
+rest ~1.7ms. CPU build 1.4ms standalone (~7ms server-side with batch
+caches). View nodes (reshape/slice) cost ~0.85us each — not the
+bottleneck. Per-kernel dispatch ~14-30us is.
+
+Attempted, measured, REVERTED:
+- Custom fused-expert Metal kernels (2-bit dequant gemv + swiglu + weighted
+ down, simdgroup-cooperative, register-staged x, uint4 loads): correct
+ (maxdiff 0.004) but ~0.7ms vs ~0.5ms eager gather_qmm chain per layer.
+ gather_qmm is already well-tuned; a naive scalar dequant-gemv cannot beat
+ it. Would need simdgroup_matrix (tensor-core) kernels — deferred.
+- mx.compile over the whole expert block: router kernel mutates a device
+ atomic counter (last-block-finish trick) — unsafe/incorrect inside a
+ compile trace (maxdiff 0.58). Compiling only switch_mlp+aggregate works
+ and is correct but measurably neutral — not retained.
+- mx.compile replay of non-fusible op chains only saves ~11% eval time;
+ whole-step compile is blocked by Python-int cache slice indices (would
+ recompile every step).
+
+Kept (neutral-to-small, zero-risk):
+- _add_rms_norm: dropped input reshapes (kernel reads flat anyway).
+- _qk_fused: consumes the flat fused qkv row, derives grid from head
+ counts; saves a slice+reshape per layer.
+- generate.py GenerationBatch._step: B==1 sampler fast path skips the
+ per-row slice+concatenate when the batch has one request.
+
+Sampler/logit audit: make_sampler(temp=0) already returns argmax lambda;
+sampling ~0.01ms — no win available. Cache deepcopy for score forks:
+~0.04ms per 8MB — not a hot spot. Prompt cache/trie path already skips
+re-prefill on hits.
+
+Gates: 50/50 test_server.py; py_compile clean on touched files.
+Greedy equivalence: temp-0 completion for a fixed prompt is IDENTICAL
+between the pre-perf2 checkpoint maple.py and the synced repo version.
+
+## moe3 — fused MoE decode kernels (retained)
+
+Replaced the single-token MoE path (gather_qmm up/gate + SwiGLU +
+gather_qmm down + score-weighted aggregate, ~6 dispatches/layer) with two
+custom Metal kernels in maple.py (`MapleSwitchGLU.fused_decode`):
+
+- `maple_moe_upgate`: all 8 selected experts' fused up+gate projection,
+ 2-bit affine dequant in registers (grouped `s*sum(q*x)+b*sum(x)`),
+ packed uint4 activation reads, clamped SwiGLU -> expert intermediates.
+- `maple_moe_down`: 8 experts' down projection + router-score weighted
+ accumulation -> final (1,1,2048) output in one dispatch.
+- One-time `_matches` probe per layer vs the reference path; falls back to
+ gather_qmm on any mismatch/failure or non-decode shape.
+
+Correctness: fused=True on all 24 layers, maxdiff <= 0.03125 bf16,
+greedy (temp-0) completion IDENTICAL fused vs reference.
+
+Measured wins:
+- in-process A/B (alternating, same session): eager eval 15.29-15.37ms vs
+ fused 14.97-15.02ms -> ~0.35ms/step saved (~2.3%).
+- live A/B back-to-back on :8081 (checkpoint maple.py swapped, 8-rep
+ decode medians): reference 72.4 tok/s -> fused 75.7 tok/s (+3.3, ~4.6%);
+ ttft ~159ms both.
+- prefill unaffected (multi-token path untouched): 385.5 -> 404.6 tok/s.
+- score: 128.1 -> 124.2 ms/candidate.
+- standalone fused_decode 244us vs reference switch+agg 264-341us/layer.
+Receipts: .optimize/runs/moe-{before,after}-{decode,prefill,score}.json
+(+ an 8-rep fused decode taken pre-swap). prof_gpu: build 1.38ms,
+eval 15.22ms (thermally drifted session; A/B is the arbiter).
+
+Attempted, measured, REVERTED:
+- Generic dequant-GEMV (`maple_qgemv`) for attention qkv_proj/o_proj
+ M==1 calls. First version (1 row/simdgroup) lost to qmm (58 vs 45us);
+ x-amortized 4-rows/simdgroup version reached parity (~49 vs 47us qkv,
+ ~39 vs 33us o_proj) — never a win; MLX's tiled qmm already handles
+ these shapes. Reverted; no diff remains.
+- Earlier same-session fused-kernel draft (threadgroup-staged x, scalar
+ dequant): eval 15.09ms vs 13.27ms baseline — slower; reworked into the
+ register/uint4/grouped-affine version that was retained.
+
+Why the win is smaller than the 7.25ms MoE line suggests: gather_qmm's
+~250-340us/layer is only ~2-3x above the ~90us packed-weight read floor;
+dequant fma is the wall for scalar kernels, not dispatch overhead. A
+simdgroup_matrix path would still pay the same per-element unpack cost.
+
+Gates: tests/test_server.py 63/63; py_compile clean; repo and checkpoint
+maple.py byte-identical. Checkpoint backup: maple.py.bak-moe3.
diff --git a/.optimize/probe_bisect.py b/.optimize/probe_bisect.py
new file mode 100644
index 0000000..4c187ea
--- /dev/null
+++ b/.optimize/probe_bisect.py
@@ -0,0 +1,160 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+sw = model.model.layers[0].mlp.switch_mlp
+ug = sw.up_gate_proj
+
+D, M, K = 2048, 512, 8
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16)
+inds = mx.arange(8, dtype=mx.int32)
+mx.eval(x, inds)
+
+
+def timeit(name, k, inputs, grid, outshape, n=300):
+ def fn():
+ return k(inputs=inputs, grid=grid, threadgroup=(256, 1, 1),
+ output_shapes=[outshape], output_dtypes=[mx.bfloat16])[0]
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:34s} build {(t1-t0)/n*1e6:7.1f} eval {(t2-t1)/n*1e6:8.1f} us")
+
+
+# A: read weights only, write row checksum (no x, no dequant)
+kA = mx.fast.metal_kernel(
+ name="bA", input_names=["w", "inds"], output_names=["h"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u; uint lane = tid % 32u;
+ constexpr uint D = 2048u; constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ const device uint* wu = w + ((ulong)e * 2u * M + i) * WPR;
+ uint acc = 0u;
+ for (uint t = 0u; t < 4u; ++t) acc += wu[lane * 4u + t];
+ acc = simd_sum(acc);
+ if (lane == 0u && acc == 0xdeadbeefu) h[e_slot * M + i] = (bfloat16_t)1.0f;
+ }
+ """,
+)
+timeit("A: w-read only", kA, [ug.weight, inds], (256, (M//16)*8, 1), (K*M,))
+
+# B: w-read + x dot, NO dequant (treat words as floats via bitcast add)
+kB = mx.fast.metal_kernel(
+ name="bB", input_names=["x", "w", "inds"], output_names=["h"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u; uint lane = tid % 32u;
+ constexpr uint D = 2048u; constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ const device uint* wu = w + ((ulong)e * 2u * M + i) * WPR;
+ float acc = 0.0f;
+ for (uint t = 0u; t < 4u; ++t) {
+ uint widx = lane * 4u + t;
+ uint p = wu[widx];
+ for (uint c = 0u; c < 16u; ++c)
+ acc += (float)(p & 3u) * (float)x[widx * 16u + c];
+ }
+ acc = simd_sum(acc);
+ if (lane == 0u) h[e_slot * M + i] = (bfloat16_t)acc;
+ }
+ """,
+)
+timeit("B: +x dot noextract", kB, [x.reshape(-1), ug.weight, inds],
+ (256, (M//16)*8, 1), (K*M,))
+
+# C: same but x in threadgroup mem
+kC = mx.fast.metal_kernel(
+ name="bC", input_names=["x", "w", "inds"], output_names=["h"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u; uint lane = tid % 32u;
+ constexpr uint D = 2048u; constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ threadgroup float xs[D];
+ for (uint j = tid; j < D; j += 256u) xs[j] = (float)x[j];
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ const device uint* wu = w + ((ulong)e * 2u * M + i) * WPR;
+ float acc = 0.0f;
+ for (uint t = 0u; t < 4u; ++t) {
+ uint widx = lane * 4u + t;
+ uint p = wu[widx];
+ for (uint c = 0u; c < 16u; ++c)
+ acc += (float)(p & 3u) * xs[widx * 16u + c];
+ }
+ acc = simd_sum(acc);
+ if (lane == 0u) h[e_slot * M + i] = (bfloat16_t)acc;
+ }
+ """,
+)
+timeit("C: tg-x + dot", kC, [x.reshape(-1), ug.weight, inds],
+ (256, (M//16)*8, 1), (K*M,))
+
+# D: full dequant (scale/bias) on one row only per sg
+kD = mx.fast.metal_kernel(
+ name="bD", input_names=["x", "w", "sc", "bi", "inds"], output_names=["h"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u; uint lane = tid % 32u;
+ constexpr uint D = 2048u; constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ constexpr uint GPW = 8u; constexpr uint NG = 16u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ ulong row = (ulong)e * 2u * M + i;
+ const device uint* wu = w + row * WPR;
+ ulong sb = row * NG;
+ float acc = 0.0f;
+ for (uint t = 0u; t < 4u; ++t) {
+ uint widx = lane * 4u + t;
+ uint g = widx / GPW;
+ float s = (float)sc[sb + g];
+ float b = (float)bi[sb + g];
+ uint p = wu[widx];
+ for (uint c = 0u; c < 16u; ++c)
+ acc += ((float)((p >> (2u*c)) & 3u) * s + b)
+ * (float)x[widx * 16u + c];
+ }
+ acc = simd_sum(acc);
+ if (lane == 0u) h[e_slot * M + i] = (bfloat16_t)acc;
+ }
+ """,
+)
+timeit("D: full dequant 1row", kD,
+ [x.reshape(-1), ug.weight, ug.scales, ug.biases, inds],
+ (256, (M//16)*8, 1), (K*M,))
diff --git a/.optimize/probe_bisect2.py b/.optimize/probe_bisect2.py
new file mode 100644
index 0000000..6e3d588
--- /dev/null
+++ b/.optimize/probe_bisect2.py
@@ -0,0 +1,131 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+sw = model.model.layers[0].mlp.switch_mlp
+ug = sw.up_gate_proj
+
+D, M, K = 2048, 512, 8
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16)
+inds = mx.arange(8, dtype=mx.int32)
+mx.eval(x, inds)
+xf = x.reshape(-1)
+
+# check bf16 unpack: bfloat2 type exists?
+ktest = mx.fast.metal_kernel(
+ name="t_bf2", input_names=["x"], output_names=["o"],
+ source="""
+ uint tid = thread_position_in_grid.x;
+ uint u = ((const constant uint*)x)[tid];
+ bfloat2 v = as_type(u);
+ o[tid * 2u] = (float)v.x;
+ o[tid * 2u + 1u] = (float)v.y;
+ """,
+)
+xx = mx.array([1.0, 2.0, 3.0, 4.0], dtype=mx.bfloat16)
+o = ktest(inputs=[xx], grid=(2,1,1), threadgroup=(2,1,1),
+ output_shapes=[(4,)], output_dtypes=[mx.float32])[0]
+mx.eval(o)
+print("bfloat2 unpack:", o.tolist())
+
+
+def timeit(name, k, inputs, grid, outshape, n=300):
+ def fn():
+ return k(inputs=inputs, grid=grid, threadgroup=(256, 1, 1),
+ output_shapes=[outshape], output_dtypes=[mx.bfloat16])[0]
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:34s} build {(t1-t0)/n*1e6:7.1f} eval {(t2-t1)/n*1e6:8.1f} us")
+
+
+# E1: full up+gate rows, uint4 x loads (bfloat2 unpack), per-lane group refactor
+kE1 = mx.fast.metal_kernel(
+ name="bE1", input_names=["x", "w", "sc", "bi", "inds"], output_names=["h"],
+ header="""
+ inline float2 bf2(uint u) {
+ bfloat2 v = as_type(u);
+ return float2((float)v.x, (float)v.y);
+ }
+ """,
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u; uint lane = tid % 32u;
+ constexpr uint D = 2048u; constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ constexpr uint NG = D / 128u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ const device uint4* x4 = (const device uint4*)x;
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ ulong up_row = (ulong)e * 2u * M + i;
+ const device uint* wu = w + up_row * WPR;
+ const device uint* wg = wu + (ulong)M * WPR;
+ ulong sbu = up_row * NG;
+ ulong sbg = sbu + (ulong)M * NG;
+ float qu = 0.0f, qg = 0.0f;
+ uint widx = lane * 4u;
+ for (uint t = 0u; t < 4u; ++t) {
+ uint4 xa = x4[(widx + t) * 2u];
+ uint4 xb = x4[(widx + t) * 2u + 1u];
+ float2 f0 = bf2(xa.x), f1 = bf2(xa.y), f2 = bf2(xa.z), f3 = bf2(xa.w);
+ float2 f4 = bf2(xb.x), f5 = bf2(xb.y), f6 = bf2(xb.z), f7 = bf2(xb.w);
+ uint pu = wu[widx + t];
+ uint pg = wg[widx + t];
+ qu += (float)(pu & 3u) * f0.x; qu += (float)((pu >> 2u) & 3u) * f0.y;
+ qu += (float)((pu >> 4u) & 3u) * f1.x; qu += (float)((pu >> 6u) & 3u) * f1.y;
+ qu += (float)((pu >> 8u) & 3u) * f2.x; qu += (float)((pu >> 10u) & 3u) * f2.y;
+ qu += (float)((pu >> 12u) & 3u) * f3.x; qu += (float)((pu >> 14u) & 3u) * f3.y;
+ qu += (float)((pu >> 16u) & 3u) * f4.x; qu += (float)((pu >> 18u) & 3u) * f4.y;
+ qu += (float)((pu >> 20u) & 3u) * f5.x; qu += (float)((pu >> 22u) & 3u) * f5.y;
+ qu += (float)((pu >> 24u) & 3u) * f6.x; qu += (float)((pu >> 26u) & 3u) * f6.y;
+ qu += (float)((pu >> 28u) & 3u) * f7.x; qu += (float)((pu >> 30u) & 3u) * f7.y;
+ qg += (float)(pg & 3u) * f0.x; qg += (float)((pg >> 2u) & 3u) * f0.y;
+ qg += (float)((pg >> 4u) & 3u) * f1.x; qg += (float)((pg >> 6u) & 3u) * f1.y;
+ qg += (float)((pg >> 8u) & 3u) * f2.x; qg += (float)((pg >> 10u) & 3u) * f2.y;
+ qg += (float)((pg >> 12u) & 3u) * f3.x; qg += (float)((pg >> 14u) & 3u) * f3.y;
+ qg += (float)((pg >> 16u) & 3u) * f4.x; qg += (float)((pg >> 18u) & 3u) * f4.y;
+ qg += (float)((pg >> 20u) & 3u) * f5.x; qg += (float)((pg >> 22u) & 3u) * f5.y;
+ qg += (float)((pg >> 24u) & 3u) * f6.x; qg += (float)((pg >> 26u) & 3u) * f6.y;
+ qg += (float)((pg >> 28u) & 3u) * f7.x; qg += (float)((pg >> 30u) & 3u) * f7.y;
+ }
+ qu += simd_shuffle_down(qu, 1u);
+ qg += simd_shuffle_down(qg, 1u);
+ float cu = 0.0f, cg = 0.0f;
+ // group = lane/2 needs per-group x sums; approximate via full dot
+ // (we fold bias by precomputed group sums is complex; keep per-group
+ // scale on the q-sum and a second pass for bias)
+ if ((lane & 1u) == 0u) {
+ uint g = lane >> 1;
+ // bias needs sum(x in group) -- compute via reduction instead:
+ cu = qu * (float)sc[sbu + g];
+ cg = qg * (float)sc[sbg + g];
+ }
+ float accu = simd_sum(cu);
+ float accg = simd_sum(cg);
+ if (lane == 0u) {
+ h[e_slot * M + i] = (bfloat16_t)(accu + accg);
+ }
+ }
+ """,
+)
+# NOTE: E1 intentionally drops bias term (test-only) -- speed probe only
+timeit("E1: uint4x + groupsum(noB)", kE1,
+ [xf, ug.weight, ug.scales, ug.biases, inds],
+ (256, (M//16)*8, 1), (K*M,))
diff --git a/.optimize/probe_bw.py b/.optimize/probe_bw.py
new file mode 100644
index 0000000..451f80f
--- /dev/null
+++ b/.optimize/probe_bw.py
@@ -0,0 +1,191 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+from mlx_lm.models.cache import make_prompt_cache
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+inner = model.model
+sw = inner.layers[0].mlp.switch_mlp
+ug, dn = sw.up_gate_proj, sw.down_proj
+print("ug.weight", ug.weight.shape, ug.weight.dtype, "contig")
+
+D, M, K, gs = 2048, 512, 8, 128
+
+# Kernel A: pure weight-streaming read — one tg per expert-row-block,
+# lanes stride words, accumulate raw packed ints. Upper bound on speed.
+kread = mx.fast.metal_kernel(
+ name="probe_rd",
+ input_names=["w", "inds"],
+ output_names=["out"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ // each tg covers 32 rows of one expert
+ uint e = (uint)inds[tgid / (M / 32u)];
+ uint row0 = (tgid % (M / 32u)) * 32u;
+ const device uint* base = w + ((ulong)e * 2u * M + row0) * WPR;
+ uint acc = 0u;
+ for (uint r = 0u; r < 32u; ++r) {
+ for (uint widx = tid; widx < WPR; widx += 256u)
+ acc += base[r * WPR + widx];
+ }
+ if (acc == 0xdeadbeefu) out[0] = 1.0f;
+ """,
+)
+
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16)
+inds = mx.arange(8, dtype=mx.int32)
+mx.eval(x, inds)
+
+
+def timeit(name, fn, n=300):
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ gb = 8 * M * D * 2 / 8 # 8 experts x M rows x D elems x 2bit
+ print(f"{name:30s} build {(t1-t0)/n*1e6:7.1f} eval {(t2-t1)/n*1e6:8.1f} us -> {gb/(t2-t1)/n/1e3:.0f} GB/s eff")
+
+
+timeit("read 8ex x 512rows(up)", lambda: kread(
+ inputs=[ug.weight, inds],
+ grid=(256, (M // 32) * 8, 1),
+ threadgroup=(256, 1, 1),
+ output_shapes=[(1,)], output_dtypes=[mx.float32],
+)[0])
+
+# Kernel B: same but dequant + x dot (no tg staging) — x read from device
+kdeq = mx.fast.metal_kernel(
+ name="probe_deq",
+ input_names=["x", "w", "sc", "bi", "inds"],
+ output_names=["h"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u;
+ uint lane = tid % 32u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ constexpr uint WPL = WPR / 32u;
+ constexpr uint GPW = 128u / 16u;
+ constexpr uint NG = D / 128u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ ulong up_row = (ulong)e * 2u * M + i;
+ const device uint* wu = w + up_row * WPR;
+ const device uint* wg = wu + (ulong)M * WPR;
+ ulong sbu = up_row * NG;
+ ulong sbg = sbu + (ulong)M * NG;
+ float accu = 0.0f, accg = 0.0f;
+ for (uint t = 0u; t < WPL; ++t) {
+ uint widx = lane * WPL + t;
+ uint g = widx / GPW;
+ float s_u = (float)sc[sbu + g];
+ float b_u = (float)bi[sbu + g];
+ float s_g = (float)sc[sbg + g];
+ float b_g = (float)bi[sbg + g];
+ uint pu = wu[widx];
+ uint pg = wg[widx];
+ for (uint c = 0u; c < 16u; ++c) {
+ float xu = (float)x[widx * 16u + c];
+ accu += ((float)((pu >> (2u * c)) & 3u) * s_u + b_u) * xu;
+ accg += ((float)((pg >> (2u * c)) & 3u) * s_g + b_g) * xu;
+ }
+ }
+ accu = simd_sum(accu);
+ accg = simd_sum(accg);
+ if (lane == 0u) {
+ float gv = metal::min(accg, 7.0f);
+ float uv = metal::clamp(accu, -7.0f, 7.0f);
+ h[e_slot * M + i] = (bfloat16_t)(gv / (1.0f + metal::exp(-gv)) * uv);
+ }
+ }
+ """,
+)
+timeit("deq gemv no-tgx", lambda: kdeq(
+ inputs=[x.reshape(-1), ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1),
+ threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16],
+)[0])
+
+# Kernel C: tg staging but VECTOR loads (uint4 = 4 words per load)
+kvec = mx.fast.metal_kernel(
+ name="probe_vec",
+ input_names=["x", "w", "sc", "bi", "inds"],
+ output_names=["h"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u;
+ uint lane = tid % 32u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ constexpr uint GPW = 128u / 16u;
+ constexpr uint NG = D / 128u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ ulong up_row = (ulong)e * 2u * M + i;
+ const device uint4* wu = (const device uint4*)(w + up_row * WPR);
+ const device uint4* wg = (const device uint4*)(w + (up_row + M) * WPR);
+ ulong sbu = up_row * NG;
+ ulong sbg = sbu + (ulong)M * NG;
+ float accu = 0.0f, accg = 0.0f;
+ // 128 words/row, uint4 -> 32 vec words, one per lane
+ {
+ uint g = (lane * 4u) / GPW;
+ float s_u = (float)sc[sbu + g];
+ float b_u = (float)bi[sbu + g];
+ float s_g = (float)sc[sbg + g];
+ float b_g = (float)bi[sbg + g];
+ uint4 pu = wu[lane];
+ uint4 pg = wg[lane];
+ uint pua[4] = {pu.x, pu.y, pu.z, pu.w};
+ uint pga[4] = {pg.x, pg.y, pg.z, pg.w};
+ for (uint t = 0u; t < 4u; ++t) {
+ uint widx = lane * 4u + t;
+ for (uint c = 0u; c < 16u; ++c) {
+ float xu = (float)x[widx * 16u + c];
+ accu += ((float)((pua[t] >> (2u * c)) & 3u) * s_u + b_u) * xu;
+ accg += ((float)((pga[t] >> (2u * c)) & 3u) * s_g + b_g) * xu;
+ }
+ }
+ }
+ accu = simd_sum(accu);
+ accg = simd_sum(accg);
+ if (lane == 0u) {
+ float gv = metal::min(accg, 7.0f);
+ float uv = metal::clamp(accu, -7.0f, 7.0f);
+ h[e_slot * M + i] = (bfloat16_t)(gv / (1.0f + metal::exp(-gv)) * uv);
+ }
+ }
+ """,
+)
+timeit("deq gemv uint4", lambda: kvec(
+ inputs=[x.reshape(-1), ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1),
+ threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16],
+)[0])
diff --git a/.optimize/probe_bw2.py b/.optimize/probe_bw2.py
new file mode 100644
index 0000000..fb26a68
--- /dev/null
+++ b/.optimize/probe_bw2.py
@@ -0,0 +1,175 @@
+import sys, time
+import mlx.core as mx
+import mlx.nn as nn
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+inner = model.model
+sw = inner.layers[0].mlp.switch_mlp
+ug, dn = sw.up_gate_proj, sw.down_proj
+
+D, M, K, gs = 2048, 512, 8, 128
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16)
+inds = mx.arange(8, dtype=mx.int32)
+mx.eval(x, inds)
+xf = x.reshape(-1)
+
+HDR = """
+#define DQ16(P, ACC, X0, X1, X2, X3) { \
+ ACC += (float)((P) & 3u) * (X0).x; \
+ ACC += (float)(((P) >> 2u) & 3u) * (X0).y; \
+ ACC += (float)(((P) >> 4u) & 3u) * (X0).z; \
+ ACC += (float)(((P) >> 6u) & 3u) * (X0).w; \
+ ACC += (float)(((P) >> 8u) & 3u) * (X1).x; \
+ ACC += (float)(((P) >> 10u) & 3u) * (X1).y; \
+ ACC += (float)(((P) >> 12u) & 3u) * (X1).z; \
+ ACC += (float)(((P) >> 14u) & 3u) * (X1).w; \
+ ACC += (float)(((P) >> 16u) & 3u) * (X2).x; \
+ ACC += (float)(((P) >> 18u) & 3u) * (X2).y; \
+ ACC += (float)(((P) >> 20u) & 3u) * (X2).z; \
+ ACC += (float)(((P) >> 22u) & 3u) * (X2).w; \
+ ACC += (float)(((P) >> 24u) & 3u) * (X3).x; \
+ ACC += (float)(((P) >> 26u) & 3u) * (X3).y; \
+ ACC += (float)(((P) >> 28u) & 3u) * (X3).z; \
+ ACC += (float)(((P) >> 30u) & 3u) * (X3).w; }
+"""
+
+k2 = mx.fast.metal_kernel(
+ name="probe_v2",
+ input_names=["x", "w", "sc", "bi", "inds"],
+ output_names=["h"],
+ header=HDR,
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u;
+ uint lane = tid % 32u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u;
+ constexpr uint NG = D / 128u;
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+
+ threadgroup float xs[D];
+ threadgroup float xg[NG];
+ for (uint j = tid; j < D; j += 256u) xs[j] = (float)x[j];
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (tid < NG) {
+ float s = 0.0f;
+ for (uint j = 0u; j < 128u; ++j) s += xs[tid * 128u + j];
+ xg[tid] = s;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ const threadgroup float4* xs4 = (const threadgroup float4*)xs;
+
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ ulong up_row = (ulong)e * 2u * M + i;
+ const device uint* wu = w + up_row * WPR;
+ const device uint* wg = wu + (ulong)M * WPR;
+ ulong sbu = up_row * NG;
+ ulong sbg = sbu + (ulong)M * NG;
+ float qu = 0.0f, qg = 0.0f;
+ uint widx = lane * 4u;
+ {
+ float4 x0 = xs4[widx * 4u];
+ float4 x1 = xs4[widx * 4u + 1u];
+ float4 x2 = xs4[widx * 4u + 2u];
+ float4 x3 = xs4[widx * 4u + 3u];
+ uint pu = wu[widx];
+ uint pg = wg[widx];
+ DQ16(pu, qu, x0, x1, x2, x3)
+ DQ16(pg, qg, x0, x1, x2, x3)
+ }
+ {
+ float4 x0 = xs4[(widx + 1u) * 4u];
+ float4 x1 = xs4[(widx + 1u) * 4u + 1u];
+ float4 x2 = xs4[(widx + 1u) * 4u + 2u];
+ float4 x3 = xs4[(widx + 1u) * 4u + 3u];
+ uint pu = wu[widx + 1u];
+ uint pg = wg[widx + 1u];
+ DQ16(pu, qu, x0, x1, x2, x3)
+ DQ16(pg, qg, x0, x1, x2, x3)
+ }
+ {
+ float4 x0 = xs4[(widx + 2u) * 4u];
+ float4 x1 = xs4[(widx + 2u) * 4u + 1u];
+ float4 x2 = xs4[(widx + 2u) * 4u + 2u];
+ float4 x3 = xs4[(widx + 2u) * 4u + 3u];
+ uint pu = wu[widx + 2u];
+ uint pg = wg[widx + 2u];
+ DQ16(pu, qu, x0, x1, x2, x3)
+ DQ16(pg, qg, x0, x1, x2, x3)
+ }
+ {
+ float4 x0 = xs4[(widx + 3u) * 4u];
+ float4 x1 = xs4[(widx + 3u) * 4u + 1u];
+ float4 x2 = xs4[(widx + 3u) * 4u + 2u];
+ float4 x3 = xs4[(widx + 3u) * 4u + 3u];
+ uint pu = wu[widx + 3u];
+ uint pg = wg[widx + 3u];
+ DQ16(pu, qu, x0, x1, x2, x3)
+ DQ16(pg, qg, x0, x1, x2, x3)
+ }
+ qu += simd_shuffle_down(qu, 1u);
+ qg += simd_shuffle_down(qg, 1u);
+ float cu = 0.0f, cg = 0.0f;
+ if ((lane & 1u) == 0u) {
+ uint g = lane >> 1;
+ cu = qu * (float)sc[sbu + g] + (float)bi[sbu + g] * xg[g];
+ cg = qg * (float)sc[sbg + g] + (float)bi[sbg + g] * xg[g];
+ }
+ float accu = simd_sum(cu);
+ float accg = simd_sum(cg);
+ if (lane == 0u) {
+ float gv = metal::min(accg, 7.0f);
+ float uv = metal::clamp(accu, -7.0f, 7.0f);
+ h[e_slot * M + i] = (bfloat16_t)(gv / (1.0f + metal::exp(-gv)) * uv);
+ }
+ }
+ """,
+)
+
+
+def timeit(name, fn, n=300):
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:30s} build {(t1-t0)/n*1e6:7.1f} eval {(t2-t1)/n*1e6:8.1f} us")
+
+
+h = k2(
+ inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1),
+ threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16],
+)[0]
+w = ug.weight; sc = ug.scales; bi = ug.biases
+e = int(inds[0])
+wdeq = mx.dequantize(w[e], sc[e], bi[e], group_size=128, bits=2).astype(mx.float32)
+xdot = wdeq @ x.reshape(-1).astype(mx.float32)
+up = xdot[:M]; gate = xdot[M:]
+h_ref = nn.silu(mx.minimum(gate, 7.0)) * mx.clip(up, -7.0, 7.0)
+mx.eval(h, h_ref)
+print("v2 maxdiff:", float(mx.abs(h[:M].astype(mx.float32) - h_ref).max()))
+
+timeit("v2 grouped+vec", lambda: k2(
+ inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1),
+ threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16],
+)[0])
diff --git a/.optimize/probe_compile.py b/.optimize/probe_compile.py
new file mode 100644
index 0000000..2f4bfc7
--- /dev/null
+++ b/.optimize/probe_compile.py
@@ -0,0 +1,44 @@
+import time, sys
+import mlx.core as mx
+import mlx.nn as nn
+from functools import partial
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.models.maple import (
+ MapleSparseMoeBlock, ModelArgs, aggregate_expert_outputs,
+)
+
+args = ModelArgs(
+ hidden_size=512, intermediate_size=1024, moe_intermediate_size=128,
+ num_hidden_layers=4, num_attention_heads=8, num_key_value_heads=2,
+ head_dim=64, num_experts=32, num_experts_per_tok=4, sliding_window=64,
+)
+blk = MapleSparseMoeBlock(args)
+mx.eval(blk.parameters())
+
+x = mx.random.normal((1, 1, 512), dtype=mx.bfloat16, key=mx.random.key(0))
+ref = blk(x)
+mx.eval(ref)
+
+def moe_block(x):
+ inds, scores = blk.gate(x)
+ y = blk.switch_mlp(x, inds)
+ return aggregate_expert_outputs(y, scores)
+
+comp = mx.compile(moe_block)
+try:
+ out = comp(x)
+ mx.eval(out)
+ print("compile ok, match:", bool(mx.allclose(out.astype(mx.float32), ref.astype(mx.float32), atol=2e-2)))
+except Exception as e:
+ import traceback; traceback.print_exc()
+
+for name, fn in [("ref", lambda: blk(x)), ("compiled", lambda: comp(x))]:
+ for _ in range(20):
+ r = fn()
+ mx.eval(r)
+ t = time.time()
+ for _ in range(200):
+ r = fn()
+ mx.eval(r)
+ print(name, "%.2f us/call-build" % ((time.time() - t) / 200 * 1e6))
diff --git a/.optimize/probe_copy.py b/.optimize/probe_copy.py
new file mode 100644
index 0000000..c2f4cb1
--- /dev/null
+++ b/.optimize/probe_copy.py
@@ -0,0 +1,43 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+
+# Does metal_kernel copy big inputs? Kernel that reads only w[0].
+k = mx.fast.metal_kernel(
+ name="probe_one",
+ input_names=["w"],
+ output_names=["out"],
+ source="""
+ uint tid = thread_position_in_grid.x;
+ if (tid == 0u) out[0] = (float)w[0];
+ """,
+)
+
+w_small = mx.zeros((1024,), dtype=mx.uint32)
+w_big = mx.zeros((256, 1024, 128), dtype=mx.uint32) # 128MB, contiguous
+mx.eval(w_small, w_big)
+
+# and a non-contiguous big view
+w_big_nc = mx.zeros((512, 1024, 128), dtype=mx.uint32)[::2]
+mx.eval(w_big_nc)
+
+
+def timeit(name, w, n=100):
+ def fn():
+ return k(inputs=[w], grid=(32, 1, 1), threadgroup=(32, 1, 1),
+ output_shapes=[(1,)], output_dtypes=[mx.float32])[0]
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:24s} build {(t1-t0)/n*1e6:7.1f} us eval {(t2-t1)/n*1e6:9.1f} us")
+
+
+timeit("small (4KB)", w_small)
+timeit("big contiguous (128MB)", w_big)
+timeit("big noncontig view", w_big_nc)
diff --git a/.optimize/probe_kernels.py b/.optimize/probe_kernels.py
new file mode 100644
index 0000000..cc8c63e
--- /dev/null
+++ b/.optimize/probe_kernels.py
@@ -0,0 +1,78 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+from mlx_lm.models.cache import make_prompt_cache
+from mlx_lm.models.maple import (
+ aggregate_expert_outputs, _moe_upgate_kernel, _moe_down_kernel,
+)
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+inner = model.model
+l = inner.layers[0]
+blk = l.mlp
+sw = blk.switch_mlp
+ug, dn = sw.up_gate_proj, sw.down_proj
+
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16, key=mx.random.key(0))
+inds, scores = blk.gate(x)
+mx.eval(inds, scores)
+inds_f = inds.reshape(-1)
+scores_f = scores.reshape(-1).astype(mx.float32)
+xf = x.reshape(-1)
+K, D, M, gs = 8, 2048, 512, 128
+
+
+def k1():
+ return _moe_upgate_kernel(
+ inputs=[xf, ug.weight, ug.scales, ug.biases, inds_f],
+ template=[("T_", x.dtype), ("DIM", D), ("MDIM", M), ("GSIZE", gs)],
+ grid=(256, (K * M) // 16, 1),
+ threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)],
+ output_dtypes=[x.dtype],
+ )[0]
+
+
+h = k1()
+
+
+def k2():
+ return _moe_down_kernel(
+ inputs=[h, dn.weight, dn.scales, dn.biases, inds_f, scores_f],
+ template=[("T_", x.dtype), ("KEXP", K), ("DIM", D), ("MDIM", M), ("GSIZE", dn.group_size)],
+ grid=(256, D // 8, 1),
+ threadgroup=(256, 1, 1),
+ output_shapes=[(D,)],
+ output_dtypes=[x.dtype],
+ )[0]
+
+
+def ref():
+ return aggregate_expert_outputs(sw(x, inds), scores)
+
+
+def timeit(name, fn, n=400):
+ for _ in range(10):
+ r = fn()
+ mx.eval(r)
+ # build only
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:34s} build {(t1-t0)/n*1e6:7.1f} us eval {(t2-t1)/n*1e6:7.1f} us")
+
+
+timeit("k1 upgate+swiglu", k1)
+timeit("k2 down+agg", k2)
+timeit("k1+k2", lambda: (k1(), k2()))
+timeit("reference switch+agg", ref)
+timeit("fused_decode", lambda: sw.fused_decode(x, inds, scores))
diff --git a/.optimize/probe_moe.py b/.optimize/probe_moe.py
new file mode 100644
index 0000000..b31cf98
--- /dev/null
+++ b/.optimize/probe_moe.py
@@ -0,0 +1,34 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+from mlx_lm.models.cache import make_prompt_cache
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+prompt = mx.array(tokenizer.encode("List the first six prime numbers."))
+cache = make_prompt_cache(model)
+logits = model(prompt[None, :-1], cache=cache)
+mx.eval(logits)
+y = prompt[-1:]
+for _ in range(4):
+ logits = model(y[None], cache=cache)
+ mx.eval(logits)
+ y = mx.argmax(logits[:, -1, :], axis=-1)
+
+inner = model.model if hasattr(model, "model") else model
+layers = inner.layers
+blk = layers[0].mlp
+print("gate fused:", blk.gate._fused)
+print("switch fused_decode:", blk.switch_mlp._fused_decode)
+n_fused = sum(
+ 1
+ for l in layers
+ if getattr(l.mlp.switch_mlp, "_fused_decode", None) is True
+)
+print("layers with fused moe:", n_fused, "/", len(layers))
diff --git a/.optimize/probe_moe_detail.py b/.optimize/probe_moe_detail.py
new file mode 100644
index 0000000..b36db3d
--- /dev/null
+++ b/.optimize/probe_moe_detail.py
@@ -0,0 +1,78 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+from mlx_lm.models.cache import make_prompt_cache
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+prompt = mx.array(tokenizer.encode("List the first six prime numbers."))
+cache = make_prompt_cache(model)
+logits = model(prompt[None, :-1], cache=cache)
+mx.eval(logits)
+
+inner = model.model if hasattr(model, "model") else model
+layer = inner.layers[0]
+mlp = layer.mlp
+sw = mlp.switch_mlp
+ug = sw.up_gate_proj
+dn = sw.down_proj
+
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16, key=mx.random.key(0))
+inds, scores = mlp.gate(x)
+mx.eval(inds, scores)
+xe = mx.expand_dims(x, (-2, -3))
+h = mx.random.normal((1, 1, 1, 8, 512), dtype=mx.bfloat16)
+
+
+def timeit(name, fn, n=300):
+ """Throughput: issue n calls, one eval -> amortized GPU cost per call."""
+ for _ in range(10):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ mx.eval(rs)
+ dt = (time.perf_counter() - t0) / n * 1e3
+ print(f"{name:44s} {dt*1e3:9.1f} us")
+ return dt
+
+
+timeit("mlp block (gate+switch+agg)", lambda: mlp(x))
+timeit(" gate (fused router)", lambda: mlp.gate(x))
+timeit(" up_gate gather_qmm 8e", lambda: ug(xe, inds))
+timeit(" down gather_qmm 8e", lambda: dn(h, inds))
+timeit(" switch_mlp total", lambda: sw(x, inds))
+yo = sw(x, inds)
+timeit(
+ " aggregate",
+ lambda: (yo.astype(mx.float32) * scores[..., None])
+ .sum(axis=-2)
+ .astype(yo.dtype),
+)
+
+# comparisons
+one_idx = mx.zeros((1, 1, 1, 1), dtype=mx.int32)
+timeit(" up_gate qmm 1 expert", lambda: ug(xe, one_idx))
+idx16 = mx.concatenate([inds, inds], axis=-1) # 16 rows, 8 unique
+timeit(" up_gate qmm 16 rows", lambda: ug(mx.broadcast_to(xe, (1,1,2,8,2048)) if False else mx.concatenate([xe, xe], axis=2), idx16))
+
+# attention
+attn = layer.self_attn
+from mlx_lm.models.cache import RotatingKVCache
+c2 = RotatingKVCache(max_size=512)
+xa = mx.random.normal((1, 20, 2048), dtype=mx.bfloat16)
+attn(xa, cache=c2)
+mx.eval(c2.keys)
+timeit("attention L=1 ctx~20", lambda: attn(x[:, :1], cache=c2))
+
+c_all = make_prompt_cache(model)
+model(prompt[None, :-1], cache=c_all)
+mx.eval([c.keys for c in c_all if c is not None])
+y = mx.array([1])
+timeit("full decode step", lambda: model(y[None], cache=c_all), n=100)
diff --git a/.optimize/probe_stream.py b/.optimize/probe_stream.py
new file mode 100644
index 0000000..17ed42d
--- /dev/null
+++ b/.optimize/probe_stream.py
@@ -0,0 +1,74 @@
+import sys, time
+import mlx.core as mx
+
+# Raw streaming bandwidth: read a big buffer with different per-thread loads.
+w = mx.zeros((64 * 1024 * 1024,), dtype=mx.uint32) # 256MB
+mx.eval(w)
+
+HDR = ""
+
+
+def mk(name, body):
+ return mx.fast.metal_kernel(
+ name=name, input_names=["w"], output_names=["o"], source=body)
+
+
+k_scalar = mk("s_scalar", """
+ uint tid = thread_position_in_grid.x;
+ uint acc = 0u;
+ for (uint i = 0u; i < 64u; ++i) acc += w[tid * 64u + i];
+ if (acc == 0xdeadbeefu) o[0] = 1.0f;
+""")
+
+k_u4 = mk("s_u4", """
+ uint tid = thread_position_in_grid.x;
+ const device uint4* w4 = (const device uint4*)w;
+ uint acc = 0u;
+ for (uint i = 0u; i < 16u; ++i) {
+ uint4 v = w4[tid * 16u + i];
+ acc += v.x + v.y + v.z + v.w;
+ }
+ if (acc == 0xdeadbeefu) o[0] = 1.0f;
+""")
+
+# strided: thread reads w[tid + i*N] - coalesced across lanes
+k_stride = mk("s_stride", """
+ uint tid = thread_position_in_grid.x;
+ uint n = threads_per_grid.x;
+ uint acc = 0u;
+ for (uint i = 0u; i < 64u; ++i) acc += w[tid + i * n];
+ if (acc == 0xdeadbeefu) o[0] = 1.0f;
+""")
+
+k_stride4 = mk("s_stride4", """
+ uint tid = thread_position_in_grid.x;
+ uint n = threads_per_grid.x;
+ const device uint4* w4 = (const device uint4*)w;
+ uint acc = 0u;
+ for (uint i = 0u; i < 16u; ++i) {
+ uint4 v = w4[tid + i * n];
+ acc += v.x + v.y + v.z + v.w;
+ }
+ if (acc == 0xdeadbeefu) o[0] = 1.0f;
+""")
+
+
+def timeit(name, k, nthreads, n=50):
+ def fn():
+ return k(inputs=[w], grid=(nthreads, 1, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(1,)], output_dtypes=[mx.float32])[0]
+ for _ in range(3):
+ fn()
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ mx.eval(rs)
+ dt = (time.perf_counter() - t0) / n
+ print(f"{name:28s} {dt*1e6:8.1f} us -> {256e6/dt/1e9:6.0f} GB/s")
+
+
+timeit("scalar x64", k_scalar, 1024 * 1024)
+timeit("uint4 x16", k_u4, 1024 * 1024)
+timeit("strided scalar x64", k_stride, 1024 * 1024)
+timeit("strided uint4 x16", k_stride4, 1024 * 1024)
diff --git a/.optimize/probe_v3.py b/.optimize/probe_v3.py
new file mode 100644
index 0000000..cf6996d
--- /dev/null
+++ b/.optimize/probe_v3.py
@@ -0,0 +1,247 @@
+import sys, time
+import mlx.core as mx
+import mlx.nn as nn
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+sw = model.model.layers[0].mlp.switch_mlp
+ug, dn = sw.up_gate_proj, sw.down_proj
+
+D, M, K, gs = 2048, 512, 8, 128
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16)
+h_in = mx.random.normal((8, 512), dtype=mx.bfloat16)
+inds = mx.array([3, 17, 42, 77, 100, 150, 201, 255], dtype=mx.int32)
+scores = mx.array([0.2, 0.15, 0.1, 0.15, 0.1, 0.1, 0.1, 0.1], dtype=mx.float32)
+mx.eval(x, h_in, inds, scores)
+xf = x.reshape(-1)
+
+# ---------------- kernel 1 v3 ----------------
+k1 = mx.fast.metal_kernel(
+ name="v3_upgate",
+ input_names=["x", "w", "sc", "bi", "inds"],
+ output_names=["h"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u;
+ uint lane = tid % 32u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u; // 128
+ constexpr uint NG = D / 128u; // 16
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+
+ // transposed x staging: element of lane l at local index j -> bank l
+ threadgroup float xs[64 * 32];
+ threadgroup float xg[NG];
+ for (uint idx = tid; idx < D; idx += 256u) {
+ uint l = idx / 64u;
+ uint j = idx % 64u;
+ xs[j * 32u + l] = (float)x[idx];
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (tid < NG) {
+ // group g = lanes 2g,2g+1 all j
+ float s = 0.0f;
+ for (uint j = 0u; j < 64u; ++j)
+ s += xs[j * 32u + 2u * tid] + xs[j * 32u + 2u * tid + 1u];
+ xg[tid] = s;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ ulong up_row = (ulong)e * 2u * M + i;
+ const device uint4* wu = (const device uint4*)(w + up_row * WPR);
+ const device uint4* wg = (const device uint4*)(w + (up_row + M) * WPR);
+ ulong sbu = up_row * NG;
+ ulong sbg = sbu + (ulong)M * NG;
+ uint4 pu = wu[lane]; // words lane*4..lane*4+3
+ uint4 pg = wg[lane];
+ float qu = 0.0f, qg = 0.0f;
+ {
+ uint pw = pu.x, gw = pg.x;
+ const threadgroup float* xe = xs + lane;
+ for (uint c = 0u; c < 16u; ++c) {
+ float xv = xe[(0u * 16u + c) * 32u];
+ qu += (float)(pw & 3u) * xv; pw >>= 2u;
+ qg += (float)(gw & 3u) * xv; gw >>= 2u;
+ }
+ }
+ {
+ uint pw = pu.y, gw = pg.y;
+ const threadgroup float* xe = xs + lane;
+ for (uint c = 0u; c < 16u; ++c) {
+ float xv = xe[(1u * 16u + c) * 32u];
+ qu += (float)(pw & 3u) * xv; pw >>= 2u;
+ qg += (float)(gw & 3u) * xv; gw >>= 2u;
+ }
+ }
+ {
+ uint pw = pu.z, gw = pg.z;
+ const threadgroup float* xe = xs + lane;
+ for (uint c = 0u; c < 16u; ++c) {
+ float xv = xe[(2u * 16u + c) * 32u];
+ qu += (float)(pw & 3u) * xv; pw >>= 2u;
+ qg += (float)(gw & 3u) * xv; gw >>= 2u;
+ }
+ }
+ {
+ uint pw = pu.w, gw = pg.w;
+ const threadgroup float* xe = xs + lane;
+ for (uint c = 0u; c < 16u; ++c) {
+ float xv = xe[(3u * 16u + c) * 32u];
+ qu += (float)(pw & 3u) * xv; pw >>= 2u;
+ qg += (float)(gw & 3u) * xv; gw >>= 2u;
+ }
+ }
+ qu += simd_shuffle_down(qu, 1u);
+ qg += simd_shuffle_down(qg, 1u);
+ float cu = 0.0f, cg = 0.0f;
+ if ((lane & 1u) == 0u) {
+ uint g = lane >> 1;
+ cu = qu * (float)sc[sbu + g] + (float)bi[sbu + g] * xg[g];
+ cg = qg * (float)sc[sbg + g] + (float)bi[sbg + g] * xg[g];
+ }
+ float accu = simd_sum(cu);
+ float accg = simd_sum(cg);
+ if (lane == 0u) {
+ float gv = metal::min(accg, 7.0f);
+ float uv = metal::clamp(accu, -7.0f, 7.0f);
+ h[e_slot * M + i] = (bfloat16_t)(gv / (1.0f + metal::exp(-gv)) * uv);
+ }
+ }
+ """,
+)
+
+# ---------------- kernel 2 v3 ----------------
+k2 = mx.fast.metal_kernel(
+ name="v3_down",
+ input_names=["h", "w", "sc", "bi", "inds", "scores"],
+ output_names=["out"],
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u; // expert slot
+ uint lane = tid % 32u;
+ constexpr uint K = 8u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = M / 16u; // 32
+ constexpr uint NG = M / 128u; // 4
+ constexpr uint ROWS = 8u;
+ uint j0 = tgid * ROWS;
+ uint e = (uint)inds[sg];
+
+ // transposed h staging: h[sg][lane*16+j] -> ht[sg][j*32+lane]
+ threadgroup float ht[K * 512];
+ threadgroup float hg[K * NG];
+ for (uint idx = tid; idx < K * 512u; idx += 256u) {
+ uint e2 = idx / 512u;
+ uint r = idx % 512u;
+ uint l = r / 16u;
+ uint j = r % 16u;
+ ht[e2 * 512u + j * 32u + l] = (float)h[idx];
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (tid < K * NG) {
+ uint e2 = tid / NG;
+ uint g = tid % NG;
+ float s = 0.0f;
+ for (uint l = 0u; l < 8u; ++l)
+ for (uint j = 0u; j < 16u; ++j)
+ s += ht[e2 * 512u + j * 32u + g * 8u + l];
+ hg[tid] = s;
+ }
+ threadgroup float pe[ROWS][K];
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+
+ for (uint r = 0u; r < ROWS; ++r) {
+ uint j = j0 + r;
+ ulong row = (ulong)e * D + j;
+ uint p = w[row * WPR + lane]; // lane = word index
+ ulong sb = row * NG;
+ float ql = 0.0f;
+ const threadgroup float* he = ht + sg * 512u + lane;
+ for (uint c = 0u; c < 16u; ++c) {
+ ql += (float)(p & 3u) * he[c * 32u];
+ p >>= 2u;
+ }
+ // octet reduction (group = lane/8)
+ ql += simd_shuffle_down(ql, 4u);
+ ql += simd_shuffle_down(ql, 2u);
+ ql += simd_shuffle_down(ql, 1u);
+ float contrib = 0.0f;
+ if ((lane & 7u) == 0u) {
+ uint g = lane >> 3;
+ contrib = ql * (float)sc[sb + g] + (float)bi[sb + g] * hg[sg * NG + g];
+ }
+ float acc = simd_sum(contrib);
+ if (lane == 0u) pe[r][sg] = acc;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (tid < ROWS) {
+ float o = 0.0f;
+ for (uint kk = 0u; kk < K; ++kk) o += pe[tid][kk] * scores[kk];
+ out[j0 + tid] = (bfloat16_t)o;
+ }
+ """,
+)
+
+
+def timeit(name, fn, n=300):
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:30s} build {(t1-t0)/n*1e6:7.1f} eval {(t2-t1)/n*1e6:8.1f} us")
+
+
+h = k1(inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16])[0]
+
+# correctness of k1
+e = int(inds[0])
+wdeq = mx.dequantize(ug.weight[e], ug.scales[e], ug.biases[e],
+ group_size=128, bits=2).astype(mx.float32)
+xdot = wdeq @ x.reshape(-1).astype(mx.float32)
+h_ref = nn.silu(mx.minimum(xdot[M:], 7.0)) * mx.clip(xdot[:M], -7.0, 7.0)
+mx.eval(h, h_ref)
+print("k1 maxdiff:", float(mx.abs(h[:M].astype(mx.float32) - h_ref).max()))
+
+out = k2(inputs=[h, dn.weight, dn.scales, dn.biases, inds, scores],
+ grid=(256, D // 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(D,)], output_dtypes=[mx.bfloat16])[0]
+
+# correctness of k2: reference = sum_e scores[e] * deq(w_e) @ h_e
+acc = mx.zeros((D,), dtype=mx.float32)
+for s_i in range(8):
+ e = int(inds[s_i])
+ wdeq = mx.dequantize(dn.weight[e], dn.scales[e], dn.biases[e],
+ group_size=128, bits=2).astype(mx.float32)
+ acc = acc + scores[s_i] * (wdeq @ h[s_i * M:(s_i + 1) * M].astype(mx.float32))
+mx.eval(out, acc)
+print("k2 maxdiff:", float(mx.abs(out.astype(mx.float32) - acc).max()))
+
+timeit("k1 v3", lambda: k1(
+ inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16])[0])
+timeit("k2 v3", lambda: k2(
+ inputs=[h, dn.weight, dn.scales, dn.biases, inds, scores],
+ grid=(256, D // 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(D,)], output_dtypes=[mx.bfloat16])[0])
diff --git a/.optimize/probe_v4.py b/.optimize/probe_v4.py
new file mode 100644
index 0000000..0d625e1
--- /dev/null
+++ b/.optimize/probe_v4.py
@@ -0,0 +1,214 @@
+import sys, time
+import mlx.core as mx
+import mlx.nn as nn
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+sw = model.model.layers[0].mlp.switch_mlp
+ug, dn = sw.up_gate_proj, sw.down_proj
+
+D, M, K = 2048, 512, 8
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16)
+h_in = mx.random.normal((8, 512), dtype=mx.bfloat16)
+inds = mx.array([3, 17, 42, 77, 100, 150, 201, 255], dtype=mx.int32)
+scores = mx.array([0.2, 0.15, 0.1, 0.15, 0.1, 0.1, 0.1, 0.1], dtype=mx.float32)
+mx.eval(x, h_in, inds, scores)
+xf = x.reshape(-1)
+
+# k1 v4: word-local affine (no tg mem, no barriers), uint4 x loads
+k1 = mx.fast.metal_kernel(
+ name="v4_upgate",
+ input_names=["x", "w", "sc", "bi", "inds"],
+ output_names=["h"],
+ header="""
+ inline float2 bf2(uint u) {
+ bfloat2 v = as_type(u);
+ return float2((float)v.x, (float)v.y);
+ }
+ """,
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u;
+ uint lane = tid % 32u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u; // 128
+ constexpr uint NG = D / 128u; // 16
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ const device uint4* x4 = (const device uint4*)x;
+
+ for (uint o = 0u; o < 2u; ++o) {
+ uint i = blk * 16u + o * 8u + sg;
+ ulong up_row = (ulong)e * 2u * M + i;
+ const device uint4* wu = (const device uint4*)(w + up_row * WPR);
+ const device uint4* wg = (const device uint4*)(w + (up_row + M) * WPR);
+ ulong sbu = up_row * NG;
+ ulong sbg = sbu + (ulong)M * NG;
+ float cu = 0.0f, cg = 0.0f;
+ // lane covers words lane*4..lane*4+3 -> group lane/2
+ {
+ uint g = lane >> 1;
+ float s_u = (float)sc[sbu + g];
+ float b_u = (float)bi[sbu + g];
+ float s_g = (float)sc[sbg + g];
+ float b_g = (float)bi[sbg + g];
+ uint4 pu = wu[lane];
+ uint4 pg = wg[lane];
+ uint pwu[4] = {pu.x, pu.y, pu.z, pu.w};
+ uint pwg[4] = {pg.x, pg.y, pg.z, pg.w};
+ for (uint t = 0u; t < 4u; ++t) {
+ uint widx = lane * 4u + t;
+ uint4 xa = x4[widx * 2u];
+ uint4 xb = x4[widx * 2u + 1u];
+ float xv[16];
+ xv[0] = bf2(xa.x).x; xv[1] = bf2(xa.x).y;
+ xv[2] = bf2(xa.y).x; xv[3] = bf2(xa.y).y;
+ xv[4] = bf2(xa.z).x; xv[5] = bf2(xa.z).y;
+ xv[6] = bf2(xa.w).x; xv[7] = bf2(xa.w).y;
+ xv[8] = bf2(xb.x).x; xv[9] = bf2(xb.x).y;
+ xv[10] = bf2(xb.y).x; xv[11] = bf2(xb.y).y;
+ xv[12] = bf2(xb.z).x; xv[13] = bf2(xb.z).y;
+ xv[14] = bf2(xb.w).x; xv[15] = bf2(xb.w).y;
+ float qu = 0.0f, qg = 0.0f, xs = 0.0f;
+ uint pu2 = pwu[t];
+ uint pg2 = pwg[t];
+ for (uint c = 0u; c < 16u; ++c) {
+ xs += xv[c];
+ qu += (float)(pu2 & 3u) * xv[c]; pu2 >>= 2u;
+ qg += (float)(pg2 & 3u) * xv[c]; pg2 >>= 2u;
+ }
+ cu += s_u * qu + b_u * xs;
+ cg += s_g * qg + b_g * xs;
+ }
+ }
+ float accu = simd_sum(cu);
+ float accg = simd_sum(cg);
+ if (lane == 0u) {
+ float gv = metal::min(accg, 7.0f);
+ float uv = metal::clamp(accu, -7.0f, 7.0f);
+ h[e_slot * M + i] = (bfloat16_t)(gv / (1.0f + metal::exp(-gv)) * uv);
+ }
+ }
+ """,
+)
+
+
+def timeit(name, fn, n=300):
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:30s} build {(t1-t0)/n*1e6:7.1f} eval {(t2-t1)/n*1e6:8.1f} us")
+
+
+h = k1(inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16])[0]
+
+e = int(inds[0])
+wdeq = mx.dequantize(ug.weight[e], ug.scales[e], ug.biases[e],
+ group_size=128, bits=2).astype(mx.float32)
+xdot = wdeq @ x.reshape(-1).astype(mx.float32)
+h_ref = nn.silu(mx.minimum(xdot[M:], 7.0)) * mx.clip(xdot[:M], -7.0, 7.0)
+mx.eval(h, h_ref)
+print("k1v4 maxdiff:", float(mx.abs(h[:M].astype(mx.float32) - h_ref).max()))
+
+timeit("k1 v4", lambda: k1(
+ inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16])[0])
+
+# k2 v4: same trick, no tg staging
+k2 = mx.fast.metal_kernel(
+ name="v4_down",
+ input_names=["h", "w", "sc", "bi", "inds", "scores"],
+ output_names=["out"],
+ header="""
+ inline float2 bf2(uint u) {
+ bfloat2 v = as_type(u);
+ return float2((float)v.x, (float)v.y);
+ }
+ """,
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u; // expert slot
+ uint lane = tid % 32u;
+ constexpr uint K = 8u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = M / 16u; // 32
+ constexpr uint NG = M / 128u; // 4
+ constexpr uint ROWS = 8u;
+ uint j0 = tgid * ROWS;
+ uint e = (uint)inds[sg];
+ const device uint4* h4 = (const device uint4*)(h + sg * M);
+ threadgroup float pe[ROWS][K];
+ for (uint r = 0u; r < ROWS; ++r) {
+ uint j = j0 + r;
+ ulong row = (ulong)e * D + j;
+ uint p = w[row * WPR + lane];
+ ulong sb = row * NG;
+ // lane's word = elements lane*16..+15 -> h4[lane*2, lane*2+1]
+ uint4 ha = h4[lane * 2u];
+ uint4 hb = h4[lane * 2u + 1u];
+ float hv[16];
+ hv[0] = bf2(ha.x).x; hv[1] = bf2(ha.x).y;
+ hv[2] = bf2(ha.y).x; hv[3] = bf2(ha.y).y;
+ hv[4] = bf2(ha.z).x; hv[5] = bf2(ha.z).y;
+ hv[6] = bf2(ha.w).x; hv[7] = bf2(ha.w).y;
+ hv[8] = bf2(hb.x).x; hv[9] = bf2(hb.x).y;
+ hv[10] = bf2(hb.y).x; hv[11] = bf2(hb.y).y;
+ hv[12] = bf2(hb.z).x; hv[13] = bf2(hb.z).y;
+ hv[14] = bf2(hb.w).x; hv[15] = bf2(hb.w).y;
+ float ql = 0.0f, hs = 0.0f;
+ for (uint c = 0u; c < 16u; ++c) {
+ hs += hv[c];
+ ql += (float)(p & 3u) * hv[c]; p >>= 2u;
+ }
+ // per-lane affine: lane's word sits wholly in group lane/8
+ uint g = lane >> 3;
+ float contrib = ql * (float)sc[sb + g] + hs * (float)bi[sb + g];
+ float acc = simd_sum(contrib);
+ if (lane == 0u) pe[r][sg] = acc;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ if (tid < ROWS) {
+ float o = 0.0f;
+ for (uint kk = 0u; kk < K; ++kk) o += pe[tid][kk] * scores[kk];
+ out[j0 + tid] = (bfloat16_t)o;
+ }
+ """,
+)
+
+h2 = h_in
+out = k2(inputs=[h2, dn.weight, dn.scales, dn.biases, inds, scores],
+ grid=(256, D // 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(D,)], output_dtypes=[mx.bfloat16])[0]
+acc = mx.zeros((D,), dtype=mx.float32)
+for s_i in range(8):
+ e = int(inds[s_i])
+ wdeq = mx.dequantize(dn.weight[e], dn.scales[e], dn.biases[e],
+ group_size=128, bits=2).astype(mx.float32)
+ acc = acc + scores[s_i] * (wdeq @ h2[s_i].astype(mx.float32))
+mx.eval(out, acc)
+print("k2v4 maxdiff:", float(mx.abs(out.astype(mx.float32) - acc).max()))
+
+timeit("k2 v4", lambda: k2(
+ inputs=[h2, dn.weight, dn.scales, dn.biases, inds, scores],
+ grid=(256, D // 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(D,)], output_dtypes=[mx.bfloat16])[0])
diff --git a/.optimize/probe_v5.py b/.optimize/probe_v5.py
new file mode 100644
index 0000000..185e13a
--- /dev/null
+++ b/.optimize/probe_v5.py
@@ -0,0 +1,160 @@
+import sys, time
+import mlx.core as mx
+import mlx.nn as nn
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+sw = model.model.layers[0].mlp.switch_mlp
+ug = sw.up_gate_proj
+
+D, M, K = 2048, 512, 8
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16)
+inds = mx.array([3, 17, 42, 77, 100, 150, 201, 255], dtype=mx.int32)
+mx.eval(x, inds)
+xf = x.reshape(-1)
+
+HDR = """
+ inline float2 bf2(uint u) {
+ bfloat2 v = as_type(u);
+ return float2((float)v.x, (float)v.y);
+ }
+ // 16 code extracts + fma into acc, independent shifts
+ #define DQ16(P, ACC, XV) { \
+ ACC += (float)((P) & 3u) * XV[0]; \
+ ACC += (float)(((P) >> 2u) & 3u) * XV[1]; \
+ ACC += (float)(((P) >> 4u) & 3u) * XV[2]; \
+ ACC += (float)(((P) >> 6u) & 3u) * XV[3]; \
+ ACC += (float)(((P) >> 8u) & 3u) * XV[4]; \
+ ACC += (float)(((P) >> 10u) & 3u) * XV[5]; \
+ ACC += (float)(((P) >> 12u) & 3u) * XV[6]; \
+ ACC += (float)(((P) >> 14u) & 3u) * XV[7]; \
+ ACC += (float)(((P) >> 16u) & 3u) * XV[8]; \
+ ACC += (float)(((P) >> 18u) & 3u) * XV[9]; \
+ ACC += (float)(((P) >> 20u) & 3u) * XV[10]; \
+ ACC += (float)(((P) >> 22u) & 3u) * XV[11]; \
+ ACC += (float)(((P) >> 24u) & 3u) * XV[12]; \
+ ACC += (float)(((P) >> 26u) & 3u) * XV[13]; \
+ ACC += (float)(((P) >> 28u) & 3u) * XV[14]; \
+ ACC += (float)(((P) >> 30u) & 3u) * XV[15]; }
+ #define XLOAD(WIDX, XV) { \
+ uint4 xa = x4[(WIDX) * 2u]; \
+ uint4 xb = x4[(WIDX) * 2u + 1u]; \
+ float2 f0 = bf2(xa.x); float2 f1 = bf2(xa.y); \
+ float2 f2 = bf2(xa.z); float2 f3 = bf2(xa.w); \
+ float2 f4 = bf2(xb.x); float2 f5 = bf2(xb.y); \
+ float2 f6 = bf2(xb.z); float2 f7 = bf2(xb.w); \
+ XV[0]=f0.x; XV[1]=f0.y; XV[2]=f1.x; XV[3]=f1.y; \
+ XV[4]=f2.x; XV[5]=f2.y; XV[6]=f3.x; XV[7]=f3.y; \
+ XV[8]=f4.x; XV[9]=f4.y; XV[10]=f5.x; XV[11]=f5.y; \
+ XV[12]=f6.x; XV[13]=f6.y; XV[14]=f7.x; XV[15]=f7.y; }
+"""
+
+# v5: each sg computes TWO outputs (up_i,gate_i and up_i8,gate_i8); x slice
+# loaded once per word position, used by all 4 rows.
+k1 = mx.fast.metal_kernel(
+ name="v5_upgate",
+ input_names=["x", "w", "sc", "bi", "inds"],
+ output_names=["h"],
+ header=HDR,
+ source="""
+ uint tid = thread_position_in_threadgroup.x;
+ uint tgid = threadgroup_position_in_grid.y;
+ uint sg = tid / 32u;
+ uint lane = tid % 32u;
+ constexpr uint D = 2048u;
+ constexpr uint M = 512u;
+ constexpr uint WPR = D / 16u; // 128
+ constexpr uint NG = D / 128u; // 16
+ uint e_slot = tgid / (M / 16u);
+ uint blk = tgid % (M / 16u);
+ uint e = (uint)inds[e_slot];
+ const device uint4* x4 = (const device uint4*)x;
+ uint i0 = blk * 16u + sg;
+ uint i1 = i0 + 8u;
+ ulong r_up0 = (ulong)e * 2u * M + i0;
+ ulong r_up1 = (ulong)e * 2u * M + i1;
+ const device uint4* w_up0 = (const device uint4*)(w + r_up0 * WPR);
+ const device uint4* w_up1 = (const device uint4*)(w + r_up1 * WPR);
+ const device uint4* w_g0 = w_up0 + (ulong)M * (WPR / 4u);
+ const device uint4* w_g1 = w_up1 + (ulong)M * (WPR / 4u);
+ ulong sb_up0 = r_up0 * NG, sb_up1 = r_up1 * NG;
+ ulong sb_g0 = sb_up0 + (ulong)M * NG, sb_g1 = sb_up1 + (ulong)M * NG;
+
+ uint g = lane >> 1;
+ float su0 = (float)sc[sb_up0 + g], bu0 = (float)bi[sb_up0 + g];
+ float su1 = (float)sc[sb_up1 + g], bu1 = (float)bi[sb_up1 + g];
+ float sg0 = (float)sc[sb_g0 + g], bg0 = (float)bi[sb_g0 + g];
+ float sg1 = (float)sc[sb_g1 + g], bg1 = (float)bi[sb_g1 + g];
+
+ float cu0 = 0.f, cg0 = 0.f, cu1 = 0.f, cg1 = 0.f;
+ uint4 pu0 = w_up0[lane], pu1 = w_up1[lane];
+ uint4 pg0 = w_g0[lane], pg1 = w_g1[lane];
+ uint wu0[4] = {pu0.x, pu0.y, pu0.z, pu0.w};
+ uint wu1[4] = {pu1.x, pu1.y, pu1.z, pu1.w};
+ uint wg0[4] = {pg0.x, pg0.y, pg0.z, pg0.w};
+ uint wg1[4] = {pg1.x, pg1.y, pg1.z, pg1.w};
+ for (uint t = 0u; t < 4u; ++t) {
+ uint widx = lane * 4u + t;
+ float xv[16];
+ XLOAD(widx, xv)
+ float xs = xv[0]+xv[1]+xv[2]+xv[3]+xv[4]+xv[5]+xv[6]+xv[7]
+ + xv[8]+xv[9]+xv[10]+xv[11]+xv[12]+xv[13]+xv[14]+xv[15];
+ float qa = 0.f, qb = 0.f, qc = 0.f, qd = 0.f;
+ DQ16(wu0[t], qa, xv)
+ DQ16(wg0[t], qb, xv)
+ DQ16(wu1[t], qc, xv)
+ DQ16(wg1[t], qd, xv)
+ cu0 += su0 * qa + bu0 * xs;
+ cg0 += sg0 * qb + bg0 * xs;
+ cu1 += su1 * qc + bu1 * xs;
+ cg1 += sg1 * qd + bg1 * xs;
+ }
+ float au0 = simd_sum(cu0), ag0 = simd_sum(cg0);
+ float au1 = simd_sum(cu1), ag1 = simd_sum(cg1);
+ if (lane == 0u) {
+ float g0v = metal::min(ag0, 7.0f);
+ float u0v = metal::clamp(au0, -7.0f, 7.0f);
+ float g1v = metal::min(ag1, 7.0f);
+ float u1v = metal::clamp(au1, -7.0f, 7.0f);
+ h[e_slot * M + i0] = (bfloat16_t)(g0v / (1.f + metal::exp(-g0v)) * u0v);
+ h[e_slot * M + i1] = (bfloat16_t)(g1v / (1.f + metal::exp(-g1v)) * u1v);
+ }
+ """,
+)
+
+
+def timeit(name, fn, n=300):
+ for _ in range(5):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ t1 = time.perf_counter()
+ mx.eval(rs)
+ t2 = time.perf_counter()
+ print(f"{name:30s} build {(t1-t0)/n*1e6:7.1f} eval {(t2-t1)/n*1e6:8.1f} us")
+
+
+h = k1(inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16])[0]
+
+e = int(inds[0])
+wdeq = mx.dequantize(ug.weight[e], ug.scales[e], ug.biases[e],
+ group_size=128, bits=2).astype(mx.float32)
+xdot = wdeq @ x.reshape(-1).astype(mx.float32)
+h_ref = nn.silu(mx.minimum(xdot[M:], 7.0)) * mx.clip(xdot[:M], -7.0, 7.0)
+mx.eval(h, h_ref)
+print("k1v5 maxdiff:", float(mx.abs(h[:M].astype(mx.float32) - h_ref).max()))
+
+timeit("k1 v5", lambda: k1(
+ inputs=[xf, ug.weight, ug.scales, ug.biases, inds],
+ grid=(256, (M // 16) * 8, 1), threadgroup=(256, 1, 1),
+ output_shapes=[(K * M,)], output_dtypes=[mx.bfloat16])[0])
diff --git a/.optimize/probes.sh b/.optimize/probes.sh
index 97a4c4f..f9435e1 100644
--- a/.optimize/probes.sh
+++ b/.optimize/probes.sh
@@ -1,7 +1,4 @@
-build | 3 | cd dream-chat && npm run build
-agent-tests | 3 | cd dream-chat && npm run test:agent
-ui-tests | 3 | cd dream-chat && npm run test:ui
-syntax | 3 | node --check dream-chat/electron/main.cjs && node --check dream-chat/electron/preload.cjs && node --check dream-chat/electron/provider_adapters.cjs
-maple-health | 1 | zsh .optimize/maple_health_probe.sh
-maple-stream-dispatch | 3 | node .optimize/maple_stream_dispatch_probe.cjs
-artifact-autopilot | 3 | cd dream-chat && npm run test:e2e:artifact-autopilot
+# label | repetitions | command
+# Maple speed probes run against a dedicated server on 127.0.0.1:8081.
+maple-decode-tps | 3 | python3 dream-chat/scripts/maple_bench.py --url http://127.0.0.1:8081 --mode decode
+maple-prefill-ttft | 3 | python3 dream-chat/scripts/maple_bench.py --url http://127.0.0.1:8081 --mode prefill
diff --git a/.optimize/prof_ab.py b/.optimize/prof_ab.py
new file mode 100644
index 0000000..1106cbf
--- /dev/null
+++ b/.optimize/prof_ab.py
@@ -0,0 +1,47 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+from mlx_lm.models.cache import make_prompt_cache
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+prompt = mx.array(tokenizer.encode("List the first six prime numbers."))
+inner = model.model
+
+
+def measure(tag, fused):
+ cache = make_prompt_cache(model)
+ logits = model(prompt[None, :-1], cache=cache)
+ mx.eval(logits)
+ y = prompt[-1:]
+ for _ in range(8):
+ logits = model(y[None], cache=cache)
+ mx.eval(logits)
+ y = mx.argmax(logits[:, -1, :], axis=-1)
+ for l in inner.layers:
+ l.mlp.switch_mlp._fused_decode = fused
+ build_t = eval_t = 0.0
+ N = 64
+ for _ in range(N):
+ t0 = time.perf_counter()
+ logits = model(y[None], cache=cache)
+ ny = mx.argmax(logits[:, -1, :], axis=-1)
+ t1 = time.perf_counter()
+ mx.eval(ny)
+ t2 = time.perf_counter()
+ build_t += t1 - t0
+ eval_t += t2 - t1
+ y = ny
+ print(f"{tag}: build {build_t/N*1e3:.2f}ms eval {eval_t/N*1e3:.2f}ms")
+
+
+measure("eager (fused off)", False)
+measure("fused moe", True)
+measure("eager again", False)
+measure("fused again", True)
diff --git a/.optimize/prof_gpu.py b/.optimize/prof_gpu.py
new file mode 100644
index 0000000..9d8c548
--- /dev/null
+++ b/.optimize/prof_gpu.py
@@ -0,0 +1,42 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+from mlx_lm.models.cache import make_prompt_cache
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+prompt = mx.array(tokenizer.encode("List the first six prime numbers."))
+cache = make_prompt_cache(model)
+
+# prefill
+logits = model(prompt[None, :-1], cache=cache)
+mx.eval(logits)
+
+y = prompt[-1:]
+# warm
+for _ in range(8):
+ logits = model(y[None], cache=cache)
+ mx.eval(logits)
+ y = mx.argmax(logits[:, -1, :], axis=-1)
+
+# measure: build time vs eval time per step
+build_t = 0.0
+eval_t = 0.0
+N = 64
+for _ in range(N):
+ t0 = time.perf_counter()
+ logits = model(y[None], cache=cache)
+ ny = mx.argmax(logits[:, -1, :], axis=-1)
+ t1 = time.perf_counter()
+ mx.eval(ny)
+ t2 = time.perf_counter()
+ build_t += t1 - t0
+ eval_t += t2 - t1
+ y = ny
+print(f"build: {build_t/N*1e3:.2f} ms/step eval: {eval_t/N*1e3:.2f} ms/step")
diff --git a/.optimize/prof_step.py b/.optimize/prof_step.py
new file mode 100644
index 0000000..b8b845e
--- /dev/null
+++ b/.optimize/prof_step.py
@@ -0,0 +1,34 @@
+import cProfile, pstats, io, sys, time
+import mlx.core as mx
+from mlx_lm.utils import load
+from mlx_lm.generate import generate_step
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True, "trust_remote_code": True},
+)
+prompt = mx.array(tokenizer.encode("List the first six prime numbers."))
+cache = None
+
+def run(n=64):
+ from mlx_lm.models.cache import make_prompt_cache
+ c = make_prompt_cache(model)
+ g = generate_step(prompt, model, max_tokens=n, prompt_cache=c)
+ out = [t for t, _ in g]
+ return out
+
+run(8) # warmup
+pr = cProfile.Profile()
+pr.enable()
+t0 = time.perf_counter()
+run(96)
+dt = time.perf_counter() - t0
+pr.disable()
+print(f"96 tokens in {dt:.2f}s -> {96/dt:.1f} tok/s")
+s = io.StringIO()
+pstats.Stats(pr, stream=s).sort_stats("cumulative").print_stats(40)
+print(s.getvalue()[:8000])
+s2 = io.StringIO()
+pstats.Stats(pr, stream=s2).sort_stats("tottime").print_stats(40)
+print(s2.getvalue()[:8000])
diff --git a/.optimize/runs/moe-after-decode.json b/.optimize/runs/moe-after-decode.json
new file mode 100644
index 0000000..83bdf59
--- /dev/null
+++ b/.optimize/runs/moe-after-decode.json
@@ -0,0 +1,15 @@
+{
+ "mode": "decode",
+ "ttft_ms": 156.6175005027617,
+ "decode_tps": 74.31758521879692,
+ "completion_tokens": 128,
+ "prompt_tokens": 25,
+ "cached_tokens": 24,
+ "context_tokens": 0,
+ "last_text": "\n\nHere are the first six prime numbers, each with a brief reason why it is prime:\n\n1. **2** \u2013 The only even prime, and it has no positive divisors other than 1 and itself. \n2. **3** \u2013 It is not divisible by 2, so the smallest possible divisor greater than 1 is 3, which is itself.",
+ "reps": 6,
+ "warmup": 2,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T14:06:40"
+}
diff --git a/.optimize/runs/moe-after-prefill.json b/.optimize/runs/moe-after-prefill.json
new file mode 100644
index 0000000..be7ae5a
--- /dev/null
+++ b/.optimize/runs/moe-after-prefill.json
@@ -0,0 +1,11 @@
+{
+ "mode": "prefill",
+ "ttft_ms": 7821.998583000095,
+ "prefill_tps": 404.628045686257,
+ "prompt_tokens": 3165,
+ "cached_tokens": 0,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T14:07:09"
+}
diff --git a/.optimize/runs/moe-after-score.json b/.optimize/runs/moe-after-score.json
new file mode 100644
index 0000000..b5c110a
--- /dev/null
+++ b/.optimize/runs/moe-after-score.json
@@ -0,0 +1,12 @@
+{
+ "mode": "score",
+ "total_ms": 993.9770420023706,
+ "per_candidate_ms": 124.24713025029632,
+ "prompt_tokens": 40,
+ "cached_tokens": 38,
+ "candidates": 8,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T14:07:12"
+}
diff --git a/.optimize/runs/moe-before-decode.json b/.optimize/runs/moe-before-decode.json
new file mode 100644
index 0000000..ac4c232
--- /dev/null
+++ b/.optimize/runs/moe-before-decode.json
@@ -0,0 +1,15 @@
+{
+ "mode": "decode",
+ "ttft_ms": 155.18335399610805,
+ "decode_tps": 80.22101795487421,
+ "completion_tokens": 128,
+ "prompt_tokens": 25,
+ "cached_tokens": 24,
+ "context_tokens": 0,
+ "last_text": "\n\nHere are the first six prime numbers, each with a brief reason:\n\n1. **2** \u2013 The only even prime, and it has no positive divisors other than 1 and itself. \n2. **3** \u2013 It is not divisible by 2, so the smallest possible divisor greater than 1 is 3 itself. \n3. **5** \u2013",
+ "reps": 6,
+ "warmup": 2,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T13:29:03"
+}
diff --git a/.optimize/runs/moe-before-prefill.json b/.optimize/runs/moe-before-prefill.json
new file mode 100644
index 0000000..fd90548
--- /dev/null
+++ b/.optimize/runs/moe-before-prefill.json
@@ -0,0 +1,11 @@
+{
+ "mode": "prefill",
+ "ttft_ms": 8209.231958004239,
+ "prefill_tps": 385.54154836787546,
+ "prompt_tokens": 3165,
+ "cached_tokens": 0,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T13:29:38"
+}
diff --git a/.optimize/runs/moe-before-score.json b/.optimize/runs/moe-before-score.json
new file mode 100644
index 0000000..078c606
--- /dev/null
+++ b/.optimize/runs/moe-before-score.json
@@ -0,0 +1,12 @@
+{
+ "mode": "score",
+ "total_ms": 1024.9728750059148,
+ "per_candidate_ms": 128.12160937573935,
+ "prompt_tokens": 40,
+ "cached_tokens": 38,
+ "candidates": 8,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T13:29:44"
+}
diff --git a/.optimize/runs/perf2-after.json b/.optimize/runs/perf2-after.json
new file mode 100644
index 0000000..707bac2
--- /dev/null
+++ b/.optimize/runs/perf2-after.json
@@ -0,0 +1,27 @@
+{
+ "agent": "perf2",
+ "server": "mlx_lm.server :8081 maple-2bit-mlx --flash-head --temp 0.7 --top-p 0.95 --top-k 20 --prefill-step-size 1024",
+ "ts": "2026-09-22T02:07:00",
+ "note": "after checkpoint maple.py sync (runtime artifact update, not a repo diff) + neutral node trims (add_rms reshape drop, qk flat-qkv reuse, B=1 sampler fast path). Custom fused-expert kernels tried and REVERTED (slower than gather_qmm); compiled MoE block tried and REVERTED (neutral, plus router-kernel ctr mutation unsafe under compile).",
+ "decode": {
+ "mode": "decode",
+ "ttft_ms": 191.58,
+ "decode_tps": 88.49,
+ "completion_tokens": 128,
+ "prompt_tokens": 25,
+ "cached_tokens": 24,
+ "context_tokens": 0,
+ "reps": 4
+ },
+ "decode_postsync_reference": {
+ "mode": "decode",
+ "ttft_ms": 134.34,
+ "decode_tps": 90.78,
+ "completion_tokens": 128,
+ "prompt_tokens": 25,
+ "cached_tokens": 24,
+ "context_tokens": 0,
+ "reps": 4,
+ "note": "measured right after checkpoint sync, before any perf2 code edits"
+ }
+}
diff --git a/.optimize/runs/perf2-before.json b/.optimize/runs/perf2-before.json
new file mode 100644
index 0000000..f514ef8
--- /dev/null
+++ b/.optimize/runs/perf2-before.json
@@ -0,0 +1,43 @@
+{
+ "agent": "perf2",
+ "server": "mlx_lm.server :8081 maple-2bit-mlx --flash-head --temp 0.7 --top-p 0.95 --top-k 20 --prefill-step-size 1024",
+ "ts": "2026-09-22T01:14:30",
+ "note": "measured on current code before perf2 changes; :8080 app server had exited, 16GB box still swap-pressured",
+ "decode": {
+ "mode": "decode",
+ "ttft_ms": 144.84,
+ "decode_tps": 72.12,
+ "completion_tokens": 128,
+ "prompt_tokens": 25,
+ "cached_tokens": 24,
+ "context_tokens": 0,
+ "reps": 4
+ },
+ "decode_ctx3000": {
+ "mode": "decode",
+ "ttft_ms": 189.27,
+ "decode_tps": 58.44,
+ "completion_tokens": 96,
+ "prompt_tokens": 3665,
+ "cached_tokens": 3664,
+ "context_tokens": 3000,
+ "reps": 3
+ },
+ "score": {
+ "mode": "score",
+ "total_ms": 910.29,
+ "per_candidate_ms": 113.79,
+ "prompt_tokens": 40,
+ "cached_tokens": 38,
+ "candidates": 8,
+ "reps": 4
+ },
+ "prefill": {
+ "mode": "prefill",
+ "ttft_ms": 7401.11,
+ "prefill_tps": 427.64,
+ "prompt_tokens": 3165,
+ "cached_tokens": 0,
+ "reps": 3
+ }
+}
diff --git a/.optimize/runs/speed-after-decode.json b/.optimize/runs/speed-after-decode.json
new file mode 100644
index 0000000..ecd9e18
--- /dev/null
+++ b/.optimize/runs/speed-after-decode.json
@@ -0,0 +1,15 @@
+{
+ "mode": "decode",
+ "ttft_ms": 144.2245839999714,
+ "decode_tps": 63.008593365486384,
+ "completion_tokens": 128,
+ "prompt_tokens": 25,
+ "cached_tokens": 24,
+ "context_tokens": 0,
+ "last_text": "\n\nHere are the first six prime numbers, each with a brief explanation of why it is prime:\n\n1. **2** \u2013 The only even prime number, and it has no positive divisors other than 1 and itself. \n2. **3** \u2013 It is not divisible by 2, so it is prime. \n3. **5** \u2013",
+ "reps": 5,
+ "warmup": 1,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T00:26:43"
+}
diff --git a/.optimize/runs/speed-after-prefill.json b/.optimize/runs/speed-after-prefill.json
new file mode 100644
index 0000000..d3bb20b
--- /dev/null
+++ b/.optimize/runs/speed-after-prefill.json
@@ -0,0 +1,11 @@
+{
+ "mode": "prefill",
+ "ttft_ms": 7480.81937500001,
+ "prefill_tps": 423.08199695036694,
+ "prompt_tokens": 3165,
+ "cached_tokens": 0,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T00:27:06"
+}
diff --git a/.optimize/runs/speed-after-score.json b/.optimize/runs/speed-after-score.json
new file mode 100644
index 0000000..b550716
--- /dev/null
+++ b/.optimize/runs/speed-after-score.json
@@ -0,0 +1,12 @@
+{
+ "mode": "score",
+ "total_ms": 1009.1129170000386,
+ "per_candidate_ms": 126.13911462500482,
+ "prompt_tokens": 40,
+ "cached_tokens": 38,
+ "candidates": 8,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-22T00:27:10"
+}
diff --git a/.optimize/runs/speed-baseline.json b/.optimize/runs/speed-baseline.json
new file mode 100644
index 0000000..639a34c
--- /dev/null
+++ b/.optimize/runs/speed-baseline.json
@@ -0,0 +1,46 @@
+{
+ "label": "speed-baseline",
+ "rev": "8a5b0c8 (dirty tree: 149 pre-existing modifications)",
+ "server": "maple-2bit-mlx :8081 --flash-head --temp 0.7 --top-p 0.95 --top-k 20 prefill-step 1024 conc 1/1 (checkpoint maple.py)",
+ "note": "clean re-baseline; first pass was polluted by concurrent local mlx profiling processes (decode read ~30 t/s, prefill ~183 t/s)",
+ "probes": {
+ "decode": {
+ "mode": "decode",
+ "ttft_ms": 202.50670798122883,
+ "decode_tps": 66.9823186475278,
+ "completion_tokens": 128,
+ "prompt_tokens": 25,
+ "cached_tokens": 24,
+ "context_tokens": 0,
+ "last_text": "\n\nHere are the first six prime numbers, each with a brief explanation of why it is prime:\n\n1. **2** \u2013 The only even prime number, and it has no positive divisors other than 1 and itself. \n2. **3** \u2013 It is not divisible by 2, so it is prime. \n3. **5** \u2013",
+ "reps": 5,
+ "warmup": 1,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-21T14:54:27"
+ },
+ "prefill": {
+ "mode": "prefill",
+ "ttft_ms": 6924.981500022113,
+ "prefill_tps": 457.0409321656517,
+ "prompt_tokens": 3165,
+ "cached_tokens": 0,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-21T14:54:54"
+ },
+ "score": {
+ "mode": "score",
+ "total_ms": 921.3034159038216,
+ "per_candidate_ms": 115.1629269879777,
+ "prompt_tokens": 40,
+ "cached_tokens": 38,
+ "candidates": 8,
+ "reps": 3,
+ "url": "http://127.0.0.1:8081",
+ "model": null,
+ "ts": "2026-09-21T14:54:57"
+ }
+ }
+}
\ No newline at end of file
diff --git a/.optimize/test_moe_fused.py b/.optimize/test_moe_fused.py
new file mode 100644
index 0000000..156136a
--- /dev/null
+++ b/.optimize/test_moe_fused.py
@@ -0,0 +1,78 @@
+import sys, time
+import mlx.core as mx
+
+sys.path.insert(0, "/Users/ianzvirbulis/Code/Hemlock")
+from mlx_lm.utils import load
+from mlx_lm.models.cache import make_prompt_cache
+from mlx_lm.models.maple import aggregate_expert_outputs
+
+model, tokenizer = load(
+ "/Users/ianzvirbulis/Models/Hemlock/maple-2bit-mlx",
+ tokenizer_config={"trust_remote_code": True},
+ model_config={"use_flash_head": True},
+ trust_remote_code=True,
+)
+inner = model.model
+layers = inner.layers
+
+# Force probe on a real decode-shaped input
+x = mx.random.normal((1, 1, 2048), dtype=mx.bfloat16, key=mx.random.key(0))
+blk = layers[0].mlp
+inds, scores = blk.gate(x)
+mx.eval(inds, scores)
+
+ref = aggregate_expert_outputs(blk.switch_mlp(x, inds), scores)
+got = blk.switch_mlp.fused_decode(x, inds, scores)
+mx.eval(ref, got)
+diff = mx.abs(ref.astype(mx.float32) - got.astype(mx.float32)).max()
+print("fused_decode flag:", blk.switch_mlp._fused_decode, "maxdiff:", float(diff))
+
+# compare across all layers on same x
+bad = []
+for i, l in enumerate(layers):
+ sw = l.mlp.switch_mlp
+ inds, scores = l.mlp.gate(x)
+ ref = aggregate_expert_outputs(sw(x, inds), scores)
+ got = sw.fused_decode(x, inds, scores)
+ mx.eval(ref, got)
+ d = float(mx.abs(ref.astype(mx.float32) - got.astype(mx.float32)).max())
+ if sw._fused_decode is not True or d > 0.02:
+ bad.append((i, sw._fused_decode, d))
+ print(f"layer {i}: fused={sw._fused_decode} maxdiff={d:.5f}")
+print("bad:", bad)
+
+# throughput: fused vs reference per layer
+sw = layers[0].mlp.switch_mlp
+inds, scores = layers[0].mlp.gate(x)
+mx.eval(inds, scores)
+
+
+def timeit(name, fn, n=300):
+ for _ in range(10):
+ r = fn()
+ mx.eval(r)
+ t0 = time.perf_counter()
+ rs = [fn() for _ in range(n)]
+ mx.eval(rs)
+ print(f"{name:44s} {(time.perf_counter()-t0)/n*1e6:9.1f} us")
+
+
+timeit("reference switch+agg", lambda: aggregate_expert_outputs(sw(x, inds), scores))
+timeit("fused_decode", lambda: sw.fused_decode(x, inds, scores))
+
+# greedy-equivalence smoke: generate 32 tokens each way
+from mlx_lm.generate import generate_step
+
+prompt = mx.array(tokenizer.encode("List the first six prime numbers."))
+c1 = make_prompt_cache(model)
+g = generate_step(prompt, model, max_tokens=32, prompt_cache=c1)
+seq1 = [t for t, _ in g]
+
+# disable fused and regenerate
+for l in layers:
+ l.mlp.switch_mlp._fused_decode = False
+c2 = make_prompt_cache(model)
+g = generate_step(prompt, model, max_tokens=32, prompt_cache=c2)
+seq2 = [t for t, _ in g]
+print("greedy identical:", seq1 == seq2)
+print(tokenizer.decode(seq1)[:200])
diff --git a/AGENTS.md b/AGENTS.md
index 7fbcaa0..456b475 100644
--- a/AGENTS.md
+++ b/AGENTS.md
@@ -14,20 +14,110 @@ paths, budgets, receipts. The model supplies intent and action selection.
→ executeAction → recurse. Autonomy gating, pause/resume, repair.
- `agent_contracts.cjs` — action schema, parse/repair, `allowedNextCommands`,
`buildScoreCandidates` (RLCD-style parallel constrained scoring).
- - `agent_queue.cjs` — serialized intent queue; drains only on terminal tasks.
- - `physics_sandbox.cjs` — deterministic experiment sims (Maple's world actions).
- - `dream_train.py` — LoRA Dream training (isolated adapters, hashed base weights).
+ - `agent_queue.cjs` — serialized intent queue; drains on terminal tasks
+ (blocked/cancelled included — parked states still hold the slot); pending
+ intents persist into session state and restore as queued on boot.
+ - `thread_manager.cjs` — threads: lifecycle, checkpoints/restore, search,
+ `forkThread` (new thread with `forkedFrom` provenance — not deep history).
+ - `shell_exec.cjs` — bounded `shell.exec`: bare allowlisted executables only
+ (no paths/wrappers/shells), workspace-scoped args, 60s + 32KB caps, receipts.
+ - `health_monitor.cjs` — polls `/health` only while a task is active; an
+ outage feeds the crash budget + bounded respawn, then the task auto-resumes
+ once from its last durable action.
+ - `physics_sandbox.cjs` — deterministic experiment sims (Maple's world actions)
+ + `experimentCoverage`/`suggestExperiments` (coverage-gap-driven
+ `experiment.suggest`, hypothesis open→addressed tracking).
+ - `sips_runtime.py` / `memory_fitness.cjs` — SIPS memory ledger: `memory.list`
+ (staleness/provenance/dupe clusters via rankScore) and plan-gated
+ `memory.consolidate` (append-only merge + audit trail).
+ - `dream_train.py` — LoRA Dream training (isolated adapters, hashed base
+ weights). Dataset sources include the agentic spine (durable action +
+ observation journals → executed-action SFT rows; only completed actions
+ with passed observations). `dream.dataset.preview` audits composition
+ without training.
- `coding_workspace.cjs` / `coding_autopilot.cjs` — scoped file edits + autopilot.
+ - Context budget: `contextMaxTokens` (default 24000,
+ `HEMLOCK_CONTEXT_MAX_TOKENS`) — over-budget requests compact the volatile
+ prompt region (receipted `context.compacted`, prompt note tells Maple);
+ still over → `CONTEXT_BUDGET_EXHAUSTED` block. `conversation.trim` trims
+ a thread's log with an archived prefix. Real token accounting lives in
+ `task.contextUsage` + `agent.self.context`.
+ - Settings: `settings.get/set` (persisted `/settings.json`, honest
+ appliesOn live|next-launch), `deps.check` (python3/mlx/mlx_lm/transformers/
+ maple-model/node — no torch anywhere in this stack).
- `dream-chat/src/` — React renderer (durable projection consumer; no own task state)
- - `main.jsx` — the whole app shell. `groveScene.js`/`groveBindings.js` — 3D world.
+ - `main.jsx` — app shell only: state, ctx bag, window loop, dock, palette, dialogs.
+ - `windows/` — extracted window components ({ctx}): ChatWindow, ArtifactStudio,
+ CommandCenter, ThreadsWindow (management: search/checkpoints/fork),
+ UtilityWindows (sips/memory/dream/activity/receipts/map/grove/settings).
+ `windows/shared.jsx` holds the shared helpers — window modules import from it,
+ never from main.jsx. `groveScene.js`/`groveBindings.js` — 3D world.
+ - `ctxSlices.js` — per-window ctx manifests + memoized slices (windows don't
+ re-render on unrelated host events). `eventBuffer.js` — ~16ms event
+ coalescing before ctx updates. `workToasts.js` — quiet terminal-event
+ toasts; `agentTimeline.js` — pure transcript-receipt row placement.
- `mlx_lm/` — forked mlx-lm; `server.py` adds `POST /v1/score` (teacher-forced
- candidate scoring on a shared prefilled cache).
+ candidate scoring on a shared prefilled cache; `commit:true` teacher-forces
+ the winner + `<|im_end|>` into the KV cache — the orchestrator replays prior
+ steps verbatim so each step is a strict-prefix extension, ~20 new tokens
+ instead of a full re-prefill. `HEMLOCK_NO_PREFIX_CACHE=1` disables).
+ `POST /v1/decide` is the kev-style decision endpoint: `{state, questions}`
+ with noul/choice/score types — each question is scored on a fork of the
+ prefilled state cache (questions can't see each other), options are
+ teacher-forced labels, softmax gives calibrated probabilities + confidence;
+ `commitQuestion` commits the winner's `continuation`. The orchestrator
+ prefers decide over score for action selection (confidence < 0.35 falls
+ back to generation); `HEMLOCK_NO_DECIDE=1` disables.
+ `--prompt-cache-file` persists the top-K prompt-cache entries across
+ restarts: `{path}.manifest.json` (atomic commit) + `{path}..safetensors`
+ per entry (`HEMLOCK_PROMPT_CACHE_SLOTS`, default 4, <=1 = legacy single
+ file; manifest-corrupt falls back to the legacy file). Save rides a
+ sentinel through the request queue so it runs on the generation thread —
+ its Metal stream is thread-local; warm load mx.evals lazy safetensors on
+ the loading thread. `assistant_prefix` pre-bakes `{"command":` into
+ structured-action calls (400 → host retries without).
+ `HEMLOCK_NO_CACHE_FILE=1` disables persistence.
+ Speed: `models/maple.py` builds one shared (pos,eps) array per decode step —
+ never `float(cache.offset)` an mx.array per layer (GPU sync per layer per
+ token, was ~2× decode cost). `--kv-bits` + `--quantized-kv-start` quantize
+ ONLY true `KVCache` entries (the 6 full-attention layers) past the start
+ token — `RotatingKVCache.to_quantized` raises and used to kill the Metal
+ command buffer; sliding layers must stay exact. Host envs:
+ `HEMLOCK_KV_BITS` (0/8/4, default 0) / `HEMLOCK_KV_QUANT_START` (5000);
+ Settings exposes kvBits (next-launch). Bench: `scripts/maple_bench.py`.
+ Measured: decode ~3× (30.6→~90 tok/s), prefill ~2.3×, score ~2.5× (see
+ `.optimize/runs/` + ledger). Next decode frontier: the MoE gather_qmm
+ pair is ~7.25ms of the 11.6ms eval — needs simdgroup-matrix expert
+ kernels. IMPORTANT: the served model loads `maple.py` from the CHECKPOINT
+ dir (`~/Models/Hemlock/maple-2bit-mlx/maple.py`, trust_remote_code), not
+ the repo file — keep them in sync or live serving silently loses the
+ optimizations. Regenerated checkpoints must ship the repo maple.py.
+ Context ceiling: `resolveContextMaxTokens` — stored pin > env > kv-scaled
+ default (24000 exact / 61440 kv-bits) > clamp to modelMax−2048; unpinned
+ tasks re-resolve per request (settings knob is live).
- `state.yaml` — append-only ledger of verified changes; read the tail first.
## Core invariants — do not break
- Prose is never executed. Actions are JSON envelopes validated against the
- registry; `capability: "train"` commands always require explicit approval.
+ registry — `{"actions":[...]}` batch envelopes run sequentially, each
+ validated + receipted; `capability: "train"` always requires explicit
+ approval. `artifact.author`/`artifact.update`/`world.place`/`plan.revise`
+ are plan-gated.
+- Agent introspection: `agent.self` (task/budget/queue/approvals/receipts) and
+ `world.state` (grove/experiment projection) are auto + read; `world.place`
+ persists agent-authored grove markers.
+- Idle-dream scheduler is proposal-only: a `dream.cycle` candidate lands in the
+ ambient inbox when ≥8 new findings exist post-queue-drain; acceptance still
+ goes through explicit train approval. `HEMLOCK_NO_IDLE_DREAM=1` disables.
+- Lane fallback: when structured-action inference dies of transport-death
+ (not parse/envelope failures), the step retries once via
+ `task.fallbackLane`/`HEMLOCK_FALLBACK_LANE` (codex|claude|none) — every
+ fallback is receipted with `provider`/`fallbackFrom`; receipts never claim
+ Maple did the work. `maple.warm` + ready/switch hooks prefill the shared
+ action-prefix so the first step is cache-hot (60s cooldown; skipped while
+ any stream/task is active). `memory.note` (plan-gated) lets the agent
+ write candidate lessons.
- Workspace writes go through `assertScopedPath(thread.workspaceRoot)` — an
improvement lane for Hemlock itself = a thread whose workspaceRoot is this repo.
- Training verifies base-weight integrity (SHA-256) as provenance; mutation is
diff --git a/README.md b/README.md
index 3acd2a3..b5ec635 100644
--- a/README.md
+++ b/README.md
@@ -13,8 +13,9 @@ the agent does lands in a durable, inspectable projection — no hidden state.
## The agent loop
An intent becomes a durable plan; Maple then proposes one structured JSON
-action at a time. The host validates it against a registry of ~90 allowlisted
-commands, executes one bounded operation, records an observation with evidence
+action at a time — or a `{"actions":[...]}` batch envelope executed step by
+step. The host validates it against a registry of 116 allowlisted commands,
+executes one bounded operation, records an observation with evidence
references, and asks for the next step. Prose is never executed.
- **Graduated autonomy** — `supervised` (approve every plan) → `guided` (plans
@@ -25,15 +26,29 @@ references, and asks for the next step. Prose is never executed.
state, `steer:` injects guidance into the next decision, cancel is terminal.
- **Scored action selection** — `/v1/score` teacher-forces candidate action
envelopes against a shared prefilled KV cache and picks the argmax, instead
- of generating JSON autoregressively. Complete envelopes execute with zero
- generated tokens (structurally incapable of invalid JSON); prefix winners
- fall through to generation only for free fields.
+ of generating JSON autoregressively; `/v1/decide` scores labeled options on
+ forked caches for calibrated confidence, with confidence and top1−top2
+ margin gates falling back to generation on weak calls. Complete envelopes
+ execute with zero generated tokens (structurally incapable of invalid
+ JSON); prefix winners fall through to generation only for free fields.
- **Self-describing** — `agent.capabilities` returns the live command registry
- with input contracts; every command carries an `inputHint`; failed actions
- return actionable recovery hints, not dead ends.
+ with input contracts; `agent.self` reports task, budget, queue, pending
+ approvals, and recent failures with host-authored recovery hints; every
+ command carries an `inputHint`.
+- **Self-correcting** — a missing required input gets one bounded repair
+ inference; `plan.revise` lets Maple replan the unexecuted suffix
+ mid-task (validated, receipted); decide divergences are logged, not hidden.
- **Self-improving** — Maple can read this repo, `improve.propose` a bounded
change (receipt'd, never auto-applied), `code.apply` it under an approved
- plan, `verify` with real build/test profiles, and `remember` the lesson.
+ plan, `verify` with real build/test profiles, and `remember` the lesson via
+ `memory.note`. `shell.exec` runs bounded allowlisted executables against
+ the workspace.
+- **Resident and warm** — a persistent multi-entry prompt cache survives
+ restarts and lane switches, and the host warm-prefills the shared action
+ prefix on server-ready and thread-switch, so first steps are cache-hot.
+- **Multi-lane** — Maple is the default; Codex and Claude lanes exist, and a
+ configured fallback lane retries transport-death inference once — receipted
+ with `fallbackFrom` so the record never claims Maple did the work.
## The Grove
@@ -41,14 +56,21 @@ Maple exists as a hooded figure in a 3D world (`dream-chat/src/groveScene.js`).
`experiment.run` executes deterministic physics sims (pendulum, projectile,
orbit, spring, collision, terminal) in a bounded sandbox — each run emits a
receipt **and** a trail of real integration samples that the grove replays at
-Maple's lab bench while Maple physically works. Findings feed the Dream
+Maple's lab bench while Maple physically works. `experiment.suggest` picks
+the next run from coverage gaps and open hypotheses; `world.place` lets Maple
+leave persistent markers in the world it inhabits. Findings feed the Dream
dataset. World → experiment → receipt → finding → dataset → adapter → behavior.
## Dream training → graft → fuse
-Dream Lab prepares a token-budgeted dataset (splits oversized rows, records
-trimming metadata, deterministic holdout) and runs real MLX LoRA training with
-isolated adapters and base-weight provenance. A verified adapter can be:
+Dream Lab prepares a token-budgeted dataset — facts, conversations, coding
+examples, **and durable agentic action traces** (real system-prompt bytes,
+step context rebuilt at each action's own timestamp, only completed actions
+with passed observations) — with `dream.dataset.preview` auditing the mix
+before anything trains. Then real MLX LoRA training runs with isolated
+adapters and base-weight provenance. An idle-dream scheduler *proposes* a
+cycle when enough new findings have landed; training itself always stays
+explicitly approved. A verified adapter can be:
- **grafted** — served via `--adapter-path`, detachable (`dream.detach`),
- **fused** — `dream.fuse` merges it into a *new* checkpoint and serves that;
@@ -61,9 +83,14 @@ explicit, auditable act — never silent.
## Surfaces
Chat · Artifact Studio (versioned live previews, sandboxed) · Command Center
-(host loop, intent queue, metrics) · Understory Grove · Dream Lab · Memory
-Garden · SIPS · Activity · Receipts · Project Map · Settings. Ordinary
-conversation stays ordinary — creation verbs alone don't trigger plans.
+(host loop, intent queue, metrics, notification history) · **Threads**
+(search, checkpoints, fork, full lifecycle) · Understory Grove · Dream Lab ·
+Memory Garden (staleness/provenance-ranked recall, near-dupe consolidation) ·
+SIPS · Activity · Receipts (expandable, cross-window deep links) · Project
+Map · Settings (runtime tuning: KV-cache quantization, context ceiling,
+prefill step, prompt-cache slots, fallback lane, dependency probes).
+Ordinary conversation stays ordinary — creation verbs alone don't trigger
+plans.
## Run it
@@ -93,14 +120,22 @@ python -m mlx_lm chat --model ./maple-2bit-mlx --trust-remote-code --max-tokens
```sh
cd dream-chat
-npm run test:agent # host harness: orchestrator, queue, contracts, receipts
+npm run test:agent # host harness: orchestrator, queue, contracts, receipts,
+ # durable-io recovery, session restore, real-main soak
npm run test:ui # renderer projection tests
npm run verify:agent # all of the above + production build
node scripts/gui-electron-smoke.mjs # real-Electron end-to-end checks
+python -m pytest mlx_lm/tests/test_server.py # inference server incl. kernels
```
Maple is a 20B-A1B ternary MoE — 24 layers, 256 experts, top-8, 512-token
sliding window on 3 of 4 layers, 2-bit packed `{-α, 0, +α}` weights.
+Inside Hemlock the decode path additionally runs fused custom Metal kernels
+for the expert block (one dispatch for gate+up+SwiGLU across all 8 active
+experts, another for weighted down-projection), and `HEMLOCK_KV_BITS=8`
+opt-in quantizes the six full-attention KV caches past a configurable start
+token — lifting the practical context ceiling from ~24k to ~61k tokens while
+sliding-window layers stay exact.
| chip | head | decode tok/s | prefill tok/s | peak |
| --- | --- | --- | --- | --- |
diff --git a/benchmarks/prefix_session_probe.py b/benchmarks/prefix_session_probe.py
new file mode 100644
index 0000000..8b6a273
--- /dev/null
+++ b/benchmarks/prefix_session_probe.py
@@ -0,0 +1,133 @@
+"""Continuous-KV session probe for the Hemlock Maple server.
+
+Simulates the agent loop's per-step prompts two ways against a running server:
+
+ session - prior steps replayed verbatim as chat turns (user request +
+ assistant reply), so each rendered prompt is a strict token
+ prefix extension of the server's committed cache entry.
+ rebuilt - the classic shape: system + one user message whose JSON embeds
+ the growing history (prompt diverges mid-body from every stored
+ key, so reuse stops at the last shared prefix boundary).
+
+/v1/score is called with commit=true and the fixed reasoning suffix so the
+server stores prompt+suffix+winner+eos; the assistant turn replays as
+{reasoning_content, content}, which the chat template renders byte-identical.
+
+Usage: python benchmarks/prefix_session_probe.py [base_url]
+"""
+
+import json
+import sys
+import urllib.request
+
+BASE = sys.argv[1] if len(sys.argv) > 1 else "http://127.0.0.1:8080"
+RC = "Deciding next action."
+SYSTEM = (
+ "You are Maple-Preview operating inside Hemlock. Return exactly one "
+ "compact JSON action: "
+ '{"kind":"tool","commandId":"","input":{},"shortRationale":"..."}. '
+ "Choose the next best command from request.allowedNextCommands."
+)
+CANDIDATES = [
+ '{"kind":"tool","commandId":"repo-map","input":{},"shortRationale":"Map the repo."}',
+ '{"kind":"tool","commandId":"repo.inspect","input":{},"shortRationale":"Inspect the repo."}',
+ '{"kind":"tool","commandId":"file.read","input":',
+]
+# Pad the system prompt so the shared prefix is a realistic size (~1k tok).
+SYSTEM += "\n" + "\n".join(f"Rule {i}: keep actions scoped and receipted." for i in range(120))
+
+
+def post(path, body):
+ req = urllib.request.Request(
+ BASE + path,
+ data=json.dumps(body).encode(),
+ headers={"Content-Type": "application/json"},
+ )
+ return json.loads(urllib.request.urlopen(req, timeout=600).read())
+
+
+def step_user(step, history):
+ return json.dumps(
+ {
+ "task": {"id": "probe", "objective": "inspect the project"},
+ "context": {"note": f"observation for step {step}: files scanned ok"},
+ "nextStep": {"commandId": ["repo-map", "repo.inspect", "file.read"][step - 1]},
+ "allowedNextCommands": [{"commandId": c} for c in ["repo-map", "repo.inspect", "file.read"]],
+ "progress": {"planProgress": f"step {step} of 3"},
+ "completed": {"actions": history},
+ }
+ )
+
+
+def completions(messages):
+ body = {
+ "model": "default_model",
+ "messages": messages,
+ "max_tokens": 8,
+ "temperature": 0,
+ "stream": False,
+ }
+ out = post("/v1/chat/completions", body)
+ usage = out.get("usage") or {}
+ return (
+ usage.get("prompt_tokens"),
+ (usage.get("prompt_tokens_details") or {}).get("cached_tokens"),
+ )
+
+
+def score(messages):
+ body = {
+ "model": "default_model",
+ "messages": messages,
+ "prompt_suffix": RC + "\n\n\n",
+ "candidates": CANDIDATES,
+ "commit": True,
+ }
+ return post("/v1/score", body)
+
+
+def run_session():
+ print("\n=== session mode (verbatim turn replay) ===")
+ turns = []
+ history = []
+ for step in (1, 2, 3):
+ user = step_user(step, history[-1:])
+ messages = [{"role": "system", "content": SYSTEM}, *turns, {"role": "user", "content": user}]
+ r = score(messages)
+ winner = r["candidates"][r["committedIndex"] if r.get("committedIndex") is not None else 0]
+ print(
+ f"step {step}: prompt={r['promptTokens']:>5} cached={r['cachedTokens']:>5} "
+ f"committed={r.get('committedIndex')}"
+ )
+ turns.append({"role": "user", "content": user})
+ turns.append(
+ {
+ "role": "assistant",
+ "content": r.get("committedText") or winner["text"],
+ "reasoning_content": RC,
+ }
+ )
+ history.append({"commandId": winner["text"]})
+ # And the generative continuation for the same step.
+ gen_messages = messages + [
+ turns[-1],
+ {"role": "user", "content": '{"instruction":"Emit the complete JSON action."}'},
+ ]
+ pt, ct = completions(gen_messages)
+ print(f" infer after commit: prompt={pt} cached={ct}")
+
+
+def run_rebuilt():
+ print("\n=== rebuilt mode (single user message, growing history) ===")
+ history = []
+ for step in (1, 2, 3):
+ history.append({"commandId": f"cmd{step}", "status": "completed"})
+ user = step_user(step, history)
+ messages = [{"role": "system", "content": SYSTEM}, {"role": "user", "content": user}]
+ r = score(messages)
+ print(f"step {step}: prompt={r['promptTokens']:>5} cached={r['cachedTokens']:>5}")
+
+
+if __name__ == "__main__":
+ run_session()
+ run_rebuilt()
diff --git a/dream-chat/.impeccable/live/config.json b/dream-chat/.impeccable/live/config.json
new file mode 100644
index 0000000..dba4f77
--- /dev/null
+++ b/dream-chat/.impeccable/live/config.json
@@ -0,0 +1,6 @@
+{
+ "files": ["index.html"],
+ "insertBefore": "