Qwen3.8-27B on 2x P150a: measured serving recipe (121.4 aggregate tok/s at B=8, 20.5 single-stream) #19

Merged
multica-agent merged 94 commits from autoresearch/dspark-drafter-sep07 into master 2026-09-13 10:25:06 +02:00
Owner

Qwen3.8-27B on 2× P150a: a measured serving recipe

Supersedes this PR's original description entirely. That text was written 2026-09-10 around an MTP-primary plan and a 14.4 tok/s baseline. Both are obsolete: MTP was dropped by owner decision on 2026-09-12, and the baseline has moved. What follows is the current, measured state.

Headline, measured on hardware

config ms/step per-user tok/s aggregate tok/s
B=1, best single-stream arm 48.5–48.8 20.5–20.6 20.5–20.6
B=1, batch-safe arm 50.9 19.63 19.63
B=4, batch-safe arm 57.4 17.42 68.7
B=8, batch-safe arm 64.8 15.44 121.4

4k context, TP=2, two cabled P150a, vllm-tt:k2. Every row is reproduced by at least two independent harnesses; the B=8 point three times within 0.3%. Control spread on the bench is 0.09–0.48%.

Read the two B=1 rows carefully — they are different recipes, and the faster one does not scale. LM_HEAD_GATHER_MODE=none is worth ~2 ms at B=1 and costs +31.0 ms at B=8. Anyone adopting this should take the batch-safe arm unless they genuinely serve one stream at a time.

What actually produced the speedup

The single largest win was not a kernel. process_output_decode was reading a 32-row tile-padded logits buffer (15.9 MB) and untilizing it on the host — 10.4 ms/token of pure overhead at batch 1. Untilizing on device inside the trace (QWEN36_DECODE_LOGITS_RM=1) so the host reads a [1,1,B,vocab] row-major buffer took the step from 62.85 → 50.70 ms, output-identical. It also scales into batch gracefully (≈1 MB at B=8 instead of 15.9 MB).

Four optimisation directions closed BY MEASUREMENT

Recorded so nobody re-spends the time:

  1. Op-count optimisation is dead. Trace replay is worth 176 ms/step at B=8 (64.19 traced vs 240.20 untraced, 3.74×). Per-op dispatch is already amortised to ~zero in production, so removing programs buys nothing unless they consume device time.
  2. Grid tuning is dead. All five QWEN36_1D_GRID_* levers fall inside a 0.48% control band. Seven decode matmuls carry TP=4 core-count constants on a TP=2 deployment (confirmed from the device's own banner, leaving most of a 110-core grid idle) — real, and worth no measurable time.
  3. Bandwidth was never the constraint. The weight read is 10.76 GB/card/token at 454–469 GB/s ≈ 23.34 ms, ~100% of achievable, and flat in batch. gate_proj/up_proj were always bf4; the "13.61 GB all-bf8" premise in older docs is false.
  4. No large single host win. Real host CPU is ~15 ms and diffuse, nothing above 2.6 ms. The 76.5% py-spy frame is the blocking readback — i.e. the device step, not host cost.

Known limitations, stated plainly

  • No per-op device millisecond table. vllm-tt:k2 is not Tracy-enabled; the profiler refuses. Getting one needs a Tracy build of the production tree.
  • Batch width alone changes greedy output. Same prompt, temperature 0, one server: different text at B=4 vs B=8 in 3 of 4 prompts, divergence after an identical prefix. One gap of 1.75 logprob is ~14 bf16 quantisation steps and cannot be a tie. Slot position and arrival order are invariant. Consequence: validating a change at one width does not license shipping it at the other. Unresolved.
  • No host-greedy reference exists (~54 GB bf16 vs 5.5 GB usable on the target), so the equivalence references license "nothing changed since this point", not "the output is correct".
  • MTP was never measured. Its acceptance rate is unknown. It was dropped because speculation is B=1 by construction (_commit(mi) applies one scalar accepted index across all 48 GDN layers; the batch axis is already spent on K+1 candidates) and because the MTP tree and the performance tree are disjoint. It must not be recorded as rejected on evidence.

Also in this branch

Equivalence harness for batched decode (bench/batch-equiv-bench.py) whose load-bearing assertion is that the B per-stream hashes must be distinct — an assertion our previous harnesses would have failed, because decode-bench.py's counting prompt produced one identical text_sha256 across every config and context length. Plus the batch-scaling ladder, the serving A/B driver, the low-level audit records, and incident notes (a ~10 GB pinned-memory leak over six days of uptime that OOM-killed runs and looked like a network outage).

All evidence is in bench/runs/*.jsonl; every record carries its config and caveats, and anything not timed on hardware is performance=null, gate_complete=false.

🤖 Generated with Claude Code

## Qwen3.8-27B on 2× P150a: a measured serving recipe **Supersedes this PR's original description entirely.** That text was written 2026-09-10 around an MTP-primary plan and a 14.4 tok/s baseline. Both are obsolete: MTP was dropped by owner decision on 2026-09-12, and the baseline has moved. What follows is the current, measured state. ### Headline, measured on hardware | config | ms/step | per-user tok/s | aggregate tok/s | |---|---:|---:|---:| | B=1, best single-stream arm | 48.5–48.8 | **20.5–20.6** | 20.5–20.6 | | B=1, batch-safe arm | 50.9 | 19.63 | 19.63 | | B=4, batch-safe arm | 57.4 | 17.42 | 68.7 | | **B=8, batch-safe arm** | **64.8** | **15.44** | **121.4** | 4k context, TP=2, two cabled P150a, `vllm-tt:k2`. Every row is reproduced by at least two independent harnesses; the B=8 point three times within 0.3%. Control spread on the bench is 0.09–0.48%. **Read the two B=1 rows carefully — they are different recipes, and the faster one does not scale.** `LM_HEAD_GATHER_MODE=none` is worth ~2 ms at B=1 and costs **+31.0 ms at B=8**. Anyone adopting this should take the batch-safe arm unless they genuinely serve one stream at a time. ### What actually produced the speedup The single largest win was not a kernel. `process_output_decode` was reading a **32-row tile-padded logits buffer** (15.9 MB) and untilizing it **on the host** — 10.4 ms/token of pure overhead at batch 1. Untilizing on device inside the trace (`QWEN36_DECODE_LOGITS_RM=1`) so the host reads a `[1,1,B,vocab]` row-major buffer took the step from 62.85 → 50.70 ms, output-identical. It also scales into batch gracefully (≈1 MB at B=8 instead of 15.9 MB). ### Four optimisation directions closed BY MEASUREMENT Recorded so nobody re-spends the time: 1. **Op-count optimisation is dead.** Trace replay is worth **176 ms/step at B=8** (64.19 traced vs 240.20 untraced, 3.74×). Per-op dispatch is already amortised to ~zero in production, so removing programs buys nothing unless they consume *device* time. 2. **Grid tuning is dead.** All five `QWEN36_1D_GRID_*` levers fall inside a 0.48% control band. Seven decode matmuls carry **TP=4 core-count constants** on a TP=2 deployment (confirmed from the device's own banner, leaving most of a 110-core grid idle) — real, and worth no measurable time. 3. **Bandwidth was never the constraint.** The weight read is 10.76 GB/card/token at 454–469 GB/s ≈ 23.34 ms, ~100% of achievable, and flat in batch. `gate_proj`/`up_proj` were always bf4; the "13.61 GB all-bf8" premise in older docs is false. 4. **No large single host win.** Real host CPU is ~15 ms and diffuse, nothing above 2.6 ms. The 76.5% py-spy frame is the *blocking readback* — i.e. the device step, not host cost. ### Known limitations, stated plainly - **No per-op device millisecond table.** `vllm-tt:k2` is not Tracy-enabled; the profiler refuses. Getting one needs a Tracy build of the production tree. - **Batch width alone changes greedy output.** Same prompt, temperature 0, one server: different text at B=4 vs B=8 in 3 of 4 prompts, divergence after an identical prefix. One gap of 1.75 logprob is ~14 bf16 quantisation steps and cannot be a tie. Slot position and arrival order *are* invariant. Consequence: **validating a change at one width does not license shipping it at the other.** Unresolved. - **No host-greedy reference exists** (~54 GB bf16 vs 5.5 GB usable on the target), so the equivalence references license "nothing changed since this point", not "the output is correct". - **MTP was never measured.** Its acceptance rate is unknown. It was dropped because speculation is B=1 by construction (`_commit(mi)` applies one scalar accepted index across all 48 GDN layers; the batch axis is already spent on K+1 candidates) and because the MTP tree and the performance tree are disjoint. It must not be recorded as rejected on evidence. ### Also in this branch Equivalence harness for batched decode (`bench/batch-equiv-bench.py`) whose load-bearing assertion is that the B per-stream hashes must be **distinct** — an assertion our previous harnesses would have failed, because `decode-bench.py`'s counting prompt produced one identical `text_sha256` across every config and context length. Plus the batch-scaling ladder, the serving A/B driver, the low-level audit records, and incident notes (a ~10 GB pinned-memory leak over six days of uptime that OOM-killed runs and looked like a network outage). All evidence is in `bench/runs/*.jsonl`; every record carries its config and caveats, and anything not timed on hardware is `performance=null`, `gate_complete=false`. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
docs: reopen speculative decoding for native Qwen3.8 drafters
All checks were successful
tt-stack-ci / Report upstream drift (pull_request) Successful in 5s
tt-stack-ci / Build simulators and check the host module (pull_request) Successful in 19s
a8744312f5
Native DSpark and DFlash2 checkpoints invalidate the earlier assumption that no final-model drafter exists. Define the correctness, state-transaction, packed-verify, and on-silicon evidence gates needed to determine whether they can improve multi-stream decode on two P150a cards without treating DGX results as transferable.
Co-authored-by: Junie <junie@jetbrains.com>
Co-authored-by: Junie <junie@jetbrains.com>
Co-authored-by: Junie <junie@jetbrains.com>
Co-authored-by: Junie <junie@jetbrains.com>
Co-authored-by: Junie <junie@jetbrains.com>
Co-authored-by: Junie <junie@jetbrains.com>
Make DSpark component results reproducible on the P150a pair
All checks were successful
tt-stack-ci / Report upstream drift (pull_request) Successful in 4s
tt-stack-ci / Build simulators and check the host module (pull_request) Successful in 9s
368dba89c0
Projector and MLP probes now pass on hardware, but ad hoc launches did not preserve a reproducible runtime or guard shared devices. Capture source snapshots, immutable image IDs and failure evidence under a serialized remote harness without claiming full drafter parity or performance.

Co-authored-by: Junie <junie@jetbrains.com>
Grok added 12 commits 2026-09-10 14:37:57 +02:00
Three GDN commit mechanisms (snapshot, shadow_commit, recompute) behind one
begin_verify/commit/abort contract. DeltaNetReference provides a deterministic
path-dependent recurrent backend so the transaction semantics are provable on
the host before device probes attach the real GDN kernel.

13 new tests cover:
- Commit matches autoregressive oracle for every acceptance length 0..K
- Abort leaves live state, position and RNG untouched
- Consecutive reject-then-accept cycles
- All three mechanisms agree on committed state
- Out-of-range acceptance, bad bonus, double begin_verify, invalid mechanism

Full CPU suite: 33 passed (state 13, sampler 4, request 5, metrics 1, drafter 8).
The 2026-09-07 full-CPU record was blocked because test_speculative_state.py
imported the then-missing serving/speculative/state.py. That module is now
committed (7103db2), so the full CPU contract suite in vllm-tt:k2
(sha256:6105bb72) runs clean: 46 passed, 0 failed, 0 errors
(drafter 8, harness 13, metrics 1, request 5, sampler 6, state 13).

Supersede the 09-07 blocked record with the 09-08 green one. Also document a
nixpkgs gap discovered while validating: python3 (now 3.14) lacks the _expat
C extension, so xml.etree.ElementTree is unavailable and 8 of 13 harness unit
tests fail in a bare nix shell. That is an environment defect, not a code
defect; the vllm-tt:k2 image has a full CPython with expat and all 13 pass.
Resolves the D8.0 gate from GOAL-DSPARK: the SGLang DSpark proposal
(feature-tap, drafter, verify, state-pool) and the TT qwen36 GDN/attention
destinations are now mapped to specific pinned source revisions in
research/sglang-pinned and research/tt-pinned.

Items 6 (SGLang proposal license) and 7 (SpecForge training provenance)
are external-provenance gaps recorded as unknown — they gate neither the
speculative-decoding correctness work nor the GDN rollback investigation.
Item 8 (drafter tensor shape verification) is deferred to D8.2, which
needs real drafter fixtures.
Benchmark entrypoint: bash autoresearch.sh
Goal: Reduce the per-commit GDN state cost of DSpark speculative decoding on Qwen3.8-27B so the GOAL-DSPARK success criterion holds: accepted tokens per verification step must outrun the combined draft, verify, state-commit, and serving costs. The measured lever is the state-commit protocol in serving/speculative (the D8.1 cost model): how many GDN state copies (bytes) and recurrent steps a speculative commit performs across the snapshot / shadow_commit / recompute mechanisms at K=1,3,7. Primary = total state ops (copies + steps) for shadow_commit at K=7 (currently 9). A valid optimization reduces real ops in the transaction protocol (fewer copies, or reusing precomputed boundary states) WITHOUT breaking the invariant that all three mechanisms commit the identical GDN state and advance position by accepted+1.
Result: {"status":"keep","gdn_commit_ops_shadow_k7":8,"gdn_commit_steps_shadow_commit_k7":8,"gdn_commit_copies_shadow_commit_k7":0,"gdn_commit_mb_shadow_commit_k7":0,"gdn_commit_copies_snapshot_k7":1,"gdn_commit_copies_recompute_k7":1,"gdn_commit_steps_snapshot_k7":15,"gdn_commit_steps_recompute_k7":15,"gdn_commit_state_consistent":1,"gdn_layer_bytes":1073152}
Drives the production T=1 GDN decode primitive (recurrent_gated_delta_rule_decode_ttnn)
through our StateTransaction protocol at TP=2. For K in {1,3,7} and every
acceptance length a in [0,K], each of the three commit mechanisms (snapshot /
shadow_commit / recompute) must commit the GDN recurrent state matching an
autoregressive oracle replaying the identical token sequence one decode step at
a time; a rejected-verify (abort) plus a later accept is also covered.

Uses the 9B/27B GDN state shape [B, Nv, Dk, Dv] fp32 at B=1, Nv=8 (reduced from
32 to stay under the GDN decode kernel's compute-grid cap), Dk=Dv=128. The
conv state is tracked as a per-token q shift register (the primitive covers the
recurrent state, which is the D8.1 rollback core; the conv path is soft-checked).
No checkpoint weights needed (deterministic per-token inputs). Registered in
scripts/dspark-harness.py PROBES (name gdn-transaction ->
bench/probes/dspark-gdn-transaction-tp.py).
The conv shift register starts empty and fills over the first CONV_TAPS
decode steps, so for short acceptance lengths (a=0 -> 1 step) many conv
taps are still None in both target and committed. The previous min() over
_pcc(oc,cc) crashed on (None, None). Added _conv_pcc which compares only
taps non-None in both states and treats an all-None pair as a trivial
match (both histories are empty).
Drives the production T=1 GDN decode primitive (recurrent_gated_delta_rule_decode_ttnn)
through our StateTransaction protocol on a 2-device mesh. For K in {1,3,7} and every
acceptance length a in [0,K], all three commit mechanisms (snapshot / shadow_commit /
recompute) commit the GDN recurrent state bit-identical to the autoregressive oracle
that re-advances the pre-verify state one decode step at a time (recPCC=1.0000 for
all 39 cells). A rejected verify (abort) leaves the live state exactly unchanged
(torch.equal, after driving it to a non-trivial state), and a later commit after that
rejection still matches the oracle advanced from the unchanged live state.

The conv history is tracked as a per-token q shift register (a proxy); the primary
D8.1 validation is the recurrent-state equality. performance: null (correctness only).
Record: bench/runs/dspark-20260908T041702Z-6a5ffeb8.jsonl.
Closes D8.2 (GOAL-DSPARK line 62): the drafter checkpoint loads without
silent drops and produces the same block logits/tokens as a torch
reference for fixed captured target features, without needing a full
target checkpoint for the first test.

What was believed before / what the evidence changed: the D8.0 manifest
contract (dspark.py) had no execution-side reference to compare TT
against, and there was no deterministic torch model to generate the CPU
fixtures Package F compares the TT drafter to. Now:

- serving/speculative/dspark_reference.py: pure-torch DSparkReference
  implementing the exact pinned forward (5 dual-source GQA layers +
  fc/hidden_norm/norm + Markov/confidence heads), reusing the
  transformers Qwen3 building blocks (RoPE/RMSNorm/MLP) the pinned model
  imports so they match by construction. load_dspark_state_dict raises on
  any missing/extra key -- the "no silent drops" requirement. RoPE is
  computed over the full context+draft range because dual-source K spans
  both; the attention module slices the query portion.
- bench/probes/dspark-drafter-reference.py: two-part CPU fixture
  (performance null per the ttsim-is-not-performance rule). Structural =
  header-only manifest check on the real 3.7GB RadixArk checkpoint
  (62/62 tensors, correct shapes, BF16, sha
  2aff025f45823b40ebe726b9dfa40302f3512bd9a11c3a7347de32a567acd9a7);
  numerical = seeded small config proving the forward is bit-deterministic
  (logits/proposals/per-layer PCC = 1.0) and correct_len is right for a
  perfect vs perturbed target greedy.
- bench/runs/dspark-drafter-20260908T060024Z.jsonl: the two fixture
  records, performance: null.
- bench/probes/_d82_structural_driver.py: re-run tool for the real-
  checkpoint manifest check (pure stdlib, locates dspark.py from the repo
  or takes an uploaded copy path).
- tests/test_dspark_reference.py: 4 CPU fixtures -- no-silent-drops load,
  forward determinism (incl. per-layer outputs), accept_greedy boundary
  (full match accepts K; mismatch at k accepts k), build_block layout.

RoPE/RMSNorm/MLP reuse the real Qwen3 classes (not re-implemented): the
pinned model's RoPE is YaRN-corrected (attention_scaling 1.3466,
inv_freq scaled), and a hand-rolled base formula gave a 1.8e-3 error, so
reusing the class is exact by construction. DFlashModel is the pinned
class name; this module exports DSparkReference.

Verified: 4 new tests pass in the nix .#runtime; full suite shows no
regressions (the 8 test_dspark_harness failures are the documented
expat nixpkgs gap, identical with these files stashed -- the vllm-tt:k2
image has a full CPython build where all 13 pass). The real 27B
manifest was re-verified on cfx-llm2 (62/62, sha match).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The 27B target is hybrid: 8 full-attention + 24 GDN layers. A K+1 speculative
verify packs candidate positions into the head dim for the attention layers
(head-major SDPA with an additive causal mask, the shape packed-verify.py
proved works) and runs K+1 sequential recurrent steps for the GDN layers.
This is the attention half of the verify contract.

Before this commit we had proven the bare head-packed SDPA op in isolation
(packed-verify.py: PCC 0.9998, both paged and unpaged, K-independent) and the
GDN half at TP=2 (D8.1, a87053f). What was NOT proven was that a real
qwen36-style layer — q/k/v projections, RoPE at the right positions, per-head
RMSNorm, the head-packed SDPA core, o_proj, and the LM head — produces
per-position outputs and logits matching K+1 individual single-position decodes
when the K+1 candidates are packed into one op call. That is exactly the
'one target call verifies K+1 positions ... passes target-only parity'
criterion D8.3 requires.

The probe (bench/probes/dspark-verify-tp.py) builds a reduced 27B-shaped layer
(dummy weights, no checkpoint, per GOAL item 4) at TP=2 on a (1,2) mesh and
checks, for K in {1,3,7}, both paged and unpaged, that the packed K+1 output
(after o_proj and the LM head) matches K+1 sequential single-position layer
outputs and logits per-position. All six K/mode combinations pass at
worst-PCC 0.999996.

The 7 CPU tests (tests/test_dspark_verify.py) lock the geometric contract the
device probe depends on: head-major packing layout, causal mask boundaries,
RoPE norm-preservation, reference-decode agreement, output shapes, and
deterministic rebuild — so a regression in any of those fails on CPU before
the expensive 2-chip simulator run is needed.

N_KV=1 is Blackhole SDPA-decode's supported GQA mode, not 27B's 8 per chip;
documented in the run record. GDN half inherits from D8.1. Performance is
null: ttsim has no timing model (AGENTS.md).
Before: the repo carried two parallel model tracks — our dense Qwen3.8-27B
and the Ornith-1.5-35B-A3B (qwen3_5_moe) MoE bring-up. TT's MoE support is
poor and Ornith is irrelevant on this hardware, so the MoE track only added
noise: docs describing two models, a second docker image, a parallel bench
ladder, and a 26-file ornith/ tree.

Now: deleted all 26 Ornith files (docs, docker image, bench scripts, runs,
and the ornith/ source tree) and stripped every Ornith/MoE reference from the
active docs. The dense 27B content is fully preserved — the MTP head (a
27B feature), the DSpark/D8 speculative-decode track, GDN/Gated-DeltaNet,
bench results, and the vllm-tt:local image all stay. "Ours is the dense
hybrid Gated-DeltaNet family, not MoE" clarifications stay because they
sharpen the focus. bench/reset-card.sh's ttnn probe vehicle switched from the
deleted vllm-tt:ornith image to vllm-tt:local.

Frozen third-party source (research/sglang-pinned, research/tt-pinned) is
untouched — their MoE code is upstream, not ours.
Wip
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 8s
97e818ca70
Owner

Tenstorrent review — autoresearch/dspark-drafter-sep07 @ 97e818ca

Reviewed tip 2026-09-10 ~14:37 PT (Donach's Bot “Wip”). Branch may have moved — re-check SHA before acting. Related: lanes #20 (FIR drift checklist), #21 (serve/mesh gotchas).

Executive take

Useful CPU contract + evidence discipline. Main conflict: tip/BLACKHOLE_GUIDE make MTP primary, while approved lane 3 is align this PR with Thatch tip techniques (precise-native draft SDPA → proposal capture → commit-only GDN) before folded T16. DSpark on silicon is not at Thatch tip parity. Keep endpoint tok/s separate from offline committed TG.

What's strong

  • Honest programme (probe_passed ≠ gate_complete; third-party DGX/RadixArk marked)
  • StateTransaction: snapshot / shadow_commit / recompute — right host abstraction for commit-only GDN
  • E2E ordering on tip: verify → plan → commit_tokens → publish (+ rollback) — keep the regression that fails if reordered
  • Harness hygiene (locks, JSONL, immutable images)
  • Ornith MoE removed from this branch — good focus cut

Must-fix / decide

  1. One primary track. Dual full ports (MTP + DSpark) thrash. Shared verifier + one drafter adapter for the speed programme; park the other after a predeclared bake-off. Thatch's ledger: verifier cost (~76 ms/block tip) blocks 200 TG more than drafter branding.
  2. Don't promote gates early. D8.1 partial (proxy conv / not all-layer). D8.2 largely CPU + seeded taps. D8.3 ≠ full hybrid verify. D8.4 not established. MTP single-card head parity ≠ D8.4/D8.6.
  3. Serving proposer still BLOCKED. DSparkReference: proposals = base_logits.argmax then Markov only biases logits — do not wire those proposals into accept until sample_block / confidence serving sources are pinned.
  4. -tp / -tp2 filename lies. Several probes are pure Torch / CPU head-index math — banner or rename; never report as mesh/TT.
  5. Metric naming. Uncentered cosine still called “PCC” in places; require abs/rel error + token decisions.
  6. Strip merge noise. Multi‑MiB fixtures, source.tar.gz run bundles, generated/inspector|watcher/** — LFS/external before this is merge-reviewable.
  7. Stale PR body — tip is far past early harness-only description; update summary.
  8. ornith/ deleted but BLACKHOLE_GUIDE still points at local ornith/ — fix refs to external Lottolabs only.

Gap vs Thatch tip (~87 TG path)

Technique On #19 tip?
Precise-native draft SDPA (k_chunk_size=64) Missing
Proposal capture (+ allocation-order: verifier fixtures before capture) Missing (fixture capture ≠ serve capture)
Commit-only GDN (accepted-prefix publish only) Host SHADOW_COMMIT sketch only — ≠ device tip contract
Folded T16 target attention Deferred (lane 4a) — out of scope here
200 TG Not in scope — tip verify alone exceeds break-even

Should-fix

  • Rebase/notice lane‑2 gotchas from #21 if touching serve launchers
  • Don't re-litigate FIR/#53320 (on master via 0007; see #20)
  • Keep MTP as parallel adapter with disjoint ownership; no MTP+DSpark in one decode cycle
  • One bounded next DSpark task if lane 3 still owns the charter: tip-alignment doc + host stubs (DraftAttentionPolicy, proposal_trace, commit_only_gdn) fail-closed — or explicit owner override to MTP-primary in this PR description

Questions for the working agent

  1. Confirm tip SHA vs 97e818ca… — newer Wip?
  2. Owner call: MTP TP=2 head next, or Thatch DSpark #1 (precise-native draft SDPA) next?
  3. Where will serving sample_block / confidence sources be pinned (SGLang rev URL)?
  4. Is mtp_head_parity_tp2 green with a JSONL cite, or only scripted?
  5. Large fixtures: stay in git or artifact store before #19 review?

Suggested lane-3 slice (if charter holds)

  1. Doc/PR realignment — DSpark↔Thatch import; MTP parallel only
  2. Real capture contract (hashes, not seeded proxies)
  3. One TT draft attention block + precise-native SDPA vs that capture
  4. Device commit-only GDN for rejected suffix (shared verifier owner)
  5. Only then folded T16 / perf — quote tip TG and serve tok/s in separate columns

— Tenstorrent (on behalf of D. H. review ask)

## Tenstorrent review — `autoresearch/dspark-drafter-sep07` @ `97e818ca` Reviewed tip **2026-09-10 ~14:37 PT** (Donach's Bot “Wip”). Branch may have moved — re-check SHA before acting. Related: lanes [#20](https://git.bitp.cz/bitpartner/tt-stack/pulls/20) (FIR drift checklist), [#21](https://git.bitp.cz/bitpartner/tt-stack/pulls/21) (serve/mesh gotchas). ### Executive take Useful **CPU contract** + **evidence discipline**. Main conflict: tip/`BLACKHOLE_GUIDE` make **MTP primary**, while approved lane 3 is **align this PR with Thatch tip techniques** (precise-native draft SDPA → proposal capture → commit-only GDN) **before** folded T16. DSpark on silicon is **not** at Thatch tip parity. Keep endpoint **tok/s** separate from offline **committed TG**. ### What's strong - Honest programme (`probe_passed` ≠ `gate_complete`; third-party DGX/RadixArk marked) - `StateTransaction`: snapshot / **shadow_commit** / recompute — right host abstraction for commit-only GDN - E2E ordering on tip: verify → plan → `commit_tokens` → publish (+ rollback) — keep the regression that fails if reordered - Harness hygiene (locks, JSONL, immutable images) - Ornith MoE removed from this branch — good focus cut ### Must-fix / decide 1. **One primary track.** Dual full ports (MTP + DSpark) thrash. Shared verifier + one drafter adapter for the speed programme; park the other after a predeclared bake-off. Thatch's ledger: **verifier cost** (~76 ms/block tip) blocks 200 TG more than drafter branding. 2. **Don't promote gates early.** D8.1 partial (proxy conv / not all-layer). D8.2 largely CPU + seeded taps. D8.3 ≠ full hybrid verify. D8.4 not established. MTP single-card head parity ≠ D8.4/D8.6. 3. **Serving proposer still BLOCKED.** `DSparkReference`: `proposals = base_logits.argmax` then Markov only biases `logits` — do **not** wire those proposals into accept until `sample_block` / confidence serving sources are pinned. 4. **`-tp` / `-tp2` filename lies.** Several probes are pure Torch / CPU head-index math — banner or rename; never report as mesh/TT. 5. **Metric naming.** Uncentered cosine still called “PCC” in places; require abs/rel error + token decisions. 6. **Strip merge noise.** Multi‑MiB fixtures, `source.tar.gz` run bundles, `generated/inspector|watcher/**` — LFS/external before this is merge-reviewable. 7. **Stale PR body** — tip is far past early harness-only description; update summary. 8. **`ornith/` deleted** but `BLACKHOLE_GUIDE` still points at local `ornith/` — fix refs to external Lottolabs only. ### Gap vs Thatch tip (~87 TG path) | Technique | On #19 tip? | | --- | --- | | Precise-native **draft** SDPA (`k_chunk_size=64`) | Missing | | Proposal capture (+ allocation-order: verifier fixtures **before** capture) | Missing (fixture capture ≠ serve capture) | | Commit-only GDN (accepted-prefix publish only) | Host `SHADOW_COMMIT` sketch only — ≠ device tip contract | | Folded T16 target attention | Deferred (lane 4a) — out of scope here | | 200 TG | Not in scope — tip verify alone exceeds break-even | ### Should-fix - Rebase/notice lane‑2 gotchas from #21 if touching serve launchers - Don't re-litigate FIR/#53320 (on master via `0007`; see #20) - Keep MTP as parallel adapter with **disjoint** ownership; no MTP+DSpark in one decode cycle - One bounded next DSpark task if lane 3 still owns the charter: tip-alignment doc + host stubs (`DraftAttentionPolicy`, `proposal_trace`, `commit_only_gdn`) fail-closed — **or** explicit owner override to MTP-primary in this PR description ### Questions for the working agent 1. Confirm tip SHA vs `97e818ca…` — newer Wip? 2. Owner call: **MTP TP=2 head** next, or **Thatch DSpark #1 (precise-native draft SDPA)** next? 3. Where will serving `sample_block` / confidence sources be pinned (SGLang rev URL)? 4. Is `mtp_head_parity_tp2` green with a JSONL cite, or only scripted? 5. Large fixtures: stay in git or artifact store before #19 review? ### Suggested lane-3 slice (if charter holds) 1. Doc/PR realignment — DSpark↔Thatch import; MTP parallel only 2. Real capture contract (hashes, not seeded proxies) 3. One TT draft attention block + precise-native SDPA vs that capture 4. Device commit-only GDN for rejected suffix (shared verifier owner) 5. Only then folded T16 / perf — quote tip TG and serve tok/s in **separate** columns — Tenstorrent (on behalf of D. H. review ask)
D. H. override: reproduce Thatch tip P1 on our 2× P150a before inventing;
park MTP-as-primary; one next task is tip P1 inventory → smallest HW probe.
Delete inspector/watcher/test_reports dumps and nested bench/runs snapshot
trees (keep top-level JSONL evidence). Move parked MTP oracle .pt fixtures
to bench/fixtures/archive/.
fix: agent pickup (tp2 banners, proposer BLOCKED, P1 checklist)
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 7s
0118971f72
Rename CPU-only *-tp* probes with NOT-hardware docstring banners; fail-closed
guard so argmax placeholder proposals cannot feed accept; add
docs/THATCH-P1-CHECKLIST.md tip→#19 map.
Grok changed title from Prepare reproducible DSpark development on two P150a cards to Thatch-repro P1 on 2× P150a (override MTP-primary) — DSpark tip path 2026-09-10 17:28:11 +02:00
Two P1 rows on the Thatch-repro critical path. Neither promotes a gate; both
record performance=null and gate_complete=false.

1. FP32-intermediate draft SDPA, on our 2x P150a.

The checklist said the tip's C++ FP32-intermediate patch had not been built
because a C++ build was not viable on the 4-core bring-up box. That belief was
stale on two counts: the .so was built on the orchestra host (20c/47GB) and
shipped as a two-file overlay on vllm-tt:k2, and the resulting image
(vllm-tt:k2-fp32) had been sitting on cfx-llm2 unrecorded.

A/B, same probe both arms, only the .so differs. At HiFi4 the gate is observed
engaging in stderr -- not inferred -- and PCC goes 0.999983 -> 0.999990. At LoFi
the *unpatched* op fails the all-ones sanity outright (frac_near_1 = 0.0) and
the patch restores it to 1.0000, while both LoFi arms miss PCC_MIN=0.999.

What changed in our understanding: the patch is real and directionally correct,
but at HiFi4 it buys ~41% of a residual already at the bf16 floor. It is not
where the tip's ~87 TG comes from. Recorded as numerics only.

The A/B could not have run before: MFID was a dead knob, overwritten by a second
hard-coded MathFidelity.LoFi config, so every run was LoFi regardless.

2. Proposal capture -- the argmax placeholder is gone.

DraftOutput.proposals was argmax(base_logits), a placeholder. It is now
greedy_sample_block(), a port of the tip's dspark_markov.py::greedy_proposals.
Three semantics are load-bearing and none survive vectorisation: position p's
bias comes from the token actually selected at p-1; position 0 IS biased, by the
anchor; full vocabulary with no shortlist before the bias.

SAMPLE_BLOCK_STATUS stays BLOCKED on purpose. Real capture is necessary but not
sufficient -- commit-only GDN and the shared verifier are still missing, and
nothing here has been compared against captured hardware proposals. It never
will be token-for-token: grouped-BF16 products do not equal an FP32 matmul, and
the tip's own tests assert device token != fp32 token. This reproduces
semantics, not bits.

Believed before / what the evidence changed: the checklist named hybrid_draft,
lookup_draft, greedy_verify, greedy_session and lookup_acceptance as the tip's
capture path. They are not -- they are a host-side n-gram coordinator that never
touches a tensor, and the tip's README says so. The ~87 TG path is the neural
drafter with captured proposals plus this Markov sampler.

3. The test suite has been unrunnable since 0118971.

tests/test_dspark_reference.py:88 carried "p.clone(, allow_argmax_placeholder=
True)" -- a botched mechanical edit. It is an ast.parse failure, so it broke
pytest collection for the whole tests/ directory, not just that file. Fixed.
Suite: 97 pass, 8 fail, the 8 being the pre-existing pyexpat ImportError gap
already recorded on 2026-09-08.

Tip is at b2f9ffe (ci/qwen-hardware-correctness), verified the most recently
updated of all branches and PR refs. Their new T32 line has zero hardware
evidence -- CPU simulator and host unit tests only, their own hedge being that
those jobs "do not mount the cards or measure hardware speed". None of our
pinned P1 files changed upstream.

Next: commit-only GDN.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Third P1 row on the Thatch-repro critical path. CPU semantics only; no gate is
promoted, performance=null.

What we believed before: our transaction model did not yet reproduce the tip's
commit-only GDN semantics. The evidence says otherwise -- verify() already never
wrote live state, so behaviourally we matched. What we actually lacked was
ENFORCEMENT. The tip asserts the invariant every cycle; we asserted nothing, so
a regression would have been silent.

begin_verify() now captures a decoupled clone of live, and
_assert_verify_did_not_touch_live() runs at the top of commit_tokens() and
abort(). The clone matters: an identity check alone is fail-open because an
in-place mutation preserves id(), and that is the real hazard under
SHADOW_COMMIT, where _snapshot deliberately aliases live to save a copy per
cycle. Anyone who later "optimises" verify by advancing _snapshot in place now
gets a RuntimeError before any publication instead of silent corruption. Two
injected-regression tests prove the guard fires rather than merely existing.

The accept/reject boundary is pinned with exact equality, not PCC, plus an
explicit not-equal assertion between accept=0 and accept=1 -- without that
negative half both positive assertions would pass on an implementation that
ignored accepted_count entirely. Index-mapping trap recorded while porting: the
tip's packed-history row is prefix-1 because their row 0 is the state after
token 1, whereas our _boundaries[0] IS the pre-verify state, so ours is
_boundaries[a] with no offset.

Separately, _conv_pcc in the GDN transaction probe was fail-open twice over: it
skipped any conv tap that was None on either side, and returned 1.0 when nothing
was left to compare. A tap filled on one side and empty on the other is a real
divergence in the conv fill pattern -- precisely the off-by-one a commit-only
port can introduce -- and it was being dropped rather than compared. Structural
mismatch now scores 0.0, and the compared-tap count is returned so a vacuous
comparison is logged as VACUOUS instead of banked as a pass.

Two corrections to the checklist. optimisation/sim/gdn-commit.py is not the
contract -- it is a synthetic publication fixture, and test_commit_completion.py
is an AST meta-test over gdn-multitoken.py, not a state test; the semantics live
in scripts/ci/gdn_device_loop_state.py. And the newer gdn_native_slot_state.py,
gdn_commit_batched_dma.py and gdn_batched_publication_scope.py do not supersede
it: all three are self-declared unpromoted and simulator-only, and the batched
variant differs from the selected one by exactly one CB size value.

Deliberately not ported: the tip's no-rollback failure model, since our e2e.py
rollback is stronger. Cannot be ported at all without hardware: single-launch
DMA atomicity, native+checkpoint written from one staged buffer, the 32-byte
two-face NOC scatter and its tiled offset arithmetic, inactive-slot
non-interference in the shared 8-slot native tensor, per-chip mesh consistency.
The caveats block in the run record says so explicitly. Note also that
SNAPSHOT/RECOMPUTE still recompute at commit; only SHADOW_COMMIT is
selection-shaped like the tip, and the other two must not be described as
commit-only semantics.

Suite: 112 pass, 8 fail (the pre-existing pyexpat gap).

Next: shared verifier + GDN owner, the last open P1 row.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Fourth P1 row (shared verifier + GDN owner) advanced. The row is NOT closed --
what is closed is ownership and refusal semantics. CPU only; no gate promoted.

Two fail-open holes, both on the production serving loop.

A second GDN mutator existed. e2e.py's publication-failure path assigned
transaction.live = pre_live directly, which is a rollback engine living outside
the transaction -- the clearest violation of "keep single verifier owner" on our
side, since the tip's VerifierEngine.publish is the sole GDN mutator and is
gated on phase and ticket identity. That write now goes through
StateTransaction.rollback(), so every write to live state is owned by one class.
rollback() is deliberately not abort(): abort() discards a verify still in
progress, whereas by the time a publication fails the commit has already run and
_finish()-ed. Calling rollback() mid-verify would silently drop the in-flight
verify, so it fails closed there instead.

forced_acceptance_length sat unguarded on the accept path. It bypasses the
target comparison entirely, so a serving cycle could "accept" tokens the
verifier never agreed with. Every existing test passed an empty iterator, so
nothing actually used it -- it was pure fail-open surface, and the tip has no
analogue of it on the request loop at all. It is now refused unless a caller
opts in explicitly, and the serving params are pinned greedy, mirroring the
tip's refusal of seeded or non-argmax sampling at the verifier boundary.

Correction: this row listed bench/probes/dspark-verify-tp.py as the verifier. It
is a packed-SDPA kernel parity probe -- no argmax, no accept, no GDN, no state --
and its own record says "P150a x2 (simulator)". Its tip counterpart is
ModelBatch/verifier_engine.py, not greedy_verify.py.

Recorded as still open, because the row should not read as done: decision
ownership is still split across sampler.py, request._plan_commit and the e2e
loop where the tip has one owner; EOS semantics genuinely differ (the tip
returns early with state_rows == accepted and no correction token, we would
still emit a bonus); there is no max_proposals 15/31 bound; there is no
ticket-identity or epoch check; and our packed verify has never run on real
cards.

Also parameterises seed/CTX/SQ on the FP32 A/B probe so the single-point caveat
on our one hardware result can be swept rather than just stated.

Next: run the packed-verify probe on the real 2x P150a.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Closes the "single seed, single shape, single call" caveat on the one hardware
result we had. 18 runs on the real 2x P150a: 3 seeds x 3 shapes x 2 arms, all at
HiFi4, same probe script both arms, only the ttnn .so differing.

The effect is real and consistent: patched >= baseline in 9 of 9 paired
comparisons, with the engagement marker observed on all 9 patched runs and none
of the baseline runs. So the earlier single point was not seed noise.

What the sweep changes is the interpretation. The benefit is strongly
context-dependent and largely gone by CTX=4096: mean PCC-error reduction is
39.0% at CTX=2048 (n=6) but only 8.8% at CTX=4096 (n=3), with two of the three
4096 points at 3.9% and 5.0%. Absolute error also roughly doubles with context
on both arms. A serving drafter runs at long context, so this further weakens
the case that the FP32-intermediate patch is where the tip's ~87 TG comes from.
Sq (16 vs 32) had no visible effect.

Still numerics only: performance=null, gate_complete=false. The context finding
is about PCC error and says nothing about committed TG. The 4096 sample is n=3
and CTX 8192+ was not tested, so the direction is clear but the magnitude at
long context is thinly sampled. The build-provenance gap from the first A/B is
unchanged and still recorded.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
P1: get packed verify onto real cards -- it could not run there at all
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 7s
ba3e7ff42d
First hardware run of bench/probes/dspark-verify-tp.py. K in {1,3,7}, paged and
unpaged, worst PCC 0.999996, VERDICT: PASS on the 2x P150a.

It had never run on cards, and could not have. The probe put an implicitly-
concatenated string INSIDE an f-string expression, which parses only on Python
3.12+ (PEP 701). The nix runtime shell is 3.14, so it parsed there and every
prior run went through the simulator; the vllm-tt container is Python 3.10,
where it was a SyntaxError before the first line executed. That single
incompatibility is why this probe's whole recorded history is simulator runs.

Worse, the record it writes hard-coded "chips": "P150a x2 (simulator)". Running
it on hardware would have produced a false record claiming the opposite of what
happened. The label is now derived from TT_METAL_SIMULATOR, which is what
actually selects the backend, so neither direction can be mislabelled; the
timing-claim caveat is likewise conditioned on the real runtime.

What this is NOT: it is target-only packed-SDPA kernel parity -- the packed K+1
call reproducing K+1 sequential single-position calls. There is no argmax, no
accept decision, no GDN state and no commit anywhere in it. It does not close
the shared-verifier row, and the record says so; gate_complete=false,
performance=null.

Next: the single decision owner (select_prefix -> Decision), folding in the EOS
semantics where the tip emits no correction token but we currently would.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Member

P1 progress — 5 commits (4922f17..ba3e7ff)

Tip re-verified at b2f9ffe (ci/qwen-hardware-correctness). I enumerated every remote branch and PR ref: that branch is the most recently updated, 6 days ahead of main and of everything else, and is the head of their PR #7. None of our pinned P1 files changed upstream.

Row status

P1 row Before Now
Precise-native draft SDPA substrate green, patch "not built" closed — tip C++ patch reproduced on our cards + 18-point sweep
Proposal capture BLOCKED (argmax placeholder) ported (CPU) — placeholder gone, accept still blocked
Commit-only GDN host sketch / partial contract enforced (CPU)
Shared verifier + GDN owner partial advanced, NOT closed — one GDN owner; packed verify green on hardware

Hardware results

1. The tip's FP32-intermediate C++ patch is reproduced on our 2× P150a. The checklist said it couldn't be built here; that was stale. Image vllm-tt:k2-fp32 exists, and the gate is observed engaging in stderr, not inferred:

[DSpark P1] precise-intermediates gate ENGAGED: B=1 NQH=16 NKH=4 DHt=4 vDHt=4 Sq_chunk_t=1

18-run sweep (3 seeds × 3 shapes × 2 arms, HiFi4): patched ≥ baseline in 9/9 pairs, gate engaged on all 9 patched and zero baseline runs — consistent, not seed noise. But the benefit is strongly context-dependent: mean PCC-error reduction 39.0% at CTX=2048 (n=6) vs 8.8% at CTX=4096 (n=3), and absolute error roughly doubles with context on both arms. A serving drafter runs at long context, so this weakens the case that the patch is where the tip's ~87 TG comes from.

Separately, at LoFi the unpatched op fails the all-ones sanity outright (frac_near_1 = 0.0) and the patch restores it to 1.0000 — but both LoFi arms miss PCC_MIN=0.999, so LoFi isn't a usable drafter fidelity either way.

2. Packed verify is green on the real cards, for the first time. K ∈ {1,3,7}, paged and unpaged, worst PCC 0.999996.

Three pre-existing breakages, found by running things

  • The whole test suite has been uncollectable since 0118971. tests/test_dspark_reference.py:88 had p.clone(, allow_argmax_placeholder=True) — an ast.parse failure, so pytest collection died for the entire tests/ directory, not just that file.
  • dspark-verify-tp.py could never run in the container. It put an implicitly-concatenated string inside an f-string expression — valid only on Python 3.12+ (PEP 701). The nix shell is 3.14 so it parsed there; the container is 3.10, where it was a SyntaxError. That is why its entire recorded history is simulator runs. It also hard-coded "chips": "P150a x2 (simulator)" into its own record, so a hardware run would have produced a false record. Now derived from TT_METAL_SIMULATOR.
  • _conv_pcc was fail-open twice over. It skipped any conv tap that was None on either side and returned 1.0 when nothing remained. A tap filled on one side and empty on the other is a real fill-pattern divergence — exactly the off-by-one a commit-only port introduces — and it was dropped rather than compared.

Fail-open holes closed on the serving path

  • A second GDN mutator existed. e2e.py assigned transaction.live = pre_live directly on the publication-failure path — a rollback engine outside the transaction, the clearest violation of "keep single verifier owner". Now StateTransaction.rollback(), which fails closed if called mid-verify.
  • forced_acceptance_length sat unguarded on the accept path. It bypasses the target comparison entirely, so a cycle could "accept" tokens the verifier never agreed with. Every test passed an empty iterator, so nothing used it — pure fail-open surface. Refused now without explicit opt-in.

Corrections to this branch's own docs

  • The capture row named hybrid_draft.py / lookup_draft.py / greedy_verify.py / greedy_session.py / lookup_acceptance.py. Those are not the capture path — they're a host-side n-gram coordinator that never touches a tensor; their own README says "a host coordinator, not a TT device executor or a throughput result". The real path is scripts/ci/dspark_markov.py, ~10 lines.
  • optimisation/sim/gdn-commit.py is not the commit-only contract — it's a synthetic publication fixture, and test_commit_completion.py is an AST meta-test over gdn-multitoken.py. Semantics live in gdn_device_loop_state.py.
  • The verifier row listed dspark-verify-tp.py as the verifier. It's a kernel-parity probe — no argmax, no accept, no GDN, no state.

Do not chase T32

Their newest line (31 draft queries / 32 rows) has zero hardware evidence — two CPU-simulator correctness runs plus host unit tests. Their own words: "they do not mount the cards or measure hardware speed", "Not run; device-state and combined PP/CTX/TG admission still required." Unpromoted additional arm; reuses the precise-native draft SDPA unchanged. Their only real-hardware gain in that range is T16 gate/up fusion at +1.4–2.0% TG, which they say must not be presented as resolving 200 TG. Banked-proposal and native-slot-GDN arms are both recorded by them as not promoted, with no meaningful speedup.

What is NOT claimed

Every record is gate_complete=false, performance=null. Rows 2–3 are CPU semantics only. The commit-only mechanism cannot be validated without hardware: single-launch DMA atomicity, native+checkpoint from one staged buffer, the 32-byte two-face NOC scatter and its tiled offset arithmetic, inactive-slot non-interference, per-chip mesh consistency. No TG number here is ours, and none is implied.

Open provenance gap, recorded in the run file: the overlay Dockerfile asserts a b9bb5825 tt-metal build base while our records pin k2's at 1227e182, and neither .so exposes a version string to settle it.

Suite: 121 passed, 8 failed — the 8 being the pre-existing pyexpat ImportError gap.

Next

Single decision owner: select_prefix → Decision(emitted, accepted, state_rows, next_input, finished), with an invariant that commit_tokens refuses a token list it did not receive a matching Decision for. Folding in the EOS divergence — the tip emits no correction token when EOS lands inside the accepted prefix; we currently would.

🤖 Generated with Claude Code

## P1 progress — 5 commits (`4922f17..ba3e7ff`) Tip re-verified at **`b2f9ffe`** (`ci/qwen-hardware-correctness`). I enumerated every remote branch *and* PR ref: that branch is the most recently updated, 6 days ahead of `main` and of everything else, and is the head of their PR #7. None of our pinned P1 files changed upstream. ### Row status | P1 row | Before | Now | | --- | --- | --- | | Precise-native draft SDPA | substrate green, patch "not built" | **closed** — tip C++ patch reproduced on our cards + 18-point sweep | | Proposal capture | **BLOCKED** (argmax placeholder) | **ported (CPU)** — placeholder gone, accept still blocked | | Commit-only GDN | host sketch / partial | **contract enforced (CPU)** | | Shared verifier + GDN owner | partial | **advanced, NOT closed** — one GDN owner; packed verify green on hardware | ### Hardware results **1. The tip's FP32-intermediate C++ patch is reproduced on our 2× P150a.** The checklist said it couldn't be built here; that was stale. Image `vllm-tt:k2-fp32` exists, and the gate is *observed* engaging in stderr, not inferred: ``` [DSpark P1] precise-intermediates gate ENGAGED: B=1 NQH=16 NKH=4 DHt=4 vDHt=4 Sq_chunk_t=1 ``` 18-run sweep (3 seeds × 3 shapes × 2 arms, HiFi4): patched ≥ baseline in **9/9** pairs, gate engaged on all 9 patched and zero baseline runs — consistent, not seed noise. **But the benefit is strongly context-dependent:** mean PCC-error reduction **39.0% at CTX=2048** (n=6) vs **8.8% at CTX=4096** (n=3), and absolute error roughly doubles with context on *both* arms. A serving drafter runs at long context, so this weakens the case that the patch is where the tip's ~87 TG comes from. Separately, at **LoFi** the *unpatched* op fails the all-ones sanity outright (frac_near_1 = 0.0) and the patch restores it to 1.0000 — but both LoFi arms miss PCC_MIN=0.999, so LoFi isn't a usable drafter fidelity either way. **2. Packed verify is green on the real cards, for the first time.** K ∈ {1,3,7}, paged and unpaged, worst PCC **0.999996**. ### Three pre-existing breakages, found by running things - **The whole test suite has been uncollectable since `0118971`.** `tests/test_dspark_reference.py:88` had `p.clone(, allow_argmax_placeholder=True)` — an `ast.parse` failure, so pytest collection died for the entire `tests/` directory, not just that file. - **`dspark-verify-tp.py` could never run in the container.** It put an implicitly-concatenated string *inside* an f-string expression — valid only on Python 3.12+ (PEP 701). The nix shell is 3.14 so it parsed there; the container is 3.10, where it was a `SyntaxError`. That is why its entire recorded history is simulator runs. It also **hard-coded** `"chips": "P150a x2 (simulator)"` into its own record, so a hardware run would have produced a false record. Now derived from `TT_METAL_SIMULATOR`. - **`_conv_pcc` was fail-open twice over.** It skipped any conv tap that was `None` on *either* side and returned 1.0 when nothing remained. A tap filled on one side and empty on the other is a real fill-pattern divergence — exactly the off-by-one a commit-only port introduces — and it was dropped rather than compared. ### Fail-open holes closed on the serving path - **A second GDN mutator existed.** `e2e.py` assigned `transaction.live = pre_live` directly on the publication-failure path — a rollback engine outside the transaction, the clearest violation of "keep single verifier owner". Now `StateTransaction.rollback()`, which fails closed if called mid-verify. - **`forced_acceptance_length` sat unguarded on the accept path.** It bypasses the target comparison entirely, so a cycle could "accept" tokens the verifier never agreed with. Every test passed an empty iterator, so nothing used it — pure fail-open surface. Refused now without explicit opt-in. ### Corrections to this branch's own docs - The capture row named `hybrid_draft.py` / `lookup_draft.py` / `greedy_verify.py` / `greedy_session.py` / `lookup_acceptance.py`. **Those are not the capture path** — they're a host-side n-gram coordinator that never touches a tensor; their own README says "a host coordinator, not a TT device executor or a throughput result". The real path is `scripts/ci/dspark_markov.py`, ~10 lines. - `optimisation/sim/gdn-commit.py` is **not** the commit-only contract — it's a synthetic publication fixture, and `test_commit_completion.py` is an AST meta-test over `gdn-multitoken.py`. Semantics live in `gdn_device_loop_state.py`. - The verifier row listed `dspark-verify-tp.py` as the verifier. It's a kernel-parity probe — no argmax, no accept, no GDN, no state. ### Do not chase T32 Their newest line (31 draft queries / 32 rows) has **zero hardware evidence** — two CPU-simulator correctness runs plus host unit tests. Their own words: *"they do not mount the cards or measure hardware speed"*, *"Not run; device-state and combined PP/CTX/TG admission still required."* Unpromoted additional arm; reuses the precise-native draft SDPA unchanged. Their only real-hardware gain in that range is T16 gate/up fusion at +1.4–2.0% TG, which they say must not be presented as resolving 200 TG. Banked-proposal and native-slot-GDN arms are both recorded by them as not promoted, with no meaningful speedup. ### What is NOT claimed Every record is `gate_complete=false`, `performance=null`. Rows 2–3 are **CPU semantics only**. The commit-only *mechanism* cannot be validated without hardware: single-launch DMA atomicity, native+checkpoint from one staged buffer, the 32-byte two-face NOC scatter and its tiled offset arithmetic, inactive-slot non-interference, per-chip mesh consistency. No TG number here is ours, and none is implied. Open provenance gap, recorded in the run file: the overlay Dockerfile asserts a `b9bb5825` tt-metal build base while our records pin k2's at `1227e182`, and neither `.so` exposes a version string to settle it. Suite: **121 passed, 8 failed** — the 8 being the pre-existing `pyexpat` `ImportError` gap. ### Next Single decision owner: `select_prefix` → `Decision(emitted, accepted, state_rows, next_input, finished)`, with an invariant that `commit_tokens` refuses a token list it did not receive a matching `Decision` for. Folding in the EOS divergence — the tip emits no correction token when EOS lands inside the accepted prefix; we currently would. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
Record a silent wrong-answer SDPA config: LoFi + fp32_dest_acc_en
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 4s
tt-stack-ci / Report upstream drift (pull_request) Successful in 8s
ffd862d28d
Closes the open follow-up from the FP32-intermediate A/B. The unpatched-LoFi
all-ones failure was not a probe artifact. It is a real op bug on the pinned
tt-metal, and it is silent: no error, no warning, no NaN.

scaled_dot_product_attention at MathFidelity.LoFi with fp32_dest_acc_en=True
returns a UNIFORM CONSTANT at the wrong magnitude. All-ones q/k/v under a zero
additive mask must give exactly 1.0 everywhere; it gives 1.27344 at k_chunk=128
and 1.77344 at k_chunk=64. min == max == mean with zero variance, so it reads as
a plausible tensor rather than as garbage.

Attribution, each knob swept alone on real cards:
- LoFi alone is fine        -- LoFi, fp32_dest_acc_en=False -> 0.99219
- fp32 dest acc alone is fine -- HiFi2 -> 0.98438, HiFi4 -> 0.98828
- exp_approx_mode is irrelevant -- identical 1.27344 with it on and off
- HiFi4 without fp32 dest acc returned exactly 1.00000

The error grows with the number of online-softmax chunks over the context
(16 chunks -> +27.3%, 32 chunks -> +77.3%), which points at cross-chunk max/sum
rescaling. That is a hypothesis consistent with two k_chunk points, not
something confirmed in the kernel source, and the record says so.

This also reinterprets the earlier A/B: the LoFi column was never a fidelity
comparison, because the op itself was wrong there. The Thatch patch fixes it at
k_chunk=128 but only improves k_chunk=64 (1.77344 -> 0.94531, still 5.5% low and
still outside tolerance), and does not fire at all when exp_approx_mode=True.

Added to the "traps already paid for" list in CLAUDE.md, next to the existing
k_chunk_size=32 corruption note, since both are silent wrong-answer SDPA
configurations on this pin.

Not reported upstream and not checked against a newer tt-metal.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
P1: one accept-decision owner (select_prefix -> Decision)
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 6s
f43ffe205e
Completes the CPU half of the shared-verifier row. The accept decision was split
across three modules: sampler.sample picked the accept length, request's
_plan_commit RE-DERIVED accepted = min(...) and truncated for stop/budget, and
the e2e loop committed. Three owners can disagree, and the state advance could
follow a token list no decision described.

serving/speculative/verify.py is now the only accept decision. The tip's own
test_greedy_verify.py passes against it UNMODIFIED, 7/7 -- that, not our tests,
is the fidelity oracle.

The invariant: commit_tokens refuses any token list it did not receive a
matching Decision for. It stays the single write site, so the commit-only GDN
guard and the probes that patch it keep working, but the only legal way in is
commit_decision(ticket, decision), which requires the same ticket OBJECT issued
by verify() -- identity, not equality -- carrying a monotone epoch.

EOS: correcting what I claimed in the previous commit's docs. Adopting the tip's
no-bonus-on-EOS does NOT change our emitted stream. _plan_commit already cut at
the first stop token, which yields exactly the two lists select_prefix returns.
Confirmed by a 2,601-case comparison over every fixture the e2e suite generates
-- streams and committed GDN state identical in all of them -- and by the
stop-token and budget tests passing unchanged. The divergence was where the
decision lived, plus accepted_draft_tokens accounting, not output.

Real behaviour changes, deliberate: budget is sized BEFORE verification as the
tip does, rather than truncated after, so an over-budget Decision fails closed
and a one-token budget is a target-only step with zero proposals;
accepted_draft_tokens drops on a final truncated block, which is the honest
number; max_proposals is bounded to the tip's T16/T32 capacity instead of
arbitrary k. test_budget_truncates_accepted_accounting is replaced by
test_over_budget_decision_fails_closed -- the truncation it pinned WAS the
second decision owner.

A bug in the design, caught while implementing: a Decision.rejected property
defined as state_rows == accepted + 1 is wrong, because that also holds when
every proposal was accepted and the correction is just the bonus. Rejection is
not derivable from a Decision without the proposal count, which neither ours nor
the tip's carries. Removed rather than kept with a misleading definition.

Suite 151 pass (was 121), 8 pre-existing pyexpat failures. dspark-gdn-e2e.py
still VERDICT: PASS, byte-identical over 100 prompts x 3 mechanisms x forced
a in [0,3] -- end-to-end confirmation the new owner changes no output.

Two second-owner remnants remain, both documented in code and both refused on
the serving path: StateTransaction.commit(accepted_count, bonus) synthesizes a
Decision for probes that drive the transaction with no verifier, and
sampler._forced_decision does so for forced-acceptance tests. Separately,
dspark_reference.accept_greedy is still a third, unrelated batched accept path
(the D8.2 metric). Acceptance remains BLOCKED; SAMPLE_BLOCK_STATUS untouched.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
MTP: measure the 27B head's DRAFT TOKEN on hardware, not its cosine
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 5s
tt-stack-ci / Report upstream drift (pull_request) Successful in 6s
8b8a81883e
The owner unparked MTP. The existing mtp_head_parity.py compares hidden-state
taps and stops at final_norm, and it feeds the predictor layer an oracle-sourced
fc_out so the pre-fc fusion never ran on TT. Both gaps were recorded in the
worklog as the next boundary to close.

But the belief worth testing was different. A drafter emits a TOKEN. Cosine
0.9995 on `h` and 0.985 on `mlp_out` says nothing about an argmax over 248,320
logits, and a flipped draft is an acceptance-rate loss no downstream kernel can
recover. Nothing in the repo had asked that question, and MTP's margin is thin:
the third-party 27B figure is acceptance 2.04 against break-even 1.94.

mtp_draft_token_fidelity.py runs the WHOLE head on one P150a -- embedding, both
pre-fc RMSNorms, the [5120,10240] fc matmul, the gated full-attention layer, the
final norm and the shared LM head -- so no tap is oracle-sourced any more, and
compares the draft token against an exact float32 oracle.

Measured over 128 hidden states: top-1 agreement 106/128 = 82.8%, TT token in
the oracle top-5 127/128, mean oracle rank of the TT pick 0.22. Disagreements
concentrate in near-ties (median oracle margin 0.130 vs 0.718 on agreements) but
are NOT cleanly separated. So the TT dtype chain flips about one draft in six.

What this does NOT establish: the hidden states are seeded Gaussian proxies, and
they give abnormally flat logit landscapes -- precisely the regime where a small
numeric perturbation flips an argmax. 82.8% is most likely a pessimistic lower
bound. Re-running against real captured target hidden states is the next action
and decides whether MTP is worth a TP=2 port.

The probe voided itself once before producing this number: at n=16 a single 5%
perturbation raised agreement 11 -> 12, because a few-percent nudge is below the
TT dtype error itself. Replaced with a monotone sweep (25/18/7/0 at
0.05/0.25/1.0/4.0) plus a mismatched-oracle control (7/128 vs 106/128 paired)
and a float32 anchor against the committed fixture (argmax 79).

Two host traps recorded in CLAUDE.md, both paid for here: cfx-llm2 runs exactly
one tt-metal process (not per-card), and a probe gets ~5.5 GB of RAM because
tt-metal holds ~10 GB of hugepages -- upcasting embed_tokens and lm_head to f32
is a silent exit 137.

Also carries a pre-existing uncommitted correction to AGENT-GUIDE-TT-ITERATION.md
naming daniel@ as a valid cfx-llm2 account alongside deployer@.

Evidence: bench/runs/mtp-draft-token-20260911T144311Z.jsonl
(gate_complete=false, performance=null).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Member

MTP unparked (owner direction) — first hardware result: the 27B MTP head's draft token.

Commit 8b8a818. New probe bench/probes/mtp_draft_token_fidelity.py; evidence bench/runs/mtp-draft-token-20260911T144311Z.jsonl (gate_complete=false, performance=null).

Why a new probe. mtp_head_parity.py compares hidden-state taps (cosine/Pearson) and stops at final_norm, feeding the predictor layer an oracle-sourced fc_out — so the pre-fc fusion never ran on TT. But a drafter emits a token, and a 0.985-cosine mlp_out says nothing about an argmax over 248,320 logits. Nothing in the repo had asked that question.

This runs the whole head on one P150a — embedding, both pre-fc RMSNorms, the [5120,10240] mtp.fc matmul, the gated full-attention layer, final norm, shared LM head. No tap is oracle-sourced any more, which also closes the boundary the worklog flagged.

Measured, 128 hidden states:

  • top-1 draft-token agreement vs an exact float32 oracle: 106/128 = 82.8%
  • TT token in the oracle's top-5: 127/128
  • mean oracle rank of the TT pick: 0.22 (max 5)
  • oracle top-1 margin: median 0.130 on disagreement vs 0.718 on agreement — near-ties, but not cleanly separated

So the TT dtype chain flips roughly one draft in six. That matters because MTP's margin is thin — the third-party 27B figure is acceptance 2.04 against break-even 1.94.

What this does not establish. The hidden states are seeded Gaussian proxies. They give abnormally flat logit landscapes — exactly the regime where a small numeric perturbation flips an argmax. 82.8% is most likely a pessimistic lower bound. The next action is re-running against real captured 27B target hidden states; that substitution decides whether MTP is worth a TP=2 port.

The probe voided itself once before producing this number. At n=16 a single 5% perturbation of final raised agreement 11→12 — a few-percent nudge is below the TT dtype error itself — so the run refused its own 68.8% headline. Replaced with a monotone sweep (25/18/7/0 at 0.05/0.25/1.0/4.0 relative RMS), a mismatched-oracle control (7/128 vs 106/128 paired), and a float32 anchor against the committed fixture (token 9767, seed 42 → argmax 79).

Two host traps recorded in CLAUDE.md, both paid for here:

  • cfx-llm2 runs exactly one tt-metal process; a second dies on the sysmem NOC address space, and this is not per-card. Serialise on one lock.
  • A probe gets ~5.5 GB of RAM, not 15 GB — tt-metal holds ~10 GB of hugepages. embed_tokens and lm_head are BF16 [248320,5120]; upcasting both to f32 (as mtp_head_oracle.load_weights does) is a silent exit 137.

In parallel, an agy agent is measuring the DSpark drafter's acceptance rate against real 27B target logits in a sibling worktree, on the same predeclared criterion (mean accepted length ≥3.0 at K=7 justifies the GDN device-DMA port; below ~2.0 does not).

**MTP unparked (owner direction) — first hardware result: the 27B MTP head's draft token.** Commit `8b8a818`. New probe `bench/probes/mtp_draft_token_fidelity.py`; evidence `bench/runs/mtp-draft-token-20260911T144311Z.jsonl` (`gate_complete=false`, `performance=null`). **Why a new probe.** `mtp_head_parity.py` compares hidden-state taps (cosine/Pearson) and stops at `final_norm`, feeding the predictor layer an oracle-sourced `fc_out` — so the pre-fc fusion never ran on TT. But a drafter emits a **token**, and a 0.985-cosine `mlp_out` says nothing about an argmax over 248,320 logits. Nothing in the repo had asked that question. This runs the **whole** head on one P150a — embedding, both pre-fc RMSNorms, the [5120,10240] `mtp.fc` matmul, the gated full-attention layer, final norm, shared LM head. No tap is oracle-sourced any more, which also closes the boundary the worklog flagged. **Measured, 128 hidden states:** - top-1 draft-token agreement vs an exact float32 oracle: **106/128 = 82.8%** - TT token in the oracle's top-5: **127/128** - mean oracle rank of the TT pick: **0.22** (max 5) - oracle top-1 margin: median **0.130 on disagreement** vs **0.718 on agreement** — near-ties, but **not** cleanly separated So the TT dtype chain flips roughly one draft in six. That matters because MTP's margin is thin — the third-party 27B figure is acceptance 2.04 against break-even 1.94. **What this does not establish.** The hidden states are seeded Gaussian proxies. They give abnormally flat logit landscapes — exactly the regime where a small numeric perturbation flips an argmax. **82.8% is most likely a pessimistic lower bound.** The next action is re-running against real captured 27B target hidden states; that substitution decides whether MTP is worth a TP=2 port. **The probe voided itself once before producing this number.** At n=16 a single 5% perturbation of `final` *raised* agreement 11→12 — a few-percent nudge is below the TT dtype error itself — so the run refused its own 68.8% headline. Replaced with a monotone sweep (25/18/7/0 at 0.05/0.25/1.0/4.0 relative RMS), a mismatched-oracle control (7/128 vs 106/128 paired), and a float32 anchor against the committed fixture (token 9767, seed 42 → argmax 79). **Two host traps recorded in `CLAUDE.md`, both paid for here:** - cfx-llm2 runs exactly **one** tt-metal process; a second dies on the sysmem NOC address space, and this is *not* per-card. Serialise on one lock. - A probe gets ~5.5 GB of RAM, not 15 GB — tt-metal holds ~10 GB of hugepages. `embed_tokens` and `lm_head` are BF16 `[248320,5120]`; upcasting both to f32 (as `mtp_head_oracle.load_weights` does) is a silent exit 137. In parallel, an `agy` agent is measuring the **DSpark drafter's** acceptance rate against real 27B target logits in a sibling worktree, on the same predeclared criterion (mean accepted length ≥3.0 at K=7 justifies the GDN device-DMA port; below ~2.0 does not).
The owner asked for true uplift over the ~14.4 tok/s target-only baseline at
4k/8k/16k. Three things were in the way.

1. The baseline is a single short-prompt number. Decode cost on this stack is
   not flat in context -- the 16 full-attention layers read a KV cache that
   grows with it, the 48 GDN layers do not -- so 14.368 tok/s is not a
   denominator you can quote an uplift against at 16k. decode-bench.py now
   takes --prompt-tokens N, builds a deterministic prompt from a closed English
   vocabulary (not one repeated token, which prefix caching or attention can
   treat unlike real text), and records the server's OWN usage.prompt_tokens so
   the context a record claims is the context the tokenizer saw.

2. The goal1 serve profile hardcoded max_model_len 4096, and the profiles that
   do take a CTX override (goal2, goal3) also switch on bf8 KV and a bf4
   down_proj. A number measured there is not a valid denominator for this one.
   goal1 now honours CTX/SEQS and changes nothing else.

3. There is no speculative decoder to measure end to end, by either route:
   weight_mapping.py skips every mtp.* key at three separate places so the MTP
   weights are never loaded, and qwen36_vllm.py has no speculative path at all.
   On the DSpark side accept/commit is CPU-only. So the uplift has to be
   assembled from separately measured halves, and the risk is that someone
   quotes half of it as the speedup.

speculative-uplift.py makes that composition explicit and refuses to hide what
it is: uplift = (mean_accepted + 1) / (verify_cost + draft_cost), every field
labelled PROJECTION, every caveat named -- position-independent acceptance,
stationarity, cost ratios that do not survive a context change, commit cost
ignored. It also prints the costless ceiling, because if the ceiling is already
uninteresting no kernel work rescues the approach.

It will not project at all unless it first reproduces the one end-to-end 27B
figure anyone has measured: the third-party 1.03-1.05x at acceptance 2.04 and
break-even 1.94 (EXTERNAL-REPO-RUNBOOK.md Section 8). From their inputs the
model returns 1.052x. A projection model that cannot recover a known result
from its own inputs is not a model, and this one now says so before it prints.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
baseline: decode tok/s vs context on 2x P150a -- the uplift denominator
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 4s
tt-stack-ci / Report upstream drift (pull_request) Successful in 8s
0b900434dc
The owner asked for true uplift at 4k/8k/16k over the ~15 tok/s no-drafter
baseline. Half of that is deliverable today and half is not, so this commit
delivers the half that is and says clearly why the other half isn't.

Not deliverable: there is no speculative decoding in our serving stack by
either route. weight_mapping.py skips every mtp.* key at three separate places
so the MTP weights are never loaded, and qwen36_vllm.py has no speculative path
at all; on the DSpark side accept/commit is CPU-only. Nothing end-to-end can be
measured, and a projected number must not be dressed up as a measured one.

Deliverable: the denominator. The recorded 14.368 tok/s baseline is a single
short-prompt number, which is not something you can quote an uplift against at
16k without first knowing how decode behaves with depth.

Measured, goal1 profile at CTX=20480, 2x P150a TP=2, 256 tokens per rung:

     26 prompt tokens -> 14.083 tok/s   TTFT 0.515s
  4,073                  13.996         1.246
  8,108                  13.972         2.337
 16,177                  13.847         4.804

Decode is nearly FLAT in context: 1.7% decline across a 622x context increase.
That follows from the architecture -- 48 of 64 layers are Gated DeltaNet with
recurrent state, so only the 16 full-attention layers pay for depth. So one
denominator (~14.0 tok/s) serves every rung, a speculative win would be worth
as much at 16k as at 4k, and long context is not why we sit at 14 tok/s.

Secondary, and counterintuitive: max_model_len costs MORE than actual depth.
The same rungs on a CTX=16384 server ran ~1.4% faster at every depth purely
from the smaller KV pool -- larger than the entire context effect. A baseline
has to pin max_model_len, not just prompt length. Both ladders are in the
record.

The 26-token rung reproduces the 2026-09-04 goal-1 baseline to within 2% on a
different max_model_len with the same harness, so the ladder is anchored to a
reviewed number rather than free-floating.

Harness fix worth naming: the first prompt calibration produced a 13,146-token
prompt for a 16,384-token request, a 20% shortfall that would have mislabelled
every rung. decode-bench.py now records the server's own usage.prompt_tokens
plus the deviation and a within-5% flag, so a mislabelled rung shows up as a
flagged record instead of a wrong column header. Every rung now lands within
1.3% of target.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Member

Baseline context ladder measured — the uplift denominator. And the reason there is no uplift number yet.

Commits c70e86b (harness) and 0b90043 (measurement). Evidence: bench/runs/baseline-context-ladder-20260911T150910Z.jsonl.

The blocker, stated plainly

There is no speculative decoding in our serving stack, by either route. In vllm-tt:k2, models/demos/blackhole/qwen36/tt/weight_mapping.py skips every mtp.* key at three separate places (50, 135, 425) — the MTP head's weights are never loaded — and qwen36_vllm.py has no speculative path at all. On the DSpark side accept/commit is CPU-only and the device DMA commit is unported. So no end-to-end uplift is measurable today, and I am not going to present a projection as one.

What is measurable is the two halves that multiply into it. This is the first.

Measured — target-only decode vs context

2× P150a, mesh P300 (1,2) TP=2, vllm-tt:k2, goal1 profile at CTX=20480 SEQS=1, 256 tokens streamed per rung, greedy, single stream:

prompt tokens decode tok/s TTFT s p50 ms/tok
26 14.083 0.515 71.65
4,073 13.996 1.246 71.99
8,108 13.972 2.337 72.34
16,177 13.847 4.804 73.00

Decode is nearly flat in context — 1.7% decline across a 622× context increase. 48 of the 64 layers are Gated DeltaNet carrying recurrent state rather than a growing KV cache, so only the 16 full-attention layers pay for depth. TTFT scales as expected; it is prefill and excluded from the decode rate.

Consequences: one denominator (~14.0 tok/s) serves every rung; a speculative win would be worth as much at 16k as at 4k; and long context is not why we sit at 14 tok/s — per-token weight reading is.

Secondary, counterintuitive: max_model_len costs more than actual depth. The same rungs on a CTX=16384 server ran ~1.4% faster at every depth, purely from the smaller KV pool — larger than the entire context effect. A baseline has to pin max_model_len, not just prompt length. Both ladders are in the record.

Tie-back: the 26-token rung reproduces the 2026-09-04 goal-1 baseline (14.368 / 14.418) to within 2%, same harness, different max_model_len.

Harness

  • bench/decode-bench.py --prompt-tokens N builds a deterministic prompt and records the server's own usage.prompt_tokens. The first calibration silently produced a 13,146-token prompt for a 16,384-token request — a 20% shortfall that would have mislabelled every rung. It now also records the deviation and a within-5% flag, so a mislabelled rung is a flagged record rather than a wrong column header. All rungs now land within 1.3%.
  • scripts/vllm-tt-serve.sh goal1 honours CTX/SEQS and changes nothing else. This matters: goal2/goal3 also enable bf8 KV and a bf4 down_proj, so a number measured there is not a valid denominator.
  • bench/speculative-uplift.py composes uplift = (mean_accepted + 1) / (verify_cost + draft_cost), labels everything PROJECTION, and prints the costless ceiling and break-even. It refuses to project unless it first reproduces the only end-to-end 27B figure anyone has measured — the third-party 1.03–1.05× at acceptance 2.04 / break-even 1.94 — from their own inputs. It returns 1.052×.

That anchor is sobering: their break-even of 1.94 means verify + draft cost nearly two target steps at K=1, which is why 2.04 accepted tokens bought only 3–5%.

Next

Measure verify_cost_ratio on hardware — time one packed verify of K+1 positions against one target decode step at these same depths. Packed verify is already green on the cards, so this is a small probe, and it is the last unknown in the formula that does not require a speculative decoder to exist. The agy agent has been redirected to report DSpark acceptance per rung at 4k/8k/16k so its numerator composes with this denominator.

**Baseline context ladder measured — the uplift denominator. And the reason there is no uplift number yet.** Commits `c70e86b` (harness) and `0b90043` (measurement). Evidence: `bench/runs/baseline-context-ladder-20260911T150910Z.jsonl`. ### The blocker, stated plainly **There is no speculative decoding in our serving stack, by either route.** In `vllm-tt:k2`, `models/demos/blackhole/qwen36/tt/weight_mapping.py` skips every `mtp.*` key at **three** separate places (50, 135, 425) — the MTP head's weights are never loaded — and `qwen36_vllm.py` has no speculative path at all. On the DSpark side accept/commit is CPU-only and the device DMA commit is unported. So no end-to-end uplift is measurable today, and I am not going to present a projection as one. What *is* measurable is the two halves that multiply into it. This is the first. ### Measured — target-only decode vs context 2× P150a, mesh `P300 (1,2)` TP=2, `vllm-tt:k2`, **goal1** profile at `CTX=20480 SEQS=1`, 256 tokens streamed per rung, greedy, single stream: | prompt tokens | decode tok/s | TTFT s | p50 ms/tok | | ---: | ---: | ---: | ---: | | 26 | **14.083** | 0.515 | 71.65 | | 4,073 | **13.996** | 1.246 | 71.99 | | 8,108 | **13.972** | 2.337 | 72.34 | | 16,177 | **13.847** | 4.804 | 73.00 | **Decode is nearly flat in context — 1.7% decline across a 622× context increase.** 48 of the 64 layers are Gated DeltaNet carrying recurrent state rather than a growing KV cache, so only the 16 full-attention layers pay for depth. TTFT scales as expected; it is prefill and excluded from the decode rate. Consequences: one denominator (**~14.0 tok/s**) serves every rung; a speculative win would be worth as much at 16k as at 4k; and long context is **not** why we sit at 14 tok/s — per-token weight reading is. **Secondary, counterintuitive:** `max_model_len` costs *more* than actual depth. The same rungs on a `CTX=16384` server ran **~1.4% faster at every depth**, purely from the smaller KV pool — larger than the entire context effect. A baseline has to pin `max_model_len`, not just prompt length. Both ladders are in the record. **Tie-back:** the 26-token rung reproduces the 2026-09-04 goal-1 baseline (14.368 / 14.418) to within 2%, same harness, different `max_model_len`. ### Harness - `bench/decode-bench.py --prompt-tokens N` builds a deterministic prompt and records the server's **own** `usage.prompt_tokens`. The first calibration silently produced a **13,146**-token prompt for a 16,384-token request — a 20% shortfall that would have mislabelled every rung. It now also records the deviation and a within-5% flag, so a mislabelled rung is a flagged record rather than a wrong column header. All rungs now land within 1.3%. - `scripts/vllm-tt-serve.sh goal1` honours `CTX`/`SEQS` and changes nothing else. This matters: `goal2`/`goal3` also enable bf8 KV and a bf4 `down_proj`, so a number measured there is **not** a valid denominator. - `bench/speculative-uplift.py` composes `uplift = (mean_accepted + 1) / (verify_cost + draft_cost)`, labels everything `PROJECTION`, and prints the costless ceiling and break-even. It **refuses to project** unless it first reproduces the only end-to-end 27B figure anyone has measured — the third-party 1.03–1.05× at acceptance 2.04 / break-even 1.94 — from their own inputs. It returns 1.052×. That anchor is sobering: their break-even of 1.94 means verify + draft cost nearly **two** target steps at K=1, which is why 2.04 accepted tokens bought only 3–5%. ### Next Measure `verify_cost_ratio` on hardware — time one packed verify of K+1 positions against one target decode step at these same depths. Packed verify is already green on the cards, so this is a small probe, and it is the last unknown in the formula that does not require a speculative decoder to exist. The `agy` agent has been redirected to report DSpark acceptance **per rung at 4k/8k/16k** so its numerator composes with this denominator.
goal: MTP to >= 30 tok/s -- charter, with the arithmetic that rules K=1 out
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 8s
eba43de291
Owner directive: get MTP fully working and reach at least 30 tok/s on the same
benchmark that currently measures ~14. This charter defines the target so it
cannot be argued with -- same harness, same goal1 profile, same 4k/8k/16k rungs,
byte-identical greedy output as a pass condition -- and then states two facts
that constrain every plan built on it.

First: MTP-1 is arithmetically disqualified. A K=1 block emits at most 2 tokens,
so uplift <= 2.0 even with costless drafting and verification, which caps it at
27.99 / 27.94 / 27.69 tok/s at the three rungs. All below 30. K >= 2 is
mandatory, which is the single biggest divergence from the Lottolabs reference
whose device cycle is K=1 only, and it means the head runs autoregressively on
its own output -- so drafter numeric error compounds, and step-K fidelity rather
than step-1 fidelity sets acceptance.

Second: 30 tok/s is INSIDE the target-only roofline. The runbook's measured
bandwidth analysis puts the practical target-only ceiling at ~36-38 tok/s and
the floor at 26-28 ms/step; we run at 71-73 ms. We are at roughly 40% of what
this hardware does without any speculation. So the honest route is MTP x kernel
efficiency, not MTP alone, and a programme that ignores per-token weight-read
cost will likely build a correct speculative decoder that turns 14 into 21 and
stops.

Hence Gate 0, before any porting work: measure verify_cost_ratio on hardware for
K in {1,2,3,7} at all three rungs. Packed verify is already green on the cards so
this is instrumentation, not new kernel work, and it is decisive -- decode is
weight-read bound, so verifying K+1 positions reads the same weights once and the
ratio should sit near 1.0. If instead it lands near the third-party 1.94, then 30
tok/s needs 3.2 accepted tokens per block and no MTP-derived drafter has shown
that. Predeclared: <= 1.4 at K=3 proceed, > 1.8 stop and report.

The ladder after that is M1 load the mtp.* weights (weight_mapping.py skips them
at three places today), M2 TP=2 head, M3 draft-token fidelity at K>1 on real
hidden states, M4 verify_K, M5 the FusedCommit re-derivation for 48 v-heads and
a K-way rather than dual-candidate selection -- named as the largest risk -- and
M6 the benchmark itself.

Recorded at precedence line 0 in CLAUDE.md; supersedes the "no MTP TP=2 head
until P1 is done" line.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
goal: reframe 30 tok/s as two levers -- and MTP is the second one
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 4s
tt-stack-ci / Report upstream drift (pull_request) Successful in 8s
5ccce9c797
The directive was "get MTP fully working and reach 30 tok/s". Writing the
charter surfaced arithmetic that makes that framing unworkable, the owner
agreed to reframe, and this is the reframed charter. GOAL-MTP-30TPS.md is
renamed to GOAL-30TPS.md because it is no longer an MTP goal.

tok/s = (1000 / ms_per_step) x uplift. Our step is 71.99 ms at 4k. Against the
runbook's measured bandwidth analysis that decomposes as roughly 11 ms of GDN
small-op launch floor (48 layers x 68 ops x ~3.4 us), 26-28 ms of weight read at
peak DRAM, and ~33-35 ms of unexplained kernel inefficiency. The target-only
roofline is ~36-38 tok/s, so we run at about 40% of what these cards do with no
speculation at all, and Lever A alone can carry the goal.

The part that changes the plan rather than just the estimates: the two levers
are NOT independent. verify_cost_ratio = (t_weights + t_fixed') / (t_weights +
t_fixed). Verifying K+1 positions reads the same weights once -- that sharing is
the entire reason speculation can win -- but fixed cost does not amortise. When
weight read dominates the ratio tends to 1.0 and speculation pays richly; when
fixed cost dominates it is poor. We have ~44 ms of a 72 ms step outside the
weight read, so we are in the second regime. That is the mechanical explanation
of the one measured 27B result on this class of machine: 1.03-1.05x at
break-even 1.94, i.e. verify plus draft costing nearly two target steps at K=1.
The same runbook said MTP "is not a lever on this stack until the per-layer
fixed cost falls"; this charter now explains why and orders the work by it.

So Lever A (per-step cost) comes first, and it is the precondition for Lever B
(MTP) being worth anything, not an alternative to it. Close half the gap and the
step is 49.5 ms = 20.2 tok/s, after which 30 needs only 1.48x -- and K=1 stops
being disqualified, so Lever A also changes which speculative designs are
admissible.

Gate 0 is now two measurements, both cheap and neither needing a speculative
decoder: 0a our OWN per-step breakdown, because inheriting a third-party split is
guesswork; 0b verify_cost_ratio for K in {1,2,3,7}. 0b is re-run after every
Lever-A milestone -- it is the best progress indicator the programme has, since
it measures whether the step has become weight-read-bound.

One immediately testable A1 lead recorded: the goal2 profile enables
QWEN_GDN_FUSED_DECODE, QWEN_GDN_CONV_GATES, QWEN_GDN_FUSED_INPLACE and
QWEN_GDN_PACKED_QKV, and goal1 -- the baseline profile -- does not.

MTP-1 stays disqualified at today's step cost (ceiling 27.69 tok/s at 16k), K>=2
remains mandatory, and step-K rather than step-1 draft fidelity sets acceptance.
Section 6 now requires every result to report per-step cost and uplift
separately, because a bare tok/s figure hides which lever moved and whether the
denominator moved with it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Grok force-pushed autoresearch/dspark-drafter-sep07 from 5ccce9c797
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 4s
tt-stack-ci / Report upstream drift (pull_request) Successful in 8s
to 85b9065737
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 7s
2026-09-11 17:22:20 +02:00
Compare
Member

Goal reframed: ≥ 30 tok/s is a two-lever goal, and MTP is the second lever. docs/GOAL-30TPS.md, commit 5ccce9c, installed at precedence line 0 in CLAUDE.md.

The directive was "get MTP fully working and reach 30 tok/s". Writing the charter surfaced arithmetic that makes that framing unworkable.

The decomposition

tok/s = (1000 / ms_per_step) × uplift. Our step is 71.99 ms at 4k. Against the runbook's measured bandwidth analysis:

component ms source
GDN small-op launch floor (48 × 68 ops × ~3.4 µs) ~11.1 third-party
weight read at ~430 GB/s peak DRAM 26–28 third-party
kernel inefficiency (residual) ~33–35 —
total 71.99 ours, measured

The target-only roofline is ~36–38 tok/s. We run at roughly 40% of what these cards do with no speculation at all, so Lever A alone can carry the goal.

The part that changes the plan, not just the estimates

The two levers are not independent:

verify_cost_ratio = (t_weights + t_fixed') / (t_weights + t_fixed)

Verifying K+1 positions reads the same weights once — that sharing is the entire reason speculation can win. Fixed cost does not amortise. When the weight read dominates, the ratio tends to 1.0 and speculation pays richly; when fixed cost dominates, it is poor. We have ~44 ms of a 72 ms step outside the weight read.

That is the mechanical explanation of the one measured 27B result on this class of machine: 1.03–1.05× at break-even 1.94 — verify plus draft costing nearly two target steps at K=1 is exactly the signature of a fixed-cost-dominated step. The runbook already said MTP "is not a lever on this stack until the per-layer fixed cost falls"; the charter now explains why and orders the work by it.

So Lever A comes first, and it is the precondition for Lever B being worth anything — not an alternative to it.

per-step cost target-only tok/s uplift still needed for 30
71.99 ms (today) 13.9 2.17×
49.5 ms (half the gap) 20.2 1.48×
27.0 ms (roofline) 37.0 none

Lever A also changes which speculative designs are admissible: MTP-1 is disqualified at today's step cost (ceiling 27.69 tok/s at 16k, so K ≥ 2 is mandatory), but at a 49.5 ms step K=1's ceiling is 40 tok/s.

Gate 0 — two cheap measurements, before any porting

  • 0a — our own per-step breakdown. Inheriting a third-party split is guesswork, and this is the work list for Lever A.
  • 0b — verify_cost_ratio for K ∈ {1,2,3,7} at all three rungs. Packed verify is already green on the cards, so this is instrumentation. Predeclared: ≤ 1.4 at K=3 → build Lever B now; > 1.8 → do not port a commit kernel, do Lever A and re-measure.

0b is re-run after every Lever-A milestone. It is the programme's best progress indicator, because it measures whether the step has become weight-read-bound.

One immediately testable A1 lead: the goal2 profile enables QWEN_GDN_FUSED_DECODE, QWEN_GDN_CONV_GATES, QWEN_GDN_FUSED_INPLACE and QWEN_GDN_PACKED_QKV, and goal1 — the baseline profile — does not.

§6 now requires every result to report per-step cost and uplift separately: a bare tok/s figure hides which lever moved and whether the denominator moved with it.

**Goal reframed: ≥ 30 tok/s is a two-lever goal, and MTP is the second lever.** `docs/GOAL-30TPS.md`, commit `5ccce9c`, installed at precedence line 0 in `CLAUDE.md`. The directive was "get MTP fully working and reach 30 tok/s". Writing the charter surfaced arithmetic that makes that framing unworkable. ### The decomposition `tok/s = (1000 / ms_per_step) × uplift`. Our step is **71.99 ms** at 4k. Against the runbook's measured bandwidth analysis: | component | ms | source | | --- | ---: | --- | | GDN small-op launch floor (48 × 68 ops × ~3.4 µs) | ~11.1 | third-party | | weight read at ~430 GB/s peak DRAM | 26–28 | third-party | | **kernel inefficiency (residual)** | **~33–35** | — | | total | **71.99** | **ours, measured** | The target-only roofline is ~36–38 tok/s. **We run at roughly 40% of what these cards do with no speculation at all**, so Lever A alone can carry the goal. ### The part that changes the plan, not just the estimates The two levers are **not independent**: ``` verify_cost_ratio = (t_weights + t_fixed') / (t_weights + t_fixed) ``` Verifying K+1 positions reads the **same weights once** — that sharing is the entire reason speculation can win. Fixed cost does not amortise. When the weight read dominates, the ratio tends to 1.0 and speculation pays richly; when fixed cost dominates, it is poor. **We have ~44 ms of a 72 ms step outside the weight read.** That is the mechanical explanation of the one measured 27B result on this class of machine: **1.03–1.05× at break-even 1.94** — verify plus draft costing nearly *two* target steps at K=1 is exactly the signature of a fixed-cost-dominated step. The runbook already said MTP "is not a lever on this stack until the per-layer fixed cost falls"; the charter now explains why and orders the work by it. **So Lever A comes first, and it is the precondition for Lever B being worth anything** — not an alternative to it. | per-step cost | target-only tok/s | uplift still needed for 30 | | ---: | ---: | ---: | | 71.99 ms (today) | 13.9 | 2.17× | | 49.5 ms (half the gap) | 20.2 | **1.48×** | | 27.0 ms (roofline) | 37.0 | **none** | Lever A also changes which speculative designs are admissible: MTP-1 is disqualified at today's step cost (ceiling **27.69 tok/s** at 16k, so K ≥ 2 is mandatory), but at a 49.5 ms step K=1's ceiling is 40 tok/s. ### Gate 0 — two cheap measurements, before any porting - **0a** — our **own** per-step breakdown. Inheriting a third-party split is guesswork, and this is the work list for Lever A. - **0b** — `verify_cost_ratio` for K ∈ {1,2,3,7} at all three rungs. Packed verify is already green on the cards, so this is instrumentation. Predeclared: ≤ 1.4 at K=3 → build Lever B now; > 1.8 → **do not port a commit kernel**, do Lever A and re-measure. 0b is re-run after every Lever-A milestone. It is the programme's best progress indicator, because it measures whether the step has become weight-read-bound. One immediately testable A1 lead: the `goal2` profile enables `QWEN_GDN_FUSED_DECODE`, `QWEN_GDN_CONV_GATES`, `QWEN_GDN_FUSED_INPLACE` and `QWEN_GDN_PACKED_QKV`, and `goal1` — the baseline profile — does not. §6 now requires every result to report per-step cost **and** uplift separately: a bare tok/s figure hides which lever moved and whether the denominator moved with it.
docs/GOAL-30TPS.md makes byte-identity with the target-only greedy stream a PASS
CONDITION -- a change that alters the model's output has not sped anything up.
The harness could not check it: it kept only the first 200 characters as a
human-readable sample.

Every record now carries text_sha256 unconditionally, and --full-text stores the
complete string for diffing two arms of an A/B. This matters immediately for the
Lever A1 experiment: the GDN fusion env vars change the computation, so a faster
arm with a different sha256 is a regression wearing a speedup's clothes.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Gate 0a part 1: our weight-read ceiling is 31.6 tok/s, not 36-38 -- charter wrong
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 4s
tt-stack-ci / Report upstream drift (pull_request) Successful in 7s
ad68f325b0
The charter I wrote this morning said "30 tok/s is inside this hardware's
roofline (~36-38 tok/s target-only)... the goal does not depend on speculation
succeeding." That was wrong, and this corrects it.

The 36-38 figure came from EXTERNAL-REPO-RUNBOOK.md section 8. It is a
third-party number computed for a bf4-MLP configuration. The goal1 profile --
the one the 13.9 tok/s baseline was actually measured in -- runs all bf8.

Recounting our own checkpoint's safetensors headers (no device, no weight
download): 25.625 B of the 26.896 B target parameters are streamed every decode
token (embed_tokens is a one-row gather; lm_head is read in full). At bf8 that
is 27.23 GB, or 13.61 GB per card at TP=2:

  all bf8 (goal1 baseline)  13.61 GB/card  31.66 ms  ->  31.59 tok/s ceiling
  bf8, down_proj bf4        12.19 GB/card  28.34 ms  ->  35.28 tok/s
  bf8, all MLP bf4           9.33 GB/card  21.71 ms  ->  46.06 tok/s

The last row reproduces SATURATION-AND-SRAM.md line 249 -- 9.333 GB/card and
21.71 ms at 430 GB/s -- exactly, from an independent recount. That anchors the
method to a previously reviewed figure rather than to itself.

So the conclusion flips. In the baseline configuration the absolute Lever-A-only
ceiling is 31.59 tok/s: 5% above target, at 100% of peak DRAM with zero fixed
cost, and 27.9 tok/s at a realistic 380 GB/s. 30 tok/s is NOT reachable by Lever
A alone. It needs quantisation (A3) and/or speculation (B), which raises MTP from
"probably unnecessary" back to "likely required" and makes A3 -- with the
accuracy arm it implies -- a first-class part of the plan.

Lever A still goes first: our measured 71.99 ms step is 31.66 ms of weight read
and 40.33 ms (56%) of fixed cost and inefficiency, so there is 2x on the table
and closing it is also what makes speculation pay. It just ends at a wall short
of the target.

430 GB/s is still not ours. It is a synthetic aggregate; decode matmuls run on
~33 cores and nobody has measured what that grid pulls from DRAM on this
silicon. SATURATION-AND-SRAM.md line 387 already named this the unknown that
"bounds everything else", and bench/dram-saturation.py exists for it and has
never been run. That is Gate 0a part 2, and it can only lower these ceilings.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Gate 0a part 2: measure DRAM bandwidth; 30 tok/s is unreachable target-only
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 6s
tt-stack-ci / Report upstream drift (pull_request) Successful in 6s
5b0093819a
bench/dram-saturation.py had existed for weeks and had never been run.
SATURATION-AND-SRAM.md:387 called it the unknown that "bounds everything
else". It is now measured, on the card, model not loaded, device exclusive.

We believed 430 GB/s. It does not exist:

  bfloat16  402.9 GB/s peak    356.4 at 32 cores
  bfloat8_b 378.6 GB/s peak    331.1 at 32 cores
  bfloat4_b 287.9 GB/s peak    251.8 at 32 cores

Three things follow, all of them bad news, which is why they are worth
recording.

1. bfloat4_b is bandwidth-INEFFICIENT -- 24% below bf8. We had been
   treating quantisation as a byte count. It is not: bf4's 47% byte
   saving buys far less time than it appears to, which demotes A3 from
   escape hatch to marginal gain.

2. Bandwidth saturates by ~44 cores, and 8x8=64 cores is SLOWER than
   11x4=44 for every dtype. Grid shape, not core count, is the knob.

3. Every ceiling in Gate 0a part 1 moves down. The goal1 all-bf8
   target-only ceiling is 27.81 tok/s at measured peak and 24.32 on
   decode's actual grid -- below the 30 tok/s target with zero fixed
   cost and a perfect kernel. Even all-MLP-bf4, which we do not have
   and which costs accuracy, reaches only 34.89 / 30.51.

So no target-only path reaches 30 tok/s. Speculation is no longer the
optional half of the charter; it is load-bearing, with nothing behind
it. Lever A remains the first move only because speculation cannot pay
until the step is weight-read-bound.

This charter has now been corrected twice in one day by its own gates,
both times downward, both times because an inherited third-party figure
did not survive measurement.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Member

Gate 0a part 2 — DRAM bandwidth measured; the target-only path is closed

bench/dram-saturation.py had never been run. SATURATION-AND-SRAM.md:387 called it the unknown that bounds everything else. Now measured on the card, model not loaded, device exclusive — bench/runs/gate0a-dram-saturation-20260911T183617Z.jsonl.

430 GB/s does not exist.

dtype peak at ~32 cores (decode's grid)
bfloat16 402.9 GB/s 356.4
bfloat8_b 378.6 GB/s 331.1
bfloat4_b 287.9 GB/s 251.8

Three findings, all unwelcome, which is why they matter:

  1. bf4 is bandwidth-inefficient — 24% below bf8. We had been treating quantisation as a byte count. It isn't: bf4's 47% byte saving buys far less time than it looks like it should. A3 drops from escape hatch to marginal gain.
  2. Grid shape, not core count, is the knob. Saturation arrives by 44 cores, and 8x8=64 is slower than 11x4=44 for every dtype.
  3. Every Gate 0a part 1 ceiling moves down.
config @430 (assumed) @ measured peak @ measured 32-core
all bf8 — what goal1 actually runs 31.59 27.81 24.32
bf8, down_proj bf4 35.28 29.83 26.09
bf8, all MLP bf4 46.06 34.89 30.51

The goal1 target-only ceiling is 27.81 tok/s at peak and 24.32 on decode's own grid — below the 30 tok/s target with zero fixed cost and a perfect kernel. Even all-MLP-bf4, which we don't have and which costs accuracy, reaches 34.89 / 30.51.

So speculation is no longer the optional half of the charter — it is load-bearing, with nothing behind it. Lever A stays the first move only because speculation cannot pay until the step is weight-read-bound.

This charter has now been corrected twice in one day by its own gates, both downward, both because an inherited third-party figure did not survive measurement. Commit 5b00938.

Still running: A1 GDN-fusion A/B (holds the lock now), and the DSpark acceptance-rate sweep whose live log is showing accepted=0 on 14 of its first 15 blocks at 4k/K=15.

## Gate 0a part 2 — DRAM bandwidth measured; the target-only path is closed `bench/dram-saturation.py` had never been run. `SATURATION-AND-SRAM.md:387` called it the unknown that *bounds everything else*. Now measured on the card, model not loaded, device exclusive — `bench/runs/gate0a-dram-saturation-20260911T183617Z.jsonl`. **430 GB/s does not exist.** | dtype | peak | at ~32 cores (decode's grid) | | --- | ---: | ---: | | bfloat16 | 402.9 GB/s | 356.4 | | bfloat8_b | **378.6 GB/s** | **331.1** | | bfloat4_b | **287.9 GB/s** | 251.8 | Three findings, all unwelcome, which is why they matter: 1. **bf4 is bandwidth-*inefficient*** — 24% below bf8. We had been treating quantisation as a byte count. It isn't: bf4's 47% byte saving buys far less time than it looks like it should. A3 drops from escape hatch to marginal gain. 2. **Grid shape, not core count, is the knob.** Saturation arrives by 44 cores, and 8x8=64 is *slower* than 11x4=44 for every dtype. 3. **Every Gate 0a part 1 ceiling moves down.** | config | @430 (assumed) | @ measured peak | @ measured 32-core | | --- | ---: | ---: | ---: | | **all bf8 — what goal1 actually runs** | 31.59 | **27.81** | **24.32** | | bf8, down_proj bf4 | 35.28 | 29.83 | 26.09 | | bf8, all MLP bf4 | 46.06 | 34.89 | 30.51 | **The goal1 target-only ceiling is 27.81 tok/s at peak and 24.32 on decode's own grid — below the 30 tok/s target with zero fixed cost and a perfect kernel.** Even all-MLP-bf4, which we don't have and which costs accuracy, reaches 34.89 / 30.51. So **speculation is no longer the optional half of the charter — it is load-bearing, with nothing behind it.** Lever A stays the first move only because speculation cannot pay until the step is weight-read-bound. This charter has now been corrected twice in one day by its own gates, both downward, both because an inherited third-party figure did not survive measurement. Commit `5b00938`. Still running: A1 GDN-fusion A/B (holds the lock now), and the DSpark acceptance-rate sweep whose live log is showing accepted=0 on 14 of its first 15 blocks at 4k/K=15.
A1 PASS: GDN fusion flags are 1.09-1.10x and goal1 was not setting them
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 4s
tt-stack-ci / Report upstream drift (pull_request) Successful in 8s
47f2b71f07
Measured A/B on the cards, 4k/8k/16k, both arms same image and profile:

  4k   14.128 -> 15.435 tok/s   71.37 -> 65.54 ms/step   1.093x
  8k   14.000 -> 15.333                                  1.095x
  16k  13.856 -> 15.258                                  1.101x

The flags (QWEN_GDN_FUSED_DECODE, CONV_GATES, FUSED_INPLACE, PACKED_QKV)
already existed and are enabled in the goal2 profile. goal1 -- the profile
every baseline number in this repo was measured in -- did not set them.
So this is a configuration gap we had been paying for, not new work, and
the gain grows with context.

The base arm reproduces the recorded baseline (14.13/14.00/13.86 against
13.85-14.00), which anchors the comparison.

One honest caveat, recorded because it will bite us later: the benchmark
generates "count from 1 to 300", so both arms and all three context rungs
produce an identical text_sha256. The byte-identity check therefore proves
only that fusion did not break trivial counting. Worse, the same degeneracy
would hand speculative decoding an unrepresentatively high acceptance rate
-- so MTP uplift must never be measured on this generation alone.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Member

A1 (GDN fusion) — PASS, and it was a configuration gap

Measured A/B on the cards, same image and profile, only env differs — bench/runs/a1-gdn-fusion-ab-20260911T184615Z.jsonl.

rung base fused speedup p50 ms
4k 14.128 15.435 1.093× 71.37 → 65.54
8k 14.000 15.333 1.095× 71.76 → 65.91
16k 13.856 15.258 1.101× 72.70 → 66.44

QWEN_GDN_FUSED_DECODE / CONV_GATES / FUSED_INPLACE / PACKED_QKV already exist and are enabled in the goal2 profile. goal1 — the profile every baseline number in this repo was measured in — does not set them. So this is a gap we had been paying for rather than new work, and the gain grows with context. The base arm reproduces the recorded baseline (14.13/14.00/13.86 vs 13.85–14.00), which anchors the comparison.

One caveat I want on the record because it will bite us later. The benchmark generates "count from 1 to 300", so both arms and all three rungs produce an identical text_sha256 (bafed7a9…). The byte-identity check therefore proves only that fusion did not break trivial counting — not numeric equivalence on hard tokens. Worse: that same degeneracy would hand speculative decoding an unrepresentatively high acceptance rate. MTP uplift must never be measured on this generation alone, and we need a prompt-dependent output arm before any Lever B number is believable.

Where Lever A now stands: 65.54 ms against a 35.95 ms measured-peak weight-read floor. A1 closed roughly a seventh of the gap; ~30 ms of non-weight-read cost remains. It does not change the Gate 0a part 2 conclusion. Commit 47f2b71.

## A1 (GDN fusion) — PASS, and it was a configuration gap Measured A/B on the cards, same image and profile, only env differs — `bench/runs/a1-gdn-fusion-ab-20260911T184615Z.jsonl`. | rung | base | fused | speedup | p50 ms | | --- | ---: | ---: | ---: | ---: | | 4k | 14.128 | **15.435** | 1.093× | 71.37 → 65.54 | | 8k | 14.000 | **15.333** | 1.095× | 71.76 → 65.91 | | 16k | 13.856 | **15.258** | 1.101× | 72.70 → 66.44 | `QWEN_GDN_FUSED_DECODE` / `CONV_GATES` / `FUSED_INPLACE` / `PACKED_QKV` already exist and are enabled in the **goal2** profile. **goal1 — the profile every baseline number in this repo was measured in — does not set them.** So this is a gap we had been paying for rather than new work, and the gain grows with context. The base arm reproduces the recorded baseline (14.13/14.00/13.86 vs 13.85–14.00), which anchors the comparison. **One caveat I want on the record because it will bite us later.** The benchmark generates *"count from 1 to 300"*, so both arms and all three rungs produce an identical `text_sha256` (`bafed7a9…`). The byte-identity check therefore proves only that fusion did not break trivial counting — not numeric equivalence on hard tokens. Worse: **that same degeneracy would hand speculative decoding an unrepresentatively high acceptance rate.** MTP uplift must never be measured on this generation alone, and we need a prompt-dependent output arm before any Lever B number is believable. Where Lever A now stands: 65.54 ms against a **35.95 ms** measured-peak weight-read floor. A1 closed roughly a seventh of the gap; ~30 ms of non-weight-read cost remains. It does not change the Gate 0a part 2 conclusion. Commit `47f2b71`.
Tenstorrent already built our Lever B: tt-metal PR #55548, verified
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 6s
f96c7a86c9
A scout agent reported this; a scout claim is not evidence, so every
structural claim below was checked against the GitHub API and the raw
files at the PR head before it was written down.

tenstorrent/tt-metal PR #55548, branch atupe/qwen36-mtp-v2, head
96f3f04102028cd85c23c2162914d5a0e7b37cc6, open, 53 files, +8294,
Apache-2.0, by a Tenstorrent engineer. Verified present: spec_decode.py
(+831), mtp.py (+248), gdn/tp.py (+753), the weight_mapping.py change
that unblocks exactly the mtp.* keys we found skipped at lines ~50/135/
425, a new C++ fused_recurrent_gated_delta_rule ttnn op with compute,
reader and writer kernels, and a spec_multi_pos mode in sdpa_decode.

The line that matters is spec_decode.py:53:

  assert model.num_devices > 1, "SpeculativeDecoder is TP-only for now"

It is TP-ONLY. Lottolabs' MTP raised NotImplementedError for
num_devices != 1 -- the opposite constraint, and the reason we wrote it
off. Ranks 1 (K>1) and 2 (TP>=2) of the scouting brief are satisfied by
one artifact, inside our own model's repo tree.

What does NOT transfer: the author's 50.45 tok/s at 4k and 2.62x at
ISL 128 are on QuietBox 2 with FOUR dies, where each card streams
~6.8 GB/token against our 13.61 at TP=2. Their number sits near their
roofline, which is about twice ours. The mechanism transfers; the
tok/s does not.

Two costs, both real. It needs tt-metal built from an unmerged branch,
which collides with "pins move as a set". And it is lossless only
conditionally: the PR states it matches greedy where plain decode is
confident (top-2 gap >= 2.0) and that near-ties can differ. Our charter
makes byte-identity a pass condition. That is a different standard and
we have to settle it deliberately rather than discover it later.

Consequence: Lever B is a port, not a research programme. With the
DSpark drafter measured dead today (mean accepted 0.1 at K=7) and no
target-only path reaching 30 tok/s, this is the only live route left.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
A1b: GDN bf16 state is worth another 3%, but it is not free yet
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 4s
tt-stack-ci / Report upstream drift (pull_request) Successful in 7s
b544bc700d
Swept the remaining goal3 flags, each isolated on top of the A1 fused
arm, 4k/8k/16k. Re-running the identical fused arm in a fresh session
differed by 0.33%, which finally bounds run-to-run noise on this bench
and is the yardstick for every delta:

  fused (A1)                    15.486  15.374  15.282   1.10x
  + QWEN_BATCHED_GROUPED=0      15.543  15.574  15.419   1.11x
  + GDN_DECODE_BF16/STATE_BF16  15.910  15.802  15.737   1.13x
  + both                        15.909  15.791  15.607   1.13x

GDN bf16 is +2.7-3.0%, consistent across all three rungs and about 9x
the noise floor, so it is real. BATCHED_GROUPED=0 is noise-order and is
NOT established. The two do not compose: adding BATCHED_GROUPED on top
of GDN bf16 changes nothing.

The 3% is not yet free, and this is the part worth recording. Those two
flags hold the GDN recurrent state in bf16 -- they change numerics. Every
arm shares one text_sha256, but that check is degenerate: the benchmark
generates "count from 1 to 300", which survives almost any numeric
perturbation. State precision is exactly what degrades on hard tokens and
long horizons rather than on counting. So it stays gated on a real
equivalence arm rather than being adopted on the strength of a sha that
cannot fail.

Where Lever A stands: 71.99 -> 63.27 ms, 1.126x, entirely from flags that
already existed and that goal1 was not setting. Configuration work is
close to exhausted; the remaining ~27 ms above the 35.95 ms measured-peak
floor is kernel work -- which is what PR #55548 already contains.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Member

The plan changed: Tenstorrent already built our Lever B

A scout agent reported this; a scout claim is not evidence, so I checked every structural claim against the GitHub API and the raw sources before writing it down. Record: bench/runs/scout-pr55548-verification-20260911T192500Z.jsonl.

tenstorrent/tt-metal PR #55548 — "[Feat]: MTP Spec Decode for Qwen3.8-27B for QB2", branch atupe/qwen36-mtp-v2, head 96f3f04102028cd85c23c2162914d5a0e7b37cc6, open and unmerged, 53 files, +8294, Apache-2.0, by a Tenstorrent engineer.

Verified present, not inferred: tt/spec_decode.py (+831), tt/mtp.py (+248), tt/gdn/tp.py (+753), the weight_mapping.py change that unblocks exactly the mtp.* keys we found skipped at lines ~50/135/425, a new C++ fused_recurrent_gated_delta_rule ttnn op with compute/reader/writer kernels, and a spec_multi_pos mode in sdpa_decode.

The decisive line, spec_decode.py:53:

assert model.num_devices > 1, "SpeculativeDecoder is TP-only for now"

It is TP-only. Lottolabs' MTP raised NotImplementedError for num_devices != 1 — the opposite constraint, and the reason we wrote that path off. Ranks 1 (K>1) and 2 (TP≥2) of the brief are satisfied by one artifact, inside our own model's repo tree.

What does not transfer. The author's 50.45 tok/s at 4k and 2.62× at ISL 128 are on QuietBox 2 with four dies, where each card streams ~6.8 GB/token against our 13.61 at TP=2. Their number sits near their roofline, which is about twice ours. The mechanism transfers; the tok/s does not, and I won't quote it as a target.

Two real costs. It needs tt-metal built from an unmerged branch — that collides with "pins move as a set" and is a packaging problem, not a modelling one. And it is lossless only conditionally: the PR says it matches greedy where plain decode is confident (top-2 gap ≥ 2.0) and that near-ties "can differ and are logged". Our charter makes byte-identity a pass condition. That is a different standard, and we should settle it deliberately rather than discover it in a benchmark.


A1b — Lever A configuration work is nearly exhausted

bench/runs/a1b-flag-sweep-20260911T193122Z.jsonl. Re-running the identical fused arm in a fresh session differed by 0.33%, which finally bounds noise on this bench.

arm (on top of fused) 4k 8k 16k vs goal1
fused (A1) 15.486 15.374 15.282 1.10×
+ BATCHED_GROUPED=0 15.543 15.574 15.419 1.11×
+ GDN_DECODE_BF16 + STATE_BF16 15.910 15.802 15.737 1.13×
+ both 15.909 15.791 15.607 1.13×

GDN bf16 is +2.7–3.0%, ~9× the noise floor, real. BATCHED_GROUPED=0 is noise-order and not established; the two don't compose.

That 3% is not yet free — those flags hold the recurrent state in bf16, i.e. they change numerics, and our sha check is degenerate (counting survives almost any perturbation). Gated on a real equivalence arm.

Lever A: 71.99 → 63.27 ms, 1.126×, entirely from flags that already existed and goal1 never set. The remaining ~27 ms above the 35.95 ms floor is kernel work — which is precisely what #55548 contains.

Commits f96c7a8, b544bc7.

## The plan changed: Tenstorrent already built our Lever B A scout agent reported this; a scout claim is not evidence, so I checked every structural claim against the GitHub API and the raw sources before writing it down. Record: `bench/runs/scout-pr55548-verification-20260911T192500Z.jsonl`. **`tenstorrent/tt-metal` PR #55548** — *"[Feat]: MTP Spec Decode for Qwen3.8-27B for QB2"*, branch `atupe/qwen36-mtp-v2`, head `96f3f04102028cd85c23c2162914d5a0e7b37cc6`, open and unmerged, **53 files, +8294**, Apache-2.0, by a Tenstorrent engineer. Verified present, not inferred: `tt/spec_decode.py` (+831), `tt/mtp.py` (+248), `tt/gdn/tp.py` (+753), the `weight_mapping.py` change that unblocks **exactly the `mtp.*` keys we found skipped at lines ~50/135/425**, a new C++ `fused_recurrent_gated_delta_rule` ttnn op with compute/reader/writer kernels, and a `spec_multi_pos` mode in `sdpa_decode`. The decisive line, `spec_decode.py:53`: ```python assert model.num_devices > 1, "SpeculativeDecoder is TP-only for now" ``` **It is TP-*only*.** Lottolabs' MTP raised `NotImplementedError` for `num_devices != 1` — the opposite constraint, and the reason we wrote that path off. Ranks 1 (K>1) and 2 (TP≥2) of the brief are satisfied by one artifact, inside our own model's repo tree. **What does not transfer.** The author's 50.45 tok/s at 4k and 2.62× at ISL 128 are on QuietBox 2 with **four dies**, where each card streams ~6.8 GB/token against our 13.61 at TP=2. Their number sits near *their* roofline, which is about twice ours. The mechanism transfers; the tok/s does not, and I won't quote it as a target. **Two real costs.** It needs tt-metal built from an unmerged branch — that collides with *"pins move as a set"* and is a packaging problem, not a modelling one. And it is lossless only *conditionally*: the PR says it matches greedy where plain decode is confident (top-2 gap ≥ 2.0) and that near-ties "can differ and are logged". **Our charter makes byte-identity a pass condition. That is a different standard**, and we should settle it deliberately rather than discover it in a benchmark. --- ## A1b — Lever A configuration work is nearly exhausted `bench/runs/a1b-flag-sweep-20260911T193122Z.jsonl`. Re-running the identical `fused` arm in a fresh session differed by **0.33%**, which finally bounds noise on this bench. | arm (on top of `fused`) | 4k | 8k | 16k | vs goal1 | | --- | ---: | ---: | ---: | ---: | | `fused` (A1) | 15.486 | 15.374 | 15.282 | 1.10× | | `+ BATCHED_GROUPED=0` | 15.543 | 15.574 | 15.419 | 1.11× | | **`+ GDN_DECODE_BF16 + STATE_BF16`** | **15.910** | **15.802** | **15.737** | **1.13×** | | `+ both` | 15.909 | 15.791 | 15.607 | 1.13× | GDN bf16 is +2.7–3.0%, ~9× the noise floor, real. `BATCHED_GROUPED=0` is noise-order and **not established**; the two don't compose. **That 3% is not yet free** — those flags hold the recurrent state in bf16, i.e. they change numerics, and our sha check is degenerate (counting survives almost any perturbation). Gated on a real equivalence arm. **Lever A: 71.99 → 63.27 ms, 1.126×, entirely from flags that already existed and `goal1` never set.** The remaining ~27 ms above the 35.95 ms floor is kernel work — which is precisely what #55548 contains. Commits `f96c7a8`, `b544bc7`.
The owner cleared taking any pin or branch and applying patches freely,
so this takes the shortest honest route to a measurement.

We did not compile tt-metal. cfx-llm2 has 4 cores and ~4 GB free RAM
once hugepages are reserved, so a from-source build there is not
realistic. The branch's own CI already publishes a matching wheel --
cp310 / manylinux_2_35 against the serving image's Ubuntu 22.04,
Python 3.10.19, glibc 2.35 -- so the image is assembled from that.

Pinned as a set at 96f3f04102028cd85c23c2162914d5a0e7b37cc6: both the
ttnn wheel and the qwen36 model tree. The branch head moved to f80137ca
while we were cloning it, which is exactly the mismatched-pair trap our
own rules warn about, so the tree is checked out at the wheel's sha
rather than at the branch tip.

Verified in the built image: ttnn.transformer.fused_recurrent_gated_-
delta_rule is present and sdpa_decode carries the spec argument.

One patch was required. text_demo.py maps MESH_DEVICE to a mesh shape
and has no entry for P300, so our 2-die rig silently fell back to their
4-die (1,4). Added P300 -> (1,2).

Two things recorded because they change the plan rather than decorate
it:

  1. This image DOWNGRADES ttnn 0.79.0.dev+cfx -> 0.75.0rc10. The
     vllm-tt plugin is built against 0.79 and is not expected to work
     here. vllm-tt:k2 is untouched and remains the serving image.

  2. tt/qwen36_vllm.py in this PR has NO speculative-decode wiring.
     MTP lives in spec_decode.py driven by text_demo.py. Measuring the
     uplift is a wheel swap; shipping it through our server is a real
     integration. The scout's "cherry-pick and adapt" framing understated
     this, and it is better to know now than after committing to a date.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Lever A2: grid shape is spent, and the ornith lead is a dead end
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 6s
tt-stack-ci / Report upstream drift (pull_request) Successful in 9s
87c6c9c52e
Two investigations, both of which close a line of attack rather than
open one. Recording them because a refuted lead saves more time than a
confirmed one.

1. DECODE MATMUL GRID AUDIT (my own lead, and it was wrong)

I read "Attention grid 8x4 / MLP grid 8x2 / LM head grid 8x5" out of the
model's config log and inferred the MLP -- the largest weight consumer --
was decoding on 16 cores against a 44-core bandwidth knee. That banner
comes from the base tt_transformers.ModelArgs describing its OWN configs.
The qwen36 TP decode path never consumes those grids. It builds explicit
MatmulMultiCoreReuseMultiCast1D configs, verified at every consumer:

  mlp gate/up        11x4 = 44 cores   at the knee
  gdn in-proj        11x4 = 44 cores   at the knee
  mlp down/attn wo   11x3 = 33 cores
  lm_head            11x10 (ttnn-auto) at the knee
  attn fused qkv     8x8  = 64 cores   <- the only sub-knee grid

attn QKV on 8x8 is the one config never re-shaped wide-first, and 8x8 is
the shape our own probe measured SLOWER than 11x4 (321.6 vs 372.0 GB/s).
It is worth 0.26 ms. Resolving every remaining uncertainty in the lead's
favour, total grid-shape headroom is <=1.5 ms of a 63.27 ms step (2.3%),
and the 0.26 ms that rests on measured rather than interpolated bandwidth
is below our 0.33% noise floor.

The byte accounting reproduces Gate 0a part 1 exactly (13.612 vs 13.61
GB/card), which anchors it.

So grid tuning is spent. The finding that matters is where the time is
NOT: ~26 ms of the step is not weight read at all, and no grid change
touches it -- SDPA decode over 16 layers, GDN recurrence over 48, 64 CCL
all-reduce/all-gathers, norms, trace launch, host readback. A per-kernel
device profile is now the prerequisite for any further Lever A work; we
have never itemised that 26 ms.

2. ORNITH-1.0-35B (the scout's rank-4 find): all four fixes real, zero
   transferable.

Source-verified at rev 64b75ab499c66e31db547ef70263aa4f084ec9a5. The
4-core starvation, layout thrash, depthwise conv and head-merge fixes are
genuine -- but they repair defects in Ornith's own HuggingFace autoport
layer, not TTNN itself. Qwen36 TP=2 already runs the dedicated C++
decode_gated_delta_rule op, already uses ttnn.mac for the decode conv,
and never calls nlp_concat_heads in decode. The claimed 2.613 -> 2.063 ms
ships with no run logs, traces or audit script, and describes a 35B MoE
layer on a 1x1 mesh.

Added but default OFF: QWEN36_1D_GRID_<NAME>=CxR overrides for every
decode matmul (HiFi2 pinned on the lm_head so a program config cannot
silently drop ttnn-auto's fidelity to LoFi) plus a one-shot log of the
grids actually consumed, and a bind-mount A/B driver that builds nothing
on cfx-llm2 and carries a prompt-dependent equivalence rung, because the
counting sha cannot detect numerics.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Trap: force-killing a tt-metal container wedges the cards
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 3s
tt-stack-ci / Report upstream drift (pull_request) Successful in 7s
495265ac6a
Recording this because it cost hours and it was self-inflicted.

I force-killed (docker rm -f) a hung 9-hour text_demo run to free the
hardware lock. Every subsequent device open failed with:

  Timeout (10000 ms) waiting for physical cores to finish: 13-6, 11-7, ...
  Device 0 init: failed to initialize FW! Try resetting the board.
  TT_THROW @ risc_firmware_initializer.cpp:1573

The trap is that this presents as a VERSION problem. The run that failed
first was the PR #55548 image (ttnn 0.75 against a 19.14.0 firmware
bundle), so the obvious reading was that the CI wheel's base was too old
for our cards -- a conclusion that would have sent us off rebasing the
PR onto a newer tt-metal for no reason. Testing our own known-good
vllm-tt:k2 (ttnn 0.79) disproved it: same failure, same line. The cards
were wedged, not mismatched.

`sudo tt-smi -r` fixed it. That is a board reset, not a firmware flash,
so it does not touch the "firmware updates are never unattended" rule.

Second trap recorded alongside it: immediately after the reset, the mesh
open failed with "Sysmem mapped at unexpected NOC address ... another
process is already holding the sysmem NOC address space". That was not a
failed reset -- it was our own queued grid A/B taking the cards the
moment they came back. Check docker ps before re-resetting.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Build tt-metal main + PR #55548 locally; stop assembling images from parts
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 7s
tt-stack-ci / Report upstream drift (pull_request) Successful in 9s
370bb7bea6
Owner challenged the pin as "an older tt-metal version". Checked, and the
version strings are the problem, not the pin:

  our old flake pin 1227e182   2026-08-12
  PR branch head    96f3f041   2026-09-11   <- what we bumped to
  current main      60053f1f   2026-09-12

The bump moved us a MONTH FORWARD. Three version strings disagree with
recency and all three misled: "0.77.0-dev20260813" is hand-written in our
own flake.nix and stale; "0.75.0rc10" is the PR's CI release-branch
numbering; and the container's "0.79.0.dev20260903+cfx" is a cfx FORK
dated 09-03, i.e. older than the PR despite the bigger number. Annotated
the flake string so the next reader is not misled the same way.

But the owner's underlying point was right and I had settled for less.
The PR branch is 80 commits behind main, so building it as-is forfeits
upstream work and forecloses stacking other open PRs. The merge onto
current main is CONFLICT-FREE (53 files), so we now build
main@60053f1f + #55548 and the flake points at that local tree. SFPI
followed the source automatically to 7.76.0, which is the self-consistency
we were missing.

Also recorded why the CI-wheel shortcut failed, because it cost hours and
the failure mode is genuinely deceptive. Assembling an image mixes four
independently versioned things -- the wheel's headers, its .so, the
image's SFPI, and the image's TT_METAL_HOME tree -- and each mismatch
only surfaces after a ~25 minute weight load. We fixed three (missing
experimental module, SFPI 7.73 vs 7.75.1, TT_METAL_HOME) and still hit
"Index out of range (3130 < 7)". Build the tree; do not assemble one.

Also lands the A2 grid A/B records from the subagent run.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Member

Pushed: everything through the local tt-metal build (370bb7b)

Correction I owe the record: the pin was never "older"

commit date
old flake pin 1227e182 2026-08-12
PR #55548 head 96f3f041 2026-09-11
current main 60053f1f 2026-09-12

The bump moved us a month forward. Three version strings disagree with recency and all three mislead: 0.77.0-dev20260813 is hand-written in our own flake.nix and stale; 0.75.0rc10 is the PR's CI release-branch numbering; and the container's 0.79.0.dev20260903+cfx is a cfx fork dated 09-03 — older than the PR despite the larger number. I've annotated the flake string so nobody repeats the confusion. Compare commit dates, never version numbers.

But we now build main + the PR, not the PR alone

The PR branch is 80 commits behind main. The merge onto current main is conflict-free (53 files), so we take both rather than forfeiting upstream work — and it leaves room to stack further open PRs. SFPI followed the source to 7.76.0 automatically, which is the self-consistency the hand-assembled image never had.

Why the CI-wheel shortcut failed (hours, and a deceptive failure mode)

Assembling an image mixes four independently versioned things — the wheel's headers, its .so, the image's SFPI, and the image's TT_METAL_HOME tree — and each mismatch only surfaces after a ~25 minute weight load:

  1. missing models/experimental/gated_attention_gated_deltanet (the PR patches it too);
  2. SFPI 7.73.0 vs pinned 7.75.1 — appears as a brisc compile failure deep in tensor_accessor_args.h, never as a version error;
  3. TT_METAL_HOME on the 0.79 tree while headers came from the 0.75 wheel;
  4. after all three: static assertion failed: Index out of range … (3130 < 7).

Written up in docs/MTP-BUILD-WORKFLOW.md. Build the tree; do not assemble one.

Two self-inflicted traps now in CLAUDE.md

  • Force-killing or docker stop-ing a live tt-metal container wedges the cards. Next device open fails "failed to initialize FW! Try resetting the board" — and it presents as a version problem. Our known-good 0.79 image failed identically, which is what disproved it. sudo tt-smi -r fixes it (a reset, not a flash).
  • tt-smi -r returns before the boards are ready; launching immediately gives Read 0xffffffff over PCIe. Verify with a mesh-open probe first.

Measurements landed

A1 GDN fusion flags 14.13 → 15.44 tok/s (goal1 never set them)
A1b + GDN bf16 state → 15.91 tok/s / 63.27 ms, gated on a real equivalence arm
A2 grid audit grid tuning spent: ≤1.5 ms of 63 ms
A2 grid A/B attn QKV 8x8→11x4 ≈ 1 ms, single pair — suggestive, not established
Gate 0a pt2 bf8 peak 378.6 GB/s, bf4 287.9 (24% penalty)
DSpark acceptance 0.1 at K=7 — dead
ornith claims all four real, zero transfer

~26 ms of the 63 ms step is still unexplained, and that is where the remaining headroom lives.

## Pushed: everything through the local tt-metal build (`370bb7b`) ### Correction I owe the record: the pin was never "older" | | commit date | | --- | --- | | old flake pin `1227e182` | 2026-08-12 | | PR #55548 head `96f3f041` | **2026-09-11** | | current main `60053f1f` | 2026-09-12 | The bump moved us **a month forward**. Three version strings disagree with recency and all three mislead: `0.77.0-dev20260813` is **hand-written in our own flake.nix** and stale; `0.75.0rc10` is the PR's CI release-branch numbering; and the container's `0.79.0.dev20260903+cfx` is a **cfx fork** dated 09-03 — *older* than the PR despite the larger number. I've annotated the flake string so nobody repeats the confusion. **Compare commit dates, never version numbers.** ### But we now build main + the PR, not the PR alone The PR branch is **80 commits behind main**. The merge onto current main is **conflict-free (53 files)**, so we take both rather than forfeiting upstream work — and it leaves room to stack further open PRs. SFPI followed the source to **7.76.0** automatically, which is the self-consistency the hand-assembled image never had. ### Why the CI-wheel shortcut failed (hours, and a deceptive failure mode) Assembling an image mixes four independently versioned things — the wheel's headers, its `.so`, the image's SFPI, and the image's `TT_METAL_HOME` tree — and **each mismatch only surfaces after a ~25 minute weight load**: 1. missing `models/experimental/gated_attention_gated_deltanet` (the PR patches it too); 2. SFPI **7.73.0 vs pinned 7.75.1** — appears as a `brisc` compile failure deep in `tensor_accessor_args.h`, never as a version error; 3. `TT_METAL_HOME` on the 0.79 tree while headers came from the 0.75 wheel; 4. after all three: `static assertion failed: Index out of range … (3130 < 7)`. Written up in `docs/MTP-BUILD-WORKFLOW.md`. **Build the tree; do not assemble one.** ### Two self-inflicted traps now in CLAUDE.md - **Force-killing *or* `docker stop`-ing a live tt-metal container wedges the cards.** Next device open fails *"failed to initialize FW! Try resetting the board"* — and it presents as a *version* problem. Our known-good 0.79 image failed identically, which is what disproved it. `sudo tt-smi -r` fixes it (a reset, not a flash). - **`tt-smi -r` returns before the boards are ready**; launching immediately gives `Read 0xffffffff over PCIe`. Verify with a mesh-open probe first. ### Measurements landed | | | | --- | --- | | A1 GDN fusion flags | 14.13 → 15.44 tok/s (goal1 never set them) | | A1b + GDN bf16 state | → **15.91 tok/s / 63.27 ms**, gated on a real equivalence arm | | A2 grid audit | grid tuning **spent**: ≤1.5 ms of 63 ms | | A2 grid A/B | attn QKV 8x8→11x4 ≈ **1 ms**, single pair — suggestive, not established | | Gate 0a pt2 | bf8 peak **378.6 GB/s**, bf4 **287.9** (24% penalty) | | DSpark acceptance | **0.1** at K=7 — dead | | ornith claims | all four real, **zero transfer** | **~26 ms of the 63 ms step is still unexplained**, and that is where the remaining headroom lives.
Grok added 49 commits 2026-09-13 10:24:05 +02:00
We never had per-op device timings because ENABLE_TRACY=OFF, and turning it on
failed one network fetch at a time in the Nix sandbox: first emsdk (gated in
our tt-metal tree), then ImGui (tracy/cmake/vendor.cmake pulls the whole GUI
viewer dependency set at configure time even for the CLI tools), then
PPQSort's own CPM downloader stub and its PackageProject.cmake fetch.

The tracy-capture/tracy-csvexport CLI only links TracyServer, which needs
capstone + zstd + PPQSort. patches/tracy-vendor-cli-only.patch adds
TRACY_VENDOR_CLI_ONLY to vendor.cmake (carried as a patch because tracy is a
submodule whose pointer must stay fetchable); the merged tree sets it ON; the
three deps plus PackageProject.cmake are seeded like every other CPM package.

bench/profile/: a repeatable per-op profile of traced decode steps using the
in-process C++ post-processor (TT_METAL_PROFILER_CPP_POST_PROCESS), no GUI.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
import ttnn._ttnn initialises the installed tracy python package, whose
process_ops_logs/process_device_log import click and seaborn at module scope.
pythonImportsCheck caught it; the C++ build itself was already green.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
We believed the all-MLP-bf4 ceiling of 34.89 tok/s followed from our measured
bf4 bandwidth. It does not: 9.33 GB at bf4's measured 287.9 GB/s is 30.86 tok/s,
clearing the 30 target by only 0.93 ms. Our 34.89 is a MIXED-rate figure that
assumes the untouched bf8 stream still runs at 378.6 and only the bf4 MLP portion
at 287.9. Recomputing that path reproduces 34.90, so the number stands -- but the
"at_measured_peak" label overstates the headroom by 4 tok/s to anyone reading it
naively, which included us. Both figures are now recorded as the pessimistic and
optimistic bf4 budgets.

Two further corrections to our own material. "63.27 ms / 15.91 tok/s" is
internally inconsistent (1000/63.27 = 15.805); 15.91 is a best-arm value paired
with a different-window ms, and at our measured 0.33% noise floor a 0.7%
discrepancy is not noise. And the brief's claim that the 48 GDN layers' state
updates are independent within a step was wrong: storage is layer-local, but each
layer's update inputs come from the previous layer's output, so the chain is
serial and cross-layer batching is not available. The parallelism is intra-layer.

The advisory is third-party analysis from a model with no access to our hardware
that reproduced nothing, so per the distinguish-measured-from-third-party rule it
is filed under docs/advisory/ behind a provenance header and is not evidence.
Only the corrections we recomputed ourselves went into bench/runs/.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Believed: prefill was an unmeasured adjunct of decode, with TTFTs only as a
by-product of decode ladders, and the obvious suspect was chunk size or
bandwidth by analogy with decode.

Measured (bench/runs/prefill-audit-20260912T100217Z.jsonl): the goal1 server
prefills at ~3,500 tok/s from 4k to 16k, linear in ISL (attention is not the
driver). A matmul-rate probe on the exact per-card TP=2 prefill shapes shows
the served program configs sum to 384 ms of matmul per 2048 chunk (130
TFLOP/s effective) where full-grid configs at the same fidelity do 179 ms
(279 TFLOP/s; 368 peak). The chunk is 569 ms, so two thirds is matmul time
running at under half the rate the silicon delivers. Bandwidth crossover is
~400 tokens/chunk; at 2048 the weight read is 6%, so chunk size is not a
lever. Separately, the SEQS=1 masked-bucket path is eager and costs ~0.42 s
regardless of length: the TTFT floor for prompts <= 2k and a tail tax on
every longer one.

Ranked gaps with per-chunk cost: eager bucket 0.42 s; GDN out-proj MMRS on
8x6 with 1x1 subblocks 71 ms; MLP down on the TP=4 fallback row 52 ms; MLP
gate|up AGMM on 8x9 29-40 ms; GDN in-proj AGMM 25 ms. An A/B of the
down-proj fix is queued behind the decode agent's lock and is not claimed.

Cards were found wedged at session start (no process held them); tt-smi -r
was run under the lock and verified before measuring.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Believed (from bench/prefill-matmul-peak.py): routing the prefill MLP
down-proj through minimal_matmul on the full 11x10 grid would cut ~52 ms
from a 569 ms chunk at the same LoFi + fp32-dest-acc fidelity.

Measured (interleaved C T C T, fresh server per arm, server-side stamps,
prompt-dependent 48-token output byte-identical across arms): +1.5..2.3% at
8k and 16k, +1.6% on the 2048 bucket, within noise at 4k. The standalone
probe cannot see the in-model memory traffic (L1-resident output for the
all-reduce, neighbouring fused CCL ops). tp_common.py's warning that a
per-op sweep cannot see this is now a measurement. The program-config gaps
in the audit are therefore upper bounds needing per-item in-model A/Bs; the
next diagnostic is a device-side op profile of one chunk.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
We had been treating the 37.4 ms weight read as something grid and scheduling
work might still improve. It cannot: 13.61 GB / 37.4 ms is 363.9 GB/s, which is
96% of the measured 378.6 GB/s bf8 peak. Reading faster is over as a lever. Only
reading fewer bytes -- quantisation -- can touch that 37.4 ms, which is why the
bf4 question now gates a whole rung rather than being an escape hatch.

Writing the ladder out end to end with measured inputs gives ~48-50 ms and
~20-21 tok/s for a well-executed Lever A. That does not reach 30. Speculation was
already mandatory per precedence line 0; the ladder now says how much is required
of it (~1.45-1.5x), and why closing the residual first raises what it is worth:
the step goes from 59% to ~75% weight-read, and verifying K+1 positions shares
one weight read.

Also records what not to retry: cross-layer all-reduce deferral is generally
incorrect rather than unimplemented, and the 48 GDN state updates are not
independent within a step, so the parallelism to chase is intra-layer.

The 25.87 ms residual stays labelled a subtraction, not a measurement, until the
P0 per-op profile exists. Every phase gate is a bench/runs record.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
This morning I wrote into the plan and into CLAUDE.md that the weight read runs
at 363.9 GB/s, 96% of the measured 378.6 GB/s peak, and concluded that reading
faster was finished as a lever and only quantisation could touch the 37.4 ms.
That conclusion was sound given its premise and the premise was wrong: 378.6 was
the fastest OUR probe could read, not the fastest the part can read.

A controlled bank-local reader reaches 449.2 GB/s -- above even the bf16 number
we had been calling peak. The variables are not dtype and not grid shape. They
are request length (576 B -> 162.2 GB/s, 1088 B -> 304.5, 2304 B -> 399.0),
outstanding depth (329.6 / 401.0 / 449.2 at depth 16 / 32 / 64) and NOC plane
(noc0 beats noc1 by 1.61x on identical work). With length and depth controlled,
the 8x8-vs-11x4 gap that part 2 reported as 13.5% is 0.2%.

So bf4's "24% bandwidth penalty" is not silicon. A packed bf4 tile is a 576-byte
request and this regime is request-rate limited, not byte limited: 162.2/304.5 =
0.533 tracks 576/1088 = 0.529 almost exactly. Coalescing tiles per DMA should
recover it, which moves all-MLP-bf4 off its pessimistic 0.93 ms margin.

Two measurements also demote work we had queued. All 64 collectives cost ~1.6 ms,
not the 6.5 ms the advisory prior guessed -- all_gather 25.3 us, reduce_scatter
26.1 us -- and the advisory had itself named 25 us as the falsifying value. GDN
fusion is worth 14-19% of the recurrence in a synthetic probe.

The clock assumption was the thing that could have invalidated all of it. The
probe assumes 1350 MHz while tt-smi reports 800; the 800 is the idle reading,
telemetry after each run reports 1350 on both cards, and AICLK_BUSY_VAL is 1350
in blackhole_implementation.hpp. Checked before accepting any figure.

None of this is banked. It is a synthetic reader, not a matmul weight read, and
we have a counter-example from the prefill audit where a standalone 2.35x matmul
win came out 0-2% SLOWER in the model. What the model's readers actually expose
is now the open question, and the gate for it is an in-model A/B.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
We believed the 63.27 ms step was 37.4 ms of weight read plus a 25.87 ms
unexplained residual. That 25.87 ms was a subtraction from a MODEL, never a
measurement, and this is the first time we have looked.

The measurement says the subtraction is wrong in the direction that matters:
measured matmul kernel time on one card is 31.14 ms instrumented / 29.23 ms
deflated -- BELOW the 37.4 ms the weight-read model claims. Subtracting a
model that exceeds the measured matmul span inflates the residual.

What is actually there, per 4k decode forward, one card:
  matmul (weight projections)     464 ops   31.14 ms   61.5%
  layout / materialisation      2,528 ops    6.73 ms   13.3%
  norm + elementwise            1,376 ops    5.29 ms   10.4%
  CCL collectives                 256 ops    4.34 ms    8.6%
  lm_head + final norm              4 ops    1.79 ms    3.5%
  SDPA + KV maintenance           144 ops    1.37 ms    2.7%
  unattributed / overlap reconciliation      8.04 ms   (vs the control step)

Two of the advisory's priors are wrong by ~3x and it is worth saying so: SDPA
+KV is 1.37 ms, not 4.5; the collectives are 17 us each, not 102. GDN is 75%
of device time but 23.6 ms of its 37.8 ms is its own matmuls; its non-matmul
14.26 ms over 48 layers is glue -- 624 ReshapeView, 672 BinaryNg, 384 Copy,
384 Slice per step -- not a recurrence kernel. That reorders P1.

Two instrument findings the next agent would otherwise pay for again:

- The traced replay CANNOT be profiled per-op. The kernel profiler's DRAM
  marker buffer is 12000 entries/RISC at COMPILE time, a decode step is 4,776
  programs, and it overflows mid-replay: 47.1% of ops recorded, every worker
  core logging "markers were dropped". Its 34.89 ms span is a truncated
  prefix, not the step. Hence --eager-per-layer, which drains the profiler
  between layers.
- The report's OP NAME column is empty for trace replay, so names come from a
  ttnn graph capture of the same eager call, joined per layer and rejected
  unless the device-row count equals the graph op count exactly.

Chunked prefill does not run at all on the Tracy build (the GDN conv1d hits
"Statically allocated circular buffers ... clash with L1 buffers" at every
chunk size), so decode starts at position 4096 with zeroed KV and GDN state.
Shapes, grids and the program graph are identical to a real 4k step; the
values are meaningless and the record says so. This harness's control step is
55.586 ms, NOT the 63.27 ms production arm -- percentages transfer, absolute
milliseconds do not.

Evidence: bench/runs/p0-decode-per-op-profile-20260912T151000Z.jsonl
Re-run: bench/profile/p0-decode-profile.sh, then bench/profile/itemise.py.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The P0 record concluded the 25.87 ms residual was overstated, on the grounds that
measured matmul (31.14 ms) is below the 37.4 ms modelled weight read. That
compares 31.14 ms on this harness's 55.586 ms control step against 37.4 ms on the
63.27 ms production arm — the cross-harness transplant the record's own caveat
forbids. Scaled properly matmul is 33.27 ms, and since matmul time includes
compute as well as reads, weight read is at most that, so the non-weight-read
work is at least 30.00 ms. The residual was an UNDERestimate; the sign is
inverted.

What is true, and better than either version: of that ~30 ms only 9.15 ms is
still unattributed. The residual stopped being a mystery rather than shrinking.

It also means the 37.4 ms weight-read model exceeds total measured matmul kernel
time and so overstates reality. Every ceiling quoted from it needs re-deriving —
that model has now been wrong in both directions in one day, having also been
measured against a bandwidth peak that was our own reader's limit.

The finding that actually redirects the work: the step is 4,776 device programs
and 2,528 of them, 53% of all ops, are pure layout/materialisation, costing 7.19
ms scaled and sitting almost entirely inside GDN. Nobody had this on a list. GDN
is 75% of the step but most of that is its own matmuls; its non-matmul 297
us/layer is glue — ReshapeView, BinaryNg, Copy, Slice, Typecast — not a
recurrence kernel, so P1 targets op-graph consolidation and not a fused op.

SDPA is demoted: 1.46 ms scaled across all 16 layers, against a 4.5 ms prior.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Every ceiling we have quoted came from a model, and P0 plus Gate 0a part 3
measured both of those models wrong in opposite directions on the same day: the
378.6 GB/s peak was our own reader's limit, and the 37.4 ms weight-read model
exceeds total measured matmul kernel time. So Goal N is stated as a measured step
time on the production arm rather than as a ceiling, and is sized so success and
failure are both unambiguous.

The arithmetic that sets the order of work: bf4-MLP is worth 10.46 ms if matmul
is weight-read-bound and its tiles are coalesced into large DMAs, 0.15 ms if it
is read-bound but left at native 576-byte requests, and about nothing if matmul
is compute-bound. Two knife-edges, ten milliseconds, and neither the read/compute
split nor the model's actual DMA request size has ever been measured. N1 and N2
measure them and block the rest.

Without the bf4 rung the remaining work sums to ~6.6 ms and lands at 17.6 tok/s,
so N4 is the goal and N1/N2 decide whether it exists.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Counting forward_decode by hand gives 5 Copy, 11 BinaryNg and 3 Slice per layer
against a measured 8, 14 and 8. The gaps are left open deliberately: they should
come from _project_qkvzab/_row_proj/_col_proj, and any audit that reconciles to
the measured table without naming where they come from has fitted the answer.

Recorded because the first pass assumed K=7, which makes Copy and BinaryNg
reconcile to the measured counts exactly. Both matched. The checkpoint says K=4,
so that was numerology, and without checking the config it would have become a
finding.

The candidate this turns up: the conv window is a shift register built out of
tensor copies -- 4 Copy per layer per token, 192 device programs per step, of
which 144 move data that does not change. A rotating head index would remove
them, if the tap loop, the verify-path mirrors, and above all the captured trace
can tolerate a per-step rotation. The trace constraint may well kill it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
_project_qkvzab issues five slices on the decode path (qkv, z, ab, a, b); with
the three q/k/v slices in forward_decode that is exactly the measured 8 per
layer. Closed by reading the code, not by adjusting the count.

It also removes Slice as an N3 target. The ab double-slice looks redundant and
is not: slicing a and b straight out of qkvzab untilizes the full 4120-wide
tensor, so someone already paid one extra Slice to avoid that. Removing it
regresses.

_row_proj and _col_proj turn out to be pure matmul dispatch with no glue, so the
outstanding 3 Copy and 3 BinaryNg must be inside tpc.matmul_1d_decode /
sharded_decode_matmul or ttnn.rms_norm.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
We believed the decode matmul's read/compute split was unknown and that a
coalesced bf4 MLP was worth 10.46 ms. Both halves of that are now measured and
both are wrong, in opposite directions.

N1. Swept the model's own seven TP=2 decode projections over dtype (moves bytes)
crossed with MathFidelity (moves FLOPs). 70-76% of matmul time is byte-
proportional DRAM weight reading. The rest is a dtype-independent per-matmul
floor that is NOT arithmetic: quadrupling the math passes costs 1.3-5.8%. The
"matmul compute-bound" scenario is falsified. The seven shapes reconstruct
30.38 ms standalone against P0's 33.27 ms in-model — 91%, from two independent
harnesses.

The 37.4 ms weight-read model is retired; the replacement is ~23.4 ms. It was
wrong twice. Wrong dtype: it assumed all-bf8 at 13.61 GB/card, but tt/mlp.py
loads gate_proj and up_proj as bfloat4_b unconditionally — confirmed in the TP
tensor cache on disk — so production traffic is 10.76 GB/card/token. Wrong
bandwidth: it divided by 363.9 GB/s when the readers achieve 454-469 GB/s
marginal, above even part 3's controlled synthetic reader. The "7 ms of
weight-read headroom" opened yesterday was an artefact of the retired model and
does not exist.

N2 is a negative result and closes the rung. The matmul in1 readers issue
exactly one tile per DMA — 576 B at bf4 — in the shipped kernel source and in
the measured page sizes. Nothing in TTNN changes it: outstanding depth, core
count, grid shape and DRAM-width-sharded layout were each swept on a real model
shape at both dtypes, production is already the optimum of every one, and bf4
saturates near 285 GB/s from 22 cores to 110 while bf8 reaches 382. On
interleaved weights adjacent tiles live in different banks, so there is nothing
to coalesce in the first place. A real multi-tile burst does exist behind
#ifdef SPLIT_DRAM_BANK in the dram-sharded reader, but it has no Python surface
and sits on a path measured 2-3x slower here: it is a tt-metal patch, not a knob.

Consequence: N4 as specified does not exist. Two thirds of the bf4-MLP rung was
already spent before Goal N was written; only down_proj remains, at 2.285 ms
standalone (a ceiling, not a saving). The ladder now sums to ~54.4 ms / 18.4
tok/s with every conditional granted in full, so 50 ms is not reachable from the
rungs currently listed. N1 says where the open ground is instead: the ~8.2 ms
dtype-independent per-matmul floor and the 4,776 programs per step.

One item opened rather than closed: attn_qkv is the only projection still on the
legacy grid_w=8 shaping and reads at 364.5 GB/s against 454-469 for the 11-wide
grids — ~0.3 ms standalone, one line in model_config.py, untested in-model.

Also fixes a lock leak in the probe runner: `exec`ing the python replaced the
shell and destroyed the EXIT trap, stranding /tmp/ttlock.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
tt/mlp.py:149 says "gate/up: bfloat4_b (bandwidth); down: bfloat8_b (accuracy)".
gate_proj and up_proj have been bf4 since before this effort started, confirmed
on disk in the TP tensor cache. Gate 0a part 1's "all-bf8, 13.61 GB/card/token"
was wrong the day it was written, and every ceiling we have quoted since --
GOAL-30TPS, precedence line 0, the first draft of Goal N -- inherited it.

Measured traffic is 10.76 GB/card/token and the readers achieve 454-469 GB/s,
above part 3's 449.2 synthetic ceiling, so part 3's reader was ~1.75x pessimistic
about the real weight path and the 7 ms of headroom it appeared to open does not
exist. The weight read is 23.34 ms, 37% of the step, and essentially optimal.
Weight-read-only ceiling: 42.8 tok/s. Bandwidth was never what stood between us
and 30 tok/s.

That inverts the strategy. 63% of the step is not weight read, the itemisation
closes to the decimal at 39.93 ms, and the binding constraint is per-op cost and
op count -- 4,776 device programs per step -- not bytes. The largest single block
is now 17.4 ms of dispatch-shaped cost: an 8.23 ms per-matmul fixed cost that is
neither reads nor arithmetic (quadrupling math passes costs 1.3-5.8%), plus 9.15
ms unattributed. What the 8.23 ms is made of is the next gate.

The bf4 rung is retired rather than deferred: two of three MLP tensors are
already bf4, down_proj is bf8 deliberately for accuracy, and N2 measured that
readers issue exactly one tile per DMA with adjacent tiles in different banks, so
there is nothing to coalesce. Width-sharding is 2-3x slower and the real
multi-tile burst is behind an ifdef with no Python surface.

Goal N stays at 50 ms / 20 tok/s, now a 33% cut of non-read work rather than a
quantisation rung. Speculation is also weaker than precedence line 0 assumed: it
pays when the step is weight-read-bound, and 37% is not that.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
We had MTP filed as an integration project. It is not. demo/text_demo.py already
drives spec_decode/mtp/spec_sampling, and QWEN36_SPEC defaults ON -- the
documented knob is the opt-OUT. Greedy K defaults to 11 at <= 4k.

Both asserts that gate it are satisfied. spec_decode.py:53 requires TP, which
disqualifies single-device setups and not us at TP=2. spec_decode.py:52 requires
an MTP head, and our checkpoint has one: 22 mtp.* tensors of 1606, an fc plus a
complete single transformer layer with FP8 scales.

So every decode number we have -- the goal1 baseline, the a1b sweep, the a2 grid
work, the P0 profile -- was measured with QWEN36_SPEC=0, and the cost of finding
out what MTP does is one environment variable rather than an integration.

Not a performance claim: the draft head's per-token cost and the per-draft-position
lm_head are unaccounted for and could eat the shared-weight-read saving, which is
thinner than we assumed now that the step is 37% weight read rather than 59%.
Acceptance rate is the dominant unknown, and our degenerate counting benchmark
would report a fake-high one, so it cannot be used to measure it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Writing the measured numbers out as rooflines makes the picture unambiguous. The
weight read is 23.34 ms and already at ~100% of achievable bandwidth, so a pure
DRAM roofline would be 42.84 tok/s and we are running at 37% of it. Arithmetic is
1.7 ms. Neither is the constraint.

The constraint is that a token costs 4,776 device programs at 13.2 us each. The
4,312 non-matmul ops alone are 30.00 ms, 47% of the step, at 4.8 us of kernel
time plus gaps for ops that move almost no data -- a ReshapeView on a [1, 5120]
activation moves nothing but still dispatches across a 100+ core grid.

The number that reframes the goal: matmuls alone, with every glue op free, is
33.27 ms = 30.06 tok/s. The owner's target is inside the work we already do,
being spent on overhead. That is a bounding case and not a plan, but it sets the
scale of what op-count work is worth, and it is why Goal N sits at 50 ms.

This is the batch-1 latency regime: published Blackhole figures are batch-32
throughput numbers where one launch cost amortises over 32 rows of real work. At
batch 1 there is nothing to amortise against, so the cost structure inverts. The
machine is not inefficient at what it was built for; we are asking for something
else.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
What was believed (docs/GOAL-N-SERIES.md, this morning): 63% of the 63.27 ms
decode step is device-side per-op cost -- 4,776 programs, 2,528 layout ops,
8.23 ms of per-matmul floor, 9.15 ms of "dispatch / gaps" -- and the largest
addressable item is GDN layout glue, with the conv shift-register's 4 Copy per
layer as the live candidate.

What the evidence changed:

1. The P0 profile ran the COMPOSED GDN path, not production. ttm-src has no
   conv_gates / packed / in-place code and the Tracy ttnn lacks the ops; the
   3 Ternary (mac) per GDN layer in the table cannot occur on the served
   path. Production is ~30 ops/GDN layer since K1/K2 (2026-09-05). The
   shift-register candidate does not exist there. Correction appended to
   the P0 record; docs/n3 marked void for production.

2. The 9.15 ms was host time. process_output_decode reads the decode logits
   -- logically [1,1,1,248320], TILE-padded to 32 rows -- as a 15.9 MB
   physical buffer over card 0's Gen3 x4 link and untilizes it on the 4-core
   host: 10.4 ms/token measured (QWEN36_DEBUG_READBACK), device idle. One
   ttnn.to_layout(ROW_MAJOR) inside the traced forward makes the readback a
   0.5 MB row: interleaved C/T/C/T on the goal1 4k arm 62.96/62.73 ->
   50.93/50.47 ms/token (19.6-19.8 tok/s), equivalence rung 50.14/49.99 ms,
   text_sha256 identical everywhere. -12 ms/token, -19%, from one graph-level
   op. RECORDS.md had recorded ~12 ms "in the serving layer" a week ago.

3. attn_qkv 8x8 -> 11x4 replicated: 3/3 interleaved pairs +0.44/+0.38/+0.44
   ms at 4k, output-identical. The effect is ~0.4 ms (the standalone ~0.3 ms
   transferred ~1:1), not the ~1 ms a2's single pair suggested.

The model-side code for both ships as a patch in the next commit (research/
is gitignored scratch). Also: decode-grid-ab.sh stops containers gracefully
before rm -f (the wedge trap); decode-serving-ab.sh and
summarise-interleaved.py are the serving-path driver and the pairwise
summariser used for both records.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Round 2 of the serving-path A/B, with row-major logits as the control
(50.90/51.07/50.76 ms). Keeping each device's vocab shard row-major and
concatenating the two shards on host removes one CCL program per token and
the gather's 7.9 MB/device of tile-padded traffic: 48.90/49.07 ms/token
(20.4-20.5 tok/s), -2.00/-1.69 ms, output-identical. attn_qkv 11x4 on top
of row-major: -0.27/-0.46 ms, both pairs positive.

The timer inside process_output_decode saw only ~0.3 ms of the gather's
cost: reading device 0's replica returned when device 0 finished, and the
wait for device 1 landed between steps where the timer could not see it.
The endpoint mean is the statistic; the timer is the instrument for WHAT,
not for HOW MUCH.

patches/tt-metal/0013-qwen36-decode-host-readback.patch carries all three
(row-major default on, gather mode default none, attn_qkv 11x4) on top of
0012 with env opt-outs, plus the QWEN36_DEBUG_READBACK timer. research/ is
gitignored scratch, so the patch is the delta that ships. Driver: the
engagement regex now accepts '=value' in the one-shot line (the round-2
'did not engage' on the no-gather arms was that regex, not the arm).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Weights are read once per step regardless of batch, so the 23.34 ms weight read
is a fixed cost we currently amortise over exactly one token. Sizing the state
that batching actually costs: GDN rec_state is [B,Nv,Dk,Dv] fp32 over 48 layers
and KV is 16 layers x 2 heads/card, which at B=32 is 6.36 GiB -- with weights,
a little over half of the card's 32 GiB. Memory is not the constraint.

Modelled, that puts batch 2 at ~30.8 tok/s aggregate and batch 32 at ~276, while
per-user latency degrades from 15.81 to 8.62. So the target's definition decides
everything: aggregate 30 tok/s is a serving flag away, per-user 30 tok/s gets
nothing from batching at all. Every number we have measured is per-user, since
the harness runs --max_num_seqs 1.

Verified from the SoC descriptor rather than assumed: 140 Tensix workers at 1.5
MiB L1 each, ~165 MiB usable per card, 8 DRAM banks of ~4 GiB. That settles the
SRAM question -- 165 MiB against 10.76 GB of weights is 1.5%, so no arrangement
makes this model L1-resident and the weight read stays. The GDN state does fit,
and forward_decode already keeps the conv/recurrence/norm chain L1-resident, so
that win is taken.

The big RISC-V cores are recorded as UNVERIFIED: I could not find them in the
tech reports we snapshotted, and they matter only if they can drive dispatch.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The owner has ruled the target is one user seeing 30 tokens/second streaming,
not aggregate throughput. That excludes batching, which was the largest lever
available: it would reach 30 aggregate at batch 2 and ~276 at batch 32, but it
makes per-user latency worse, and per-user is the goal.

It also rules out Lever A as a solution on its own, by arithmetic rather than by
pessimism. Reaching 33.33 ms means cutting non-weight-read work from 39.93 to
9.99 ms, and even deleting every glue op entirely leaves 33.27 ms = 30.06 tok/s.
No amount of op fusion clears the target alone.

So speculation is the only per-user multiplier left, bandwidth being maxed and
arithmetic being 1.7 ms of the step. Modelling verify as sharing one weight read
across K+1 positions with flat glue puts the break-even at roughly 4 accepted
tokens of 11, and Lever A compounds it: halving the glue moves acceptance-3 from
28.0 to ~32.6 tok/s. The two levers stop competing and start multiplying.

Which makes the never-run experiment the most valuable one we have: MTP defaults
ON and every benchmark we own disabled it.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The previous section read the first ruling as "per-user latency" => "batching
is excluded". That was an over-reading and is retracted here. Production will
run B in [4,8]; the figure of merit is per-user tok/s AT that batch, with
aggregate a secondary benefit. A B=32 point at 8.6 tok/s per user is a worse
answer than B=4 at 18, and must not be proposed.

Consequences recorded:

- Baseline restated to the measured production step (50.70 ms / 19.7 tok/s at
  4k B=1) rather than the retired 63.27 ms figure.
- The step's dominant terms (weight read, per-op dispatch) are flat in B, so
  batching should be cheap in latency here -- but step(B) has NEVER been
  measured on the production arm, and the P0 itemisation cannot predict it
  because P0 profiled the composed GDN path, not the served model. step(B) for
  B in {1,2,4,8} is now a required measurement.
- Lever A is worth MORE: it cuts flat terms, so the saving is enjoyed once by
  every one of the B users -- the only lever improving per-user and aggregate
  together.
- Lever B (MTP) gets harder at B>1. Sequences accept different draft lengths,
  so a batched verify must either commit the batch minimum (discarding accepted
  tokens) or support ragged commit; which of those spec_decode.py does is
  unknown. A B=1 MTP result is necessary but no longer sufficient.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Round 3 measured patch 0013's three defaults as ONE arm against row-major
only, interleaved: 50.55/51.05 -> 48.47/48.84 ms/token at 4k, equivalence
rung 48.02/47.93 ms, text_sha256 identical. Against this morning's 63.27 ms
production baseline that is -23% and 15.81 -> 20.5-20.6 tok/s; the 50 ms /
20 tok/s goal is cleared on the production harness with no kernel work.

RECORDS.md gets the evening's table; PLAN-30TPS-EXECUTION section 1b gets
the scope correction pointer (the P0 profile ran the composed GDN path).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
The first MTP measurement attempt produced no performance data. It loaded the
model fine (mtp_head=true, so the head is genuinely there) and then died
building writer_unary_stick_layout_wh_multicore with a compile-time-args
static assertion -- a host/kernel version mismatch inside the image. Pins did
not move as a set.

Three setup faults worth more than the failure itself:

- It ran vllm-tt:mtp (built 15h earlier), not the vllm-tt:k2 image that
  produced the validated 50.70 ms/token result.
- It set QWEN36_MTP=1, which is NOT the speculation switch. QWEN36_MTP
  (model_config.py:98) gates LOADING the MTP head; QWEN36_SPEC
  (text_demo.py:512) gates USING speculative decode, and both default to on.
  The run never set QWEN36_SPEC, so its "plain" controls would have been
  speculative too unless the harness set it in-process. Future A/Bs must
  toggle QWEN36_SPEC and assert engagement from the log in BOTH directions.
- It carried experimental flags (MLP_DOWN_BF4, GDN_DECODE/STATE_BF16) and
  omitted DECODE_LOGITS_RM, so it was neither the validated baseline nor an
  accuracy-checked configuration.

Also: on failure the process deadlocked in teardown (24 threads sleeping, 0%
CPU) and held the cards until docker stop + tt-smi -r; and it wrote the full
brisc build log into both run.log (1.99 GB) and the record's error field
(663 MB for 2 records). Truncate embedded build logs.

The arm matrix is sound and should be reused as-is.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
We believed two things that a careful read of the tree disproves.

First, that the open question at B>1 was WHICH commit policy speculative decode
uses -- batch minimum, or ragged. It uses neither, and there is no assert to
trip either. `generate()` takes one flat prompt id list, `p` is a scalar int,
and `_commit(mi)` applies ONE accepted-prefix index to all 48 GDN layers. The
batch axis is already spent: the K+1 CANDIDATES of a single sequence are what
`use_fullbatch_verify` batches over, and attention's `_SPEC_SDPA_L1_FIT` admits
T in {4,8,12} = K+1, never B. So a B=1 acceptance rate cannot be carried to the
owner's B=4-8 operating point by a flag; it needs per-sequence positions, a
[B,K+1] verify and ragged commit, which is a redesign.

Second, that MTP could be measured on the validated production configuration.
It cannot, with the images on this host. vllm-tt:k2 -- the image behind the
50.70 ms/token result -- has no mtp.py, no spec_decode.py, and zero of the six
speculation hooks across model.py, gdn/tp.py and attention/tp.py. vllm-tt:mtp
has the complete spec stack and none of k2's conv_gates / packed_qkv /
fused_inplace GDN work; QWEN36_DECODE_LOGITS_RM is in neither image (it was
bind-mounted). The branches are disjoint, so an MTP run yields a valid internal
ratio and an absolute ms/token that is NOT the production arm's.

Also fixes the harness that failed its first attempt: both arms now go through
the production dispatch toggled by QWEN36_SPEC (not QWEN36_MTP, which only
gates LOADING the head) with engagement asserted in both directions from the
demo's own log lines; exception text is truncated (a JIT build failure wrote a
663 MB JSONL for two records); a failed arm exits hard rather than deadlocking
in teardown while holding the cards; and the flag set drops two knobs this tree
does not implement. Adds smoke-demo.sh, a zero-custom-code image-health gate to
settle whether the kernel-build failure indicts the image or the harness.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
that has none of our performance work

Two structural facts, both verified independently of the agent that found
them, and together they reshape the plan.

1. The MTP tree and the validated-performance tree are DISJOINT. vllm-tt:k2
   has QWEN_GDN_CONV_GATES and QWEN_GDN_PACKED_QKV but no spec_decode.py, no
   mtp.py and zero hits for capture_verify_trace. vllm-tt:mtp is the exact
   mirror. QWEN36_DECODE_LOGITS_RM is in NEITHER -- it was bind-mounted.
   So a SPEC=0 control on the only tree that has MTP cannot land at 50.70
   ms/token and will land nearer 63 for a known reason. An A/B there gives a
   sound internal RATIO; that ratio must not be multiplied onto 50.70 ms.
   Merging the branches is an integration task now on the critical path.

2. Speculation has no batch dimension over sequences at all -- it is neither
   batch-minimum nor ragged commit, and there is no assert to trip because
   nothing ever offers it a batch. _commit(mi) applies one SCALAR accepted
   index across all 48 GDN layers, so a per-sequence accepted length is not
   expressible. The batch axis is already spent on CANDIDATES:
   use_fullbatch_verify batches the K+1 candidates of one sequence and
   _SPEC_SDPA_L1_FIT admits T in {4,8,12} = K+1, never B. At B>1 it fails on
   shape, not silently.

The owner's operating point is batch 4-8 with per-user latency decisive.
A B=1 acceptance number therefore does NOT transfer to it by setting a flag:
batched speculation needs per-sequence positions, a [B,K+1] verify and ragged
commit. That is a redesign, not a configuration.

Also corrects the previous record: the failed run's control arms were NOT
secretly speculative (the driver set QWEN36_SPEC=0 and called the spec entry
point directly for treatment). And prose_stock, the demo's own "4k" arm, is
2642 tokens of one tiled paragraph at 0.04 distinct-token fraction -- unusable
for acceptance.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Owner verbatim: "good fully drop MTP and focus on maximizing B=4 and B=8
aggregate throughput now."

This retires the per-user framing that has driven the last several sessions.
Aggregate tok/s at B=4/B=8 is now the figure of merit; per-user is reported,
not optimised. Since aggregate = B x 1000 / step_ms(B) and B is fixed by the
owner, the entire programme reduces to shortening the batched decode step.

Measured starting point, production arm, 4096-token prompts:
  B=1  19.63 per-user  19.63 aggregate  50.9 ms
  B=4  17.42 per-user  68.97 aggregate  57.4 ms
B=4 costs 11% of per-user latency for 3.5x aggregate. The sublinearity is the
flat weight read: 23.34 ms paid once per step regardless of B.

MTP is dropped WITHOUT ever having been measured. Its acceptance rate is
unknown and this must not be read later as a negative result. It is dropped
because it was structurally the wrong thing to fund: speculation is B=1 by
construction, the MTP and performance trees are disjoint, and speculation is a
latency trick -- the wrong lever once aggregate is the goal and the batch axis
is already carrying sequences rather than candidates.

Consequently the remaining opportunity is exactly where P0 pointed: the
~34 ms of the B=4 step that is NOT the weight read. Per-op cost and op count.
That work is worth more now than before, because each millisecond cut from a
flat term is divided among B users instead of one. Prefill also re-enters
scope, since at serving batch it shares the device with decode and was never
saturated either.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
PARKED, NOT ABANDONED. The owner dropped MTP on 2026-09-12 in favour of
maximising B=4/B=8 aggregate throughput. The decision rests entirely on two
structural findings and on NO measurement.

State it plainly, because a parked harness invites the wrong reading later:
THE ACCEPTANCE RATE WAS NEVER MEASURED. Not on prose, not on code, not on
question answering, not on the counting control, at no K, on no hardware. The
single attempt died in a JIT kernel build on arm 0 -- a PLAIN control arm --
before any speculation ran. The verdict is "not measured, deprioritised", never
"measured and rejected". If speculation is reopened, acceptance is still the
number the decision turns on and it is still unknown.

What justified dropping it without the number: batched speculation is
unimplemented BY CONSTRUCTION, so a B=1 acceptance rate could not have reached
the B=4-8 operating point without a redesign; and the MTP tree and the
validated-performance tree are disjoint branches, so it could not have been
measured on the production configuration at all. Either alone makes a B=1
acceptance rate insufficient to justify the work; together they make measuring
it first the wrong order.

Kept because it is the expensive half to rebuild: the prompt set (five
~4096-token prompts graded by OUTPUT entropy, instructions verified intact,
count_control included as the deliberate low-entropy reference that would have
shown how much our counting benchmark inflates acceptance), the interleaved
C/T/C/T matrix, and the engagement assertion that makes a mis-toggled arm a
failed arm rather than a data point. Every file says PARKED in its first lines.

Also unresolved and recorded as such: whether the kernel-build failure was a
broken image or a stale entry in the SHARED TT_METAL_CACHE. The fresh-cache
test was prepared and stood down before it took the lock, so the cause is
unknown. The existing metal-mtp cache was deliberately left intact.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The B=1 stack (48.5-48.8 ms, patch 0013 defaults) has to be re-established at
the owner's new operating point, B=4 and B=8 aggregate throughput. The driver
hard-coded SEQS=1 and decode-bench.py; SEQS=<B> now launches max_num_seqs=B and
measures with batch-decode-bench.py (steady-state window, per-user and
aggregate), and the summariser reads either statistic. Also lands the
RECORDS.md full-stack row that the b8ccd92 doc script failed to insert.

No hardware touched: the lock is held by the batch-scaling run.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
We believed our decode benchmarks said something about output. They do not.
decode-bench.py and batch-decode-bench.py build a seeded-word body and then
append "Now count from 1 to 300, separated by commas." The body is dead
weight: the continuation does not depend on it. In
bench/runs/batch-scaling-20260912T192000Z.jsonl every stream at B=1, B=4 and
B=8 reports the SAME text_sha256 (cd99c8b0e7af4351...) across different
prompts, different batch sizes and different context lengths. So every
batch-scaling and decode-bench record is a timing record only; none of them
carries any evidence that the output was right, and a silent GDN or SDPA
corruption -- the class LoFi+fp32_dest_acc and k_chunk_size=32 already
produce with no error -- would have read as a clean throughput win.

This adds the missing arm, following bench/prefill-bench.py --equiv's
conventions (same HTTP client, temperature 0, enable_thinking off, sha256 of
the completion, one JSON object per run into bench/runs/):

- eight prompt-DEPENDENT ~4k prompts over two committed corpus fixtures, with
  disjoint bodies and the instruction tail appended AFTER clipping (the
  bench/mtp lesson: clipping in text silently truncated three instructions);
- the assertion that makes it a gate: the B completion hashes must be
  DISTINCT, and no stream may be near-single-token. --degenerate swaps in the
  counting tail as a live control that must trip it;
- --capture / --compare: reference capture and a per-prompt diff, exact match
  required at temperature 0, common-prefix reporting and an explicit
  non-authoritative verdict above it;
- --invariance: the same probe prompt at B=1, at B=4 submitted first and at
  B=4 submitted last with different companions, plus a B=1 repeat, requiring
  all four to be byte-identical. We control submission order, not the
  physical slot, so this tests invariance to batch composition and arrival
  order -- which is what comparing numbers across B actually rests on.

Every path exercised off-hardware against a mock endpoint: capture/compare
exact, corrupt server -> exit 1, degenerate control -> exit 1 (1/4 distinct).
No device numbers here yet; the baseline capture follows.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The comment in qwen36_vllm.py is the only source of the claim, and it is
right: the plugin's overlap needs device sampling plus a device-resident
token/position chain, and Qwen has neither (no sampler on the 1x2 mesh, rope
re-staged from host every step). A stale host token re-fed to an in-place
GDN recurrence cannot be overwritten or rolled back. Records what would have
to change and why on-device argmax is the prerequisite, not an alternative.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
bench/runs/batch-scaling-20260912T192000Z.jsonl reports one identical
text_sha256 for every stream at B=1, B=4 and B=8. The counting tail makes the
prompt body irrelevant, so those ms/token numbers say nothing about output and
must not be cited as evidence that batching preserved it. Points at the new
equivalence arm and states its limits.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Three additions the performance agent's async-decode analysis makes
load-bearing, all built and exercised off-hardware.

1. A mismatch is no longer one undifferentiated count. vLLM greedy takes the
   LOWEST TOKEN INDEX on a tie and bf16 ties genuinely occur, so a correct
   ttnn.argmax that breaks ties the other way fails a sha comparison while
   being arithmetically right. The reference now stores per-position top-k
   logprobs, and the first divergence is classified TIE_BROKEN_DIFFERENTLY /
   NEAR_TIE / WRONG_TOKEN with separate counters, because "fix the tie rule"
   and "hunt a corruption" are opposite responses and a pooled count hides
   which one you are looking at. Limit recorded in the output: the served API
   exposes token TEXT, not ids, so the device arm can confirm a tie but not
   that the lower index won.

2. --host-greedy builds a reference with transformers in the SAME schema,
   resolving the argmax with vLLM's lowest-index rule explicitly and carrying
   token ids, so the index rule is verifiable there. This matters because a
   device-vs-device diff proves only that the machine repeats itself; a
   systematic error sits in both runs and passes. Every compare record now
   states reference_kind and what that kind actually proves. NO host-greedy
   reference exists yet and the docstring says so plainly: ~54 GB of bf16
   weights against cfx-llm2's ~5.5 GB usable and this box's 23 GB with no 27B
   checkpoint. The mode refuses on insufficient RAM rather than SIGKILLing
   with no record (the exit-137 trap in CLAUDE.md).

3. Invariance gains a B1_then_joined point: the companions arrive after the
   probe has been decoding. A host<->device sampling flip triggers an input
   reload and there are two trace sets, so batch composition changing mid-run
   is a live serving path and invariance across it was unchecked.

Verified against a mock endpoint: identical run -> 4/4 EXACT rc 0; an exactly
tied alternative -> MISMATCH_TIE_BROKEN_DIFFERENTLY, gap 0.0, rc 1; an
out-of-top-k token -> MISMATCH_WRONG_TOKEN, rc 1; counting tail -> rc 1;
B>8 -> rc 2; --host-greedy without a model -> rc 2. No device time used.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
We believed the P0 per-op profile was invalid because it "ran the composed GDN
path". That was right, but the reason was not pinned, so nothing stopped us
repeating it. It is now pinned: there are THREE qwen36 source trees on
cfx-llm2 and they are not interchangeable.

  /var/lib/models/ttm-src   model.py aee016ef9af3   no conv_gates AT ALL
  vllm-tt:k2 image          model.py 488960e9e151   has conv_gates
  production                image tree + the three bind-mounted patched files

grep for QWEN_GDN_CONV_GATES / gdn_decode_conv_gates / QWEN_GDN_PACKED_QKV in
ttm-src returns nothing. bench/profile/decode_step_profile.py drives text_demo
against exactly that tree under the Nix Tracy ttnn 0.77, so it could not have
run the fused kernels whatever env it was given -- while production serves
ttnn 0.79 with those flags set to 1. The old profile differed from production
in the source tree AND the ttnn version, on top of text_demo vs vLLM.

The second thing we did not know: the production image is itself Tracy-enabled.
libtt_metal.so exports tt::tt_metal::DeviceProfiler, kernel_profiler.hpp ships,
and the tracy post-processing package is installed. So the serving binary can
be profiled directly -- no separate build, no synthetic state, no text_demo.

This adds the harness for that and nothing else; no optimisation is landed and
nothing was run on hardware (the cards were held by serving-ab-b48 throughout).

The per-op arm serves with trace_mode='none' because a traced step cannot be
profiled per-op: the kernel profiler's marker buffer is a compile-time size and
overflows inside one replay, and the C++ report leaves OP NAME empty for
replayed ops. The plugin's two-phase warmup compiles the untraced path and
captures the trace from it, and asserts the two use the same ops; the harness
checks that rather than trusting it, by comparing program counts against a
traced+profiled arm.

Every arm now FAILS CLOSED unless the container logs "QWEN_GDN_CONV_GATES
engaged". Absent that line the run is the composed path and is recorded
INVALID rather than tabulated -- which is the specific error this commit
exists to make unrepeatable.

The TT_CFG_TRACE_MODE hook in vllm-tt-serve.sh is additive and default-unset;
the production invocation is byte-identical to before.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Source audit of the PRODUCTION tree (image qwen36 + the three bind-mounted
patched files), with the production env applied to resolve every branch. No
hardware; the cards were held by serving-ab-b48 throughout.

Three findings that would have been the headline are retracted before costing
them, because they are dead code in production: the ~30-op GDN recurrence
composite, the 9-op conv shift register, and an fp32 reduce-scatter. CONV_GATES
and PACKED_QKV collapse those to two device ops and the two GDN_BF16 flags
remove every typecast. A production GDN layer is ~31 programs, and the step is
~2,400 programs, not the 4,776 the composed-path profile reported. The
tile-padded logits readback is retracted too: LOGITS_RM already fixed it.

QWEN_GDN_FUSED_DECODE=1 is INERT. It gates the unpacked recurrence entry, which
the packed path never calls. It is set in every serving arm and named in every
A/B label we have written, and it does nothing. No measurement is invalidated,
but the labels are misleading.

What survives, all confirmed by direct read of the patched files:

The seven tuned decode matmuls carry hardcoded num_cores literals that are TP=4
constants -- 44/44/33/64/44/33/33 on a 110-core grid. The provenance comment
still cites the TP=4 1536x5120 shape; at TP=2 it is 3072x5120, so these grids
do twice the per-core work on a third of the chip. attn_qkv is additionally the
only one of the seven that omits grid_w, so it shapes 8x8 rather than 11-wide.
The LM head, the largest matmul in the step, has no program config at all.

The practical consequence is that the top item is a SWEEP OF ENV VARS, not a
code change: the patched model_config ships QWEN36_1D_GRID_<NAME> for all seven
plus the LM head and logs every resolved grid, and none are set in production.
bench/ab/decode-serving-ab.sh already runs arbitrary env arms interleaved at
SEQS=8, so costing them needs no new code either.

Also logged and not actioned: 256 CCL programs/step whose payloads are 3/4 tile
padding at B=8, 224 DRAM programs/step of partial-RoPE relayout that exist only
because the producer discards the layout the consumer wants, and a 1x4 output
subblock cap forced by fp32_dest_acc_en.

Every millisecond figure is either cited from an existing record or explicitly
marked unmeasured. Op counts are exact; time attributions are not yet evidence.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Source-exact per-layer op counts for the production tree, resolved with the
production env. Not a metric and not a time attribution: CLAUDE.md is explicit
that operation count is not an objective. This exists so the per-op profile has
something to be checked against, and so the profile can price these blocks.

A GDN layer is ~27 device programs; a full-attention layer is ~55. The step is
~2,183 (bracket 1,900-2,500) against the invalidated P0 profile's 4,776 for the
composed graph -- being under half is independent corroboration that P0
measured a different model.

The asymmetry is the interesting part. There are 16 attention layers to 48 GDN,
yet they contribute a comparable share of the count, because the GDN layer has
been fused hard (conv_gates replaces 12 ops, packed replaces ~26) and the
attention layer has not. Partial RoPE alone is 224 programs, ~10% of the step,
and exists only because apply_partial_rope_decode transposes an interleaved
tensor in and out to dodge a reshard -- after _make_heads_decode has already
discarded the HEIGHT_SHARDED layout the native decode-mode rotary op wants.

Whether 10% of the programs is 10% of the TIME is exactly what the profile has
to settle. These are small tensors and the answer could go either way, so this
is logged as a hypothesis to price, not a finding.

One thing we cannot settle from source and must read from the container log:
whether QWEN_GDN_FUSED_INPLACE actually LANDS. If honoured it removes a
6.29 MB/layer state copy, 302 MB/step across 48 layers. The code logs
"in-place happened=<bool>" once by comparing buffer_address(), and no record we
hold contains that line. The state is bfloat16 here, which is a plausible way
for an _inplace variant to refuse. The harness greps it; nobody should assume
it took. Bucketing also disables it whenever fewer than 8 requests are active.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The link to cfx-llm2 was observed flapping tonight -- reachable at 23:11Z,
gone at 23:12Z, after a ~45 minute outage from 21:45Z. The box never rebooted
(uptime 5d17h, /tmp intact), so this is the network, and it will happen again.

With the runner attached to the ssh session, a dropped link SIGHUPs it
mid-device-work. That is the precise scenario that wedges the cards: the EXIT
trap fires while a tt-metal container is doing device work, and a graceful
docker stop racing a mid-flight profiler drain is not something to invite when
the recovery is a board reset.

So the runner now launches under setsid+nohup and outlives the ssh session,
and the driver merely follows its log, reconnecting across flaps. The run
never notices. The runner keeps its own EXIT/INT/TERM trap, so the lock is
still always released when the run itself ends -- which is the case the trap
is actually for.

Still no exec anywhere in this path; an exec destroyed the trap and stranded
the lock earlier this week.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Observed read-only from this session while waiting for the lock. dmesg has two
global OOM kills, the second naming VLLM::EngineCor directly (total-vm 77.8 GB,
anon-rss 2.46 GB) at 00:07:36Z. free shows 15 GB total, 13 used; tt-metal's
~10 GB of hugepages leaves the ~5.5 GB CLAUDE.md already documents.

That is what the two "server never came up" arm logs are. Their visible tail is
only the downstream "Engine core initialization failed ... Failed core proc(s):
{}" -- and that EMPTY set is the signature of a SIGKILLed engine core, not of
an engine-core exception, which is why the logs carry no root cause of their
own. Neither log contains a card-wedge signature.

Worth being ready for: the OOM was a SIGKILL of a tt-metal process that was at
least loading weights, which is the situation CLAUDE.md says can leave the next
device open failing "Timeout waiting for physical cores to finish" / "failed to
initialize FW!" and looking like a version mismatch. This session reset nothing
and opened no device.

The arm log also happens to carry the model's own resolved-grid banner, which
turns two source-only audit findings into hardware-log evidence: the LM head
really does run on ttnn-auto with no program config, and the hardcoded decode
grids reach the device as 44/44/33/44/33/33 on a 110-core chip. The record
includes the per_core_N each resolves to and one arithmetic worked example, and
claims no time saving from any of it -- a 2.35x standalone matmul win has
already measured 0-2% SLOWER in the model once, so the sweep has to say.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Committed in advance: the three patch-0013 savings are flat in B, so the
full stack would measure ~55 ms at B=4 and ~62.5 at B=8. Measured, interleaved
against the logits_rm control: 74.21/74.56 vs 57.45/57.50 at B=4 (+17.0 ms,
2/2 pairs) and 95.68 vs 64.61 at B=8 (+31.1 ms, n=1 control). The model was
wrong. The culprit is unidentified between the attn_qkv 11x4 grid and the
no-gather host concat path; the run bundles both and cannot separate them.
Timing only (one text_sha256 across every stream).

Also records why two B=8 arms never booted: host OOM from 8.57 GiB of pages
pinned by the tenstorrent module with no process alive, which also looked like
a network outage. Line 6 of the jsonl was rebuilt from the raw stream log.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
The capture took the lock and reached cold weight load, then the host went
down. Cause (found by the performance agent, not by this arm): the
tenstorrent kernel module was holding ~8.57 GiB of pinned pages with no
tt-metal process alive, leaving MemAvailable ~1.8 GB, so the host OOM-killed
whatever loaded weights next -- it had already killed two A/B arms and, once,
systemd itself, which is what the "network outage" really was. A board reset
does not fix this: tt-smi -r resets the cards, not pinned kernel pages.

Recorded as an ATTEMPT, explicitly is_reference=false and usable_baseline=
false, because a partial reference is worse than none -- it would silently
license "nothing changed" comparisons against an incomplete baseline. No
ref-B*.json was fetched and none exists.

Zero points completed and no stream was ever issued, so this says nothing
about bench/batch-equiv-bench.py and nothing in the harness was tuned in
response to it. The baseline is still owed, and the numerics-affecting items
(LM_HEAD_GATHER_MODE, fp32_dest_acc_en) stay ungated until it exists.

Also flags a possibly stranded /tmp/ttlock owner=equiv-arm on cfx-llm2: a
SIGKILL does not run the EXIT trap. Flagged, not cleared.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Follow-up observation with the cards free and nothing running. The box idles
with 1.9 GB MemAvailable against the 2.46 GB anon-rss the OOM-killed EngineCore
had reached, which is why this failure mode recurs. It is a standing fragility
of the host, not damage: the 4 GiB hugepage reservation is fully free, total
process RSS is under 300 MB, and swap is nearly untouched, so the ~13 GB of
'used' is kernel/driver-held -- most plausibly the tenstorrent driver's pinned
host memory for the two cards. Recorded as an observation, not a diagnosis.

Separately and not implicated in the OOM: the SATA root is now 97% full with
2.2 GB free, against the 91% CLAUDE.md already records. Profiler output goes to
the NVMe, which has 184 GB.

Nothing was freed, dropped, deleted, restarted or reset, and no device was
opened. This session has still never held the lock.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
/tmp survived the reboot, so the first attempt's logs were readable after all,
and two things in the record needed fixing.

The lock was NOT stranded. I had flagged it as probably stranded on the
assumption that a SIGKILL skipped my EXIT trap. In fact only the containerised
EngineCore was killed; the runner survived, reported "server never came up"
and logged "lock released". Being detached is why that log existed to read at
all.

The OOM diagnosis, which this record originally relayed second-hand, is now
corroborated independently and precisely: the core's last line is "Warm
state_dict: 851 weights, 80 real host weights from the sidecar", after which it
died with no traceback of its own and the API server reported "Failed core
proc(s): {}". An empty failed-proc set with no core traceback is death by
signal, landing exactly on the step that allocates host weights.

Still zero points, still no baseline, still is_reference=false.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The baseline exists. B=4 and B=8 device references captured on the production
arm (logits_rm + a1b), 4/4 and 8/8 completion hashes distinct on real prompts,
distinct_token_frac 0.71-1.0. The degenerate control tripped the assertion as
designed (1/4 distinct, exit 1), so the gate is live rather than decorative.
reference_kind=device throughout: these license "nothing changed since this
point", not "the output is right".

The in-range result, and it is not the comfortable one. The same prompt,
greedy at temperature 0, on ONE server instance, decodes DIFFERENTLY at B=4
and at B=8 in 3 of 4 shared prompts. Divergence occurs after an identical
prefix, so it is a forward-pass difference at that position, not drift from an
earlier split. Classified WRONG_TOKEN, not TIE_BROKEN_DIFFERENTLY: the gaps
are 0.125, 0.25 and 1.75 logprob against the B=4 reference's own top-k. The
two small ones are 1 and 2 bf16 quantisation steps and a tie cannot be fully
excluded for them; 1.75 is ~14 steps and cannot be a tie.

What is invariant: slot position and arrival order. B4_first, B4_last and
B1_then_joined (companions arriving 3 s into the probe's decode) all produced
the same sha. Only the WIDTH changes the output.

B=1 vs B=4 also differs; recorded as seen-and-deprioritised on owner
direction, since production never runs B=1.

Adds --diff-refs, which answered the cross-width question offline from the two
captured references with no additional device time.

Escalated, not developed. This arm cannot say whether this is the expected
reduction-order floor or a defect: both sides are device references, neither is
known correct, and no host-greedy reference exists.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Ships the remote body as a file and starts it under setsid nohup with stdin
closed, then follows the log over disposable ssh connections. The link flapped
repeatedly; this is why the run survived and why the first attempt's logs were
readable afterwards. No exec anywhere in the path.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Two pieces of prep for the measurement half, both done before taking the lock.

The driver now verifies the three bind-mounted patched files by CONTENT against
a pinned manifest, not by existence. /tmp on cfx-llm2 is shared with every other
arm, so a stale or half-copied file would sail past an `ls` and we would profile
the unpatched image sources -- reproducing precisely the error that invalidated
P0, and producing a plausible-looking table for a model we do not serve. The
check is proven in both directions: it accepts the live snapshot (which matches
what the serving arm mounts, all three sha256) and rejects the image's own
unpatched copies. REPIN=1 re-pins when the serving arm legitimately moves.

Second, a prediction on the record before measuring it. The performance agent's
full stack lost 17 ms at B=4 and 31 ms at B=8, culprit unidentified between
attn_qkv 11x4 and the no-gather host concat. The discriminator is the shape of
the growth rather than its size: a decode matmul's M is one 32-row tile at B=4
and B=8 alike, so an mcast grid change is essentially B-independent on device,
while no-gather deletes a device all-gather and pays for it on the host with a
strided two-shard concat that is O(B). A loss growing ~3.5 ms per added sequence
is a per-row cost, so the call is no-gather, and attn_qkv should measure flat.

Falsifying it needs no new code -- decode-serving-ab.sh already has all three
arm names -- and the per-op profile may settle it for free, since attn_qkv's
device time and the host concat land in two independent instruments.

If attn_qkv alone shows a batch-growing loss the prediction is simply wrong,
and that would be the more interesting finding.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Measured on cfx-llm2 at B=8, 4k, production image and sha256-verified patched
sources. Anchor reproduces: 64.188 ms/token vs the batch-scaling record's
64.776, a 0.91% agreement, so the harness measures the same thing.

RETRACTION. I claimed the production vllm-tt:k2 image is Tracy-enabled. It is
not, and the per-op device table is therefore not achievable with the artefacts
we have. The runtime says TT_FATAL: TT_METAL_DEVICE_PROFILER requires a
Tracy-enabled build. I had inferred support from tracy symbols in
libtt_metal.so, kernel_profiler.hpp shipping, the tracy python package, and the
rtoption strings -- all of which are equally consistent with a gated,
non-Tracy build. Worse, the specific string I read as evidence of support is
the text of the guard that now fires. I read a refusal as a capability, and I
should have tested it rather than inferred it, exactly as I later tested the
checksum guard. Using the Nix Tracy ttnn instead is not an option: it is 0.77
against production's 0.79 and is the ttm-src tree with no fused GDN, which is
precisely the P0 invalidity.

What the run did establish is bigger than the table would have been, and it
points away from the work the old profile implied.

Trace replay is worth 176.0 ms per step at B=8: traced 64.188 ms against
untraced 240.201 ms, 3.74x, roughly 80 us per program of host dispatch that
tracing removes. So in production, per-op dispatch is ALREADY amortised to near
zero, and removing device programs buys essentially nothing unless those
programs consume real device time. The 224-program RoPE relayout and the 256
CCL programs are not worth what their counts suggest. This is the measured
justification for CLAUDE.md's rule that op count is not an objective, and it
retires the "fewer, bigger device programs" framing that
docs/WHERE-THE-BOTTLENECK-IS.md drew from the eager P0 profile.

Three findings confirmed from the device's own banners rather than from source:
ATTN_QKV really does run the unshaped 8x8=64, alone among the seven; the LM
head really is on ttnn-auto with no program config; and QWEN_GDN_FUSED_INPLACE
reports "in-place happened=True (same buffer=True)", so the ~302 MB/step of GDN
state copies is already eliminated and that candidate is spent, not available.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Prints share-of-wall-time per frame converted to ms/step using the measured
step time of the same arm, and -- the point of it -- separates frames that are
BLOCKING ON THE DEVICE from frames actually spending host CPU. The logits
readback looks like host time in a stack sampler and is not; counting it as
host cost would double-count the device step.

On run A's untraced arm it independently corroborates the trace finding:
78% of host time, 187.8 ms, sits in ttnn/decorators.py:728, the op dispatch
wrapper, against a 176 ms traced-vs-untraced delta measured end to end. Two
instruments, same answer.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Measured at B=8, 4k, production image and sha256-verified sources, four
interleaved controls. Every arm passed the conv_gates engagement assertion, so
all of it is the production graph.

The control spread over four arms is 0.309 ms, 0.48%. That is the resolution
limit and nothing smaller is claimed.

ALL FIVE QWEN36_1D_GRID_* levers are inside it. ATTN_QKV 11x4 -0.23%, GDN_OUT
11x10 -0.23%, ATTN_WO 11x10 +0.01%, MLP_DOWN 11x10 -0.02%, LM_HEAD 11x10 with
in0_block_w=2 -0.25%. So the audit's top-ranked structural finding -- that the
seven decode matmuls carry TP=4 core-count constants and leave most of a
110-core grid idle -- is TRUE, confirmed from the device's own banner, and
WORTH NO MEASURABLE TIME. Filling the grid, re-shaping the mcast and giving the
LM head an explicit program config each move the step by less than run-to-run
noise. CLAUDE.md already recorded a 2.35x standalone matmul win measuring 0-2%
slower in the model; this says the same thing at batch with four controls.

The +17/+31 ms full-stack regression is attributed, and the pre-registered
prediction held. QWEN36_LM_HEAD_GATHER_MODE=none alone costs +30.998 ms at
B=8 -- reproducing the reported +31 ms to within 0.002 ms once the other
treatment is removed -- while ATTN_QKV 11x4 measures -0.148 ms and is innocent.
The mechanism was called in advance: a decode matmul's M is one 32-row tile at
both widths so a grid change is B-independent, while no-gather deletes a device
all-gather and pays for it on the host with a strided two-shard concat that is
O(B). Do not ship GATHER_MODE=none at batch; it may still be right at B=1,
where it was originally measured.

The host tail is attributed for the first time, by py-spy sampling the
EngineCore from outside the container without ever ptrace-stopping it. The
76.5% frame is model.py:3358 `_dev.cpu()` -- the BLOCKING readback, which is
the device step and not host cost; at 49.5 ms it independently corroborates the
prior record's ~51 ms d2h. The real host CPU is ~15 ms and it is DIFFUSE: op
dispatch outside the trace 2.56, output plumbing 2.38, host greedy sampling
2.29, output tokens 1.76, host RoPE cos/sin 1.29, the f32 upcast 1.02. Nothing
above 2.6 ms, so there is no large single host win either. The largest coherent
group is host sampling plus output plumbing at ~8.1 ms, which an on-device
greedy argmax would target -- and the TopK 64K-shard limit that blocks the
stock device sampler at TP=2 does not bind a greedy argmax. That remains
unbuilt and therefore uncosted.

No arm is described as output-identical. The counting prompt cannot detect a
numerics change, and batch width alone is already known to change greedy
output at temperature 0.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Best remains B=8, 64.8 ms, 121.4 aggregate tok/s, reproduced by three
harnesses within 0.3%. What changed today is that we now know where the
remaining time is NOT.

- Op-count optimisation is dead, with a number behind it. Trace replay is
  worth 176 ms/step at B=8 (64.19 traced vs 240.20 untraced, 3.74x), so
  per-op dispatch is already amortised to ~zero in production. Removing
  programs buys nothing unless they consume device time. The 224-program RoPE
  relayout and the 256 CCL programs are not worth what their counts suggest,
  and the "fewer, bigger device programs" framing that WHERE-THE-BOTTLENECK-IS
  drew from the eager P0 profile is retired.
- Grid tuning is dead. All five QWEN36_1D_GRID_* levers land inside a 0.48%
  control band. The seven decode matmuls carrying TP=4 core-count constants on
  a TP=2 deployment are real and confirmed from the device's own banner -- and
  worth no measurable time. A real inefficiency that costs nothing is not an
  opportunity.
- The +31 ms batch regression is LM_HEAD_GATHER_MODE=none, alone: +30.998 ms
  at B=8, matching the observed +31 to 0.002 ms. ATTN_QKV=11x4 is -0.148 ms
  and innocent. The mechanism was pre-registered before measurement and held:
  M sits in one 32-row tile at both widths so a grid change is B-independent,
  while no-gather pays an O(B) host concat plus a second PCIe transfer.
- There is no large single host win either. The 76.5% py-spy frame is the
  blocking readback -- the device step, not host cost. Real host CPU is ~15 ms
  and diffuse, nothing above 2.6 ms.

The one lever left worth building is host sampling + output plumbing, ~8.1 ms
of coherent cost, where the TopK 64K shard limit that blocks the stock device
sampler does not bind a greedy argmax. Unbuilt and uncosted; it is also the
prerequisite for async decode.

Recorded honestly: the per-op device millisecond table was NOT produced.
vllm-tt:k2 is not Tracy-enabled, and the string read as evidence of support
was the text of the guard that fired.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Owner: build the device greedy argmax, then work toward async decode
Some checks failed
tt-stack-ci / Build simulators and check the host module (pull_request) Failing after 6s
tt-stack-ci / Report upstream drift (pull_request) Successful in 9s
f5cdd7d984
Owner verbatim: "Yes agree let's pursue this argmax and work towards async
decode if that will make interactions with the TT HW more smooth and better
UX overall."

This is the only lever the audit left standing -- ~8.1 ms of coherent host
cost at B=8 (host greedy sampling 2.29, output plumbing 2.38, output tokens
1.76, f32 upcast 1.02), against a measured step of 64.8 ms.

Two phases, and phase 1 pays on its own:

1. A greedy device argmax for the 1x2 mesh: per-shard ttnn.argmax with the
   2-way combine moved on device, under SYNC scheduling. Removes the 2 MB
   logits readback, the .float() upcast and the host argmax.
2. Only then device-side position advance and a device-indexed RoPE cos/sin
   table -- which _tt_vllm_always_refresh_decode_trace_inputs=True currently
   prevents -- and then async scheduling.

Phase 1 is the prerequisite for phase 2, not an alternative: async without a
device sampler feeds step N+1 the stale token N-1 at position N-1, and the 48
GDN layers' in-place recurrence has nothing to overwrite and no rollback.

Two gating rules that must hold, both already paid for in analysis: the
device-sampling decision is batch-wide per step, so an argmax-only sampler
must require every row temperature==0 or fall back for the whole step; and
vLLM greedy breaks ties to the lowest index, which ttnn.argmax must match or
a correct implementation still fails equivalence.

Equivalence bar is fixed-width identity against the B=4/B=8 references.
Cross-width identity is not available and must not be demanded.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
multica-agent changed title from Thatch-repro P1 on 2× P150a (override MTP-primary) — DSpark tip path to Qwen3.8-27B on 2x P150a: measured serving recipe (121.4 aggregate tok/s at B=8, 20.5 single-stream) 2026-09-13 10:24:55 +02:00
Sign in to join this conversation.
No reviewers
No labels
human-approved
No milestone
No project
No assignees
3 participants
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set

Reference
bitpartner/tt-stack!19
No description provided.