mHC into the loop — what was wired, what it claims, what it does not

Date 2026-09-30 · worktree wt/mhc-model off 9d571c2 · lane fourth residual arm (after use_gr, use_attnres) · A/B row 7b, registered, NOT run. Revised 2026-10-01 — citation corrected and the base paper’s mechanism verified against the crate (§0, §1), mhc_streams moved 4 → 2 and the launch-count cost made first-class (§3.2, §6), two defects found in my own gates (§5.1, §5.4).

Everything below is either a citation to a paper I fetched today or a number my own gates printed. Where neither exists I say so in those words.

---

0. The papers, and the citation correction

what	id	verified how
Hyper-Connections (Defa Zhu, Hongzhi Huang, Zihao Huang, Yutao Zeng, Yunyao Mao, Banggu Wu, Qiyang Min, Xun Zhou — Seed-Foundation-Model Team, ByteDance; ICLR 2025)	arXiv:2409.19606 (v1 29 Sep 2024, v2 28 Nov 2024, v3 18 Mar 2025)	fetched arxiv.org/abs/2409.19606 and the full HTML of v3 on this revision; §2 read in full including Tab. 1-3 and App. B.
mHC: Manifold-Constrained Hyper-Connections (Xie, Wei, Cao, … Liang; DeepSeek-AI)	arXiv:2512.24880 (v1 31 Dec 2025, v2 5 Jan 2026)	fetched arxiv.org/abs/2512.24880 and the full HTML of v2. 20 authors, cs.CL. The crate’s citation is correct.
How Much Is One Recurrence Worth? Iso-Depth Scaling Laws for Looped LMs (Schwethelm, Rückert, Kaissis)	arXiv:2604.21106 v3 (7 May 2026)	fetched. The φ numbers quoted in the brief are verbatim from its abstract.

The brief’s anchor for the base mechanism was the wrong paper. arXiv:2409.05799 is “PDAF: A Phonetic Debiasing Attention Framework For Speaker Verification” (cs.SD, Massa Baali et al., accepted to SLT). It is not Hyper-Connections and has nothing to do with residual streams. The real one is arXiv:2409.19606 — re-resolved on this revision through the arXiv API and by reading v3’s §2, not by pattern-matching the digits. The misattribution also reached .bulba/goal.md’s arm-MHC line; goal.md on main already carries the owner’s correction and is not edited from this lane. grep -rn 2409.05799 over the tree now matches exactly one place — this paragraph, which is two grep lines because the id is split across them. Every other citation in this file, in docs/protocols/AB-PROTOCOL.md row 7b and in the schema.rs / loop_block.rs doc comments names 2409.19606.

The motivation, checked against the source rather than the summary. 2604.21106v3 abstract: “we measure a recurrence-equivalence exponent φ = 0.46 … Truncated backpropagation lowers φ to 0.38 … Conversely, hyperconnections raise φ to 0.65, a genuine capacity gain.” Table 3 gives the refits: baseline φ = 0.453, truncated BPTT 0.380, hyperconnections 0.646. The claim is correct as the brief states it. Three qualifiers that matter more than the number, all from the same paper:

1. It is K = 2 lanes, at the LOOP BOUNDARY. §5.2: “We replace the linear input injection of our baseline looped model with hyperconnections, which scale and mix K parallel residual lanes across loops. We use K = 2 lanes and full BPTT.” That is the placement decision in §3 below, and it is the paper’s, not mine.
2. It is unconstrained HC, not mHC. The case study cites [12] = 2409.19606. DeepSeek’s mHC is a different member of the same family (the constrained one, 2512.24880), published two months earlier and never measured for φ by anyone. No published number connects mHC to φ. The transfer is an assumption and the A/B row says so.
3. Their regime is not ours. 116 runs, 20 effective layers, d_model = 64s, FineWeb-Edu tokens, MuonH, prelude-recur-coda with (ℓ_prelude, ℓ_coda) = (2,2) so N_once > 0. We are fully looped (N_once ≈ 0), 9.2M parameters, 2-4 recurrences, bytes, vocab 256. The paper’s own Limitations section: “the link between φ and reasoning quality at scale is therefore empirically untested.”

---

1. The base mechanism (2409.19606), and whether the crate implements it

The brief asked for this explicitly. Short answer: no, and the deltas are the point — burn-mhc is a transcription of DeepSeek’s constrained member of the family (2512.24880), and the constraint plus two layout choices are what make it a different operator from the base paper’s. The Birkhoff projection is not our invention and is not “our own addition on top” — it is a third-party published mechanism with its own paper, which the crate’s header already cites correctly. What is ours is the wiring decision (§3) and the static-bias init (§2.1).

1.1 What 2409.19606 actually specifies

Its §2 is short and the mechanism is exact, so there is no reading latitude:

• §2.1: H⁰ = (h⁰ … h⁰)ᵀ ∈ R^{n×d} — the input is d-dim and is REPLICATED n times to form the hyper-hidden matrix. n is the “expansion rate”.
• Eq. 1: HC = [[0, B], [A_m, A_r]] ∈ R^{(n+1)×(n+1)}, with B ∈ R^{1×n}, A_m ∈ R^{n×1}, A_r ∈ R^{n×n}.
• Eq. 2/3/4/5, the whole operator: h₀ᵀ = A_mᵀ H (the n streams are mixed DOWN to one d-vector — this is the layer’s input), H' = A_rᵀ H (the width-connections, A_r mixes the n streams), Ĥ = Bᵀ(T(h₀))ᵀ + H'.
• Eq. 6/7: the depth-connections are DC = [B; diag(A_r)] — the residual equivalent — and the width-connections are WC = [A_m A_r].
• §2.2, DHC: H̄ = norm(H), then B(H) = s_β ∘ tanh(H̄W_β)ᵀ + B, A_m(H) = s_α ∘ tanh(H̄W_m) + A_m, A_r(H) = s_α ∘ tanh(H̄W_r) + A_r.
• §2.3 Eq. 14 init: B^k = 1_{1×n}, A_m^k = e_{k mod n}, A_r^k = I.
• The final readout SUMS the n streams (“we sum the last hyper hidden matrix row-wise to obtain the required hidden vector”).
• So: the layer body T runs at width d, the residual stream is n·d wide, and A_m is the term that collapses the stream back to d before the body. A_r is what mixes streams, unconstrained, in [−∞, ∞].

1.2 The deltas, one per row

	2409.19606 (HC/DHC)	burn-mhc (mHC, 2512.24880)	file
streams	n replicated copies of one d-vector, total width n·d	n disjoint slices of one D-vector, total width D	block.rs:180,190
body width	d, fed A_mᵀ H	D, fed the whole hidden (the caller’s body is built at D)	§2.2 below
readout	sum the n streams back to d	concatenate (disjoint slices reassemble to D)	block.rs:199
width-connections	A_r free, A_m mixes down	H_res = Sinkhorn(H̃_res), doubly stochastic; H_pre = the A_m role	Eq. 8-9
depth-connections	B free, no factor	H_post = 2σ(H̃_post), non-negative, bounded in [0,2]	Eq. 8
dynamic part	s ∘ tanh(H̄ W) + static	α (H̄ φ) + static, no tanh	§2.1 row 3
init	A_m^k = e_{k mod n} (a per-layer rotation)	b_pre = b_res = 10, b_post = 0 (uniform)	§2.1 last row
per-iteration matrices	one per layer (or per sublayer)	one shared block at the loop boundary	§3.2

Four of these are worth a sentence each, because three of them cut in a direction a reader might assume goes the other way.

Slices, not replication — and it is forced, not chosen for expressiveness. §3 of mHC builds x_l by repeating the C-dim state n times, so H_res mixes n copies of one vector: a pure re-weighting, and A_m has to mix the copies back down to C before the body. Slicing a D-dim hidden into n disjoint D/n chunks gives the same shapes with a different function — ours mixes n quarters of a real hidden state, a genuine cross-stream exchange. The reason it is forced here: our body is built at width D and a quarter-width body is a different model (that is exactly why H_pre is not wired, §3.3). Slicing is the only variant that composes with a full-width body, and the concat readout keeps D so nothing has to be re-learned at the far end.

Concat, not sum — the layout choice that makes the two consistent. If we had replicated, the base paper’s “sum the streams” readout would multiply the readout’s input by n, so the final projector would have to be retuned (their §4 does exactly that: scale the output-module stds by √n). Slicing makes the readout width-invariant, so our out_proj is untouched — which is one more thing the arms share, and one fewer confound.

No tanh is now supported by the BASE paper’s own ablation, not just by the printed equations. §2.1’s row 3 argued this from mHC’s Eq. 7 dropping it. The stronger evidence is 2409.19606 Tab. 1, where removing tanh is the better variant at every rung on OLMo-1B: V2 eval loss 2.792 vs 2.802 at n = 2 and 2.779 vs 2.781 at n = 4 (and Fig. 5’s right subfigure ablates it separately). So the crate’s unbounded linear φ_* is a transcription the source paper’s own measurements endorse, and the log-domain Sinkhorn it forces is a consequence worth paying rather than a bug.

The base paper’s single most important ablation is the axis we changed most. Tab. 3, on DHC×4: freezing the width-connections WC = [A_m A_r] costs +0.021 V2 / +0.017 V3 loss, while freezing the depth-connections B costs +0.002. We replace A_r with the Birkhoff projection — a hard constraint on exactly the matrix the base paper’s largest ablation is about — and we do not wire A_m at all. Row 7b therefore does not measure the base paper’s effect; it measures a family member with one of its two load-bearing blocks changed and the other dropped. That is a legitimate experiment, and it is the one the φ evidence motivates, but it is not “Hyper-Connections” and the A/B row now says so.

1.3 The base paper’s own expansion-rate ablation is why n = 2 is first

2409.19606 Tab. 1, OLMo-1B, 500 B tokens, V2 eval loss (lower is better):

n	1	2	4	8	baseline (no HC)
DHC	2.819	2.802	2.781	2.778	2.811

Three things follow, and all three point the same way:

1. n = 1 is WORSE than the residual baseline (2.819 vs 2.811). Their App. F says why: the seesaw effect persists at n = 1. This is also why this crate’s MhcBlock::new(0, …) no longer silently trains the n = 1 arm — a request for zero streams is now a loud refusal, because the n = 1 rung is the one the family is known to lose at.
2. Most of the gain arrives by n = 2 (2.811 → 2.802, 0.009 of the 0.030 total) and n = 4 → 8 buys 0.003 for double the stream width. The paper’s own words: “the increase to n = 8 provides minimal additional benefits.” That is the base paper’s version of 2607.14530’s collapse past N = 4, and it is why the external review’s “n = 2 first” is not a compromise between the papers and the box — it is what the base paper’s own table supports.
3. We are at T = 2-4 recurrences, not 16-20 layers. Their n buys width the depth-connection matrix can route around; ours has 2-4 slots to route around, so the argument for a wide stream is weaker here, not stronger.

1.4 The manifold constraint, in two sentences

mHC takes Hyper-Connections’ n-stream residual stream and projects the per-token stream-mixing matrix H_res onto the Birkhoff polytope — the set of doubly stochastic n × n matrices — with 20 iterations of Sinkhorn-Knopp (Eq. 9). That makes H_res x_l a convex combination of the streams, so the forward signal’s mean is conserved and ‖H_res‖₂ ≤ 1 bounds the backward gain; because the set is closed under multiplication, the composite ∏ᵢ H_res^i over depth is doubly stochastic too, which is the identity-mapping property plain HC destroys (HC’s own composite reaches Amax gain ≈ 3000 on their 27B, mHC’s ≈ 1.6).

H_pre = σ(·) and H_post = 2σ(·) (Eq. 8) are the non-negativity half of the constraint — the paper calls it “a special manifold projection” — and they are what stop a branch from being cancelled by a negative coefficient. mHC’s measured evidence for the constraint is stability at 27B (Fig. 2’s loss surge and gradient-norm spike for HC at ~12k steps), not a quality win over HC: the paper’s own Tab. 4 has mHC ahead of HC on 7 of 8 benchmarks and behind on MATH (26.0 vs 26.4).

---

2. The crate, read in full (1 332 lines, 4 files)

vendor/dormouse-fused/crates/burn-mhc was 1 011 LOC with zero call sites in the model crates (its own README says so; docs/architecture/library-crate-fate.md calls it a reference port). What it is: MhcBlock (Eq. 7 hyper-network + Eq. 8 constraints + Eq. 3 residual), a fused CUDA Sinkhorn forward, and a hand-derived exact backward for it. There is an audit trail on it already (e884bc6 withdrew its “verified” claims, 6e9624c fixed a NoCheckpointing-only seam, 63e01c3 retracted its speedups).

2.1 Formula map, every equation ↔ file:line

n = n_branches (streams), C = D/n, D = d_model.

paper	crate	verdict
Eq. 3 x' = H_res x + H_postᵀ F(H_pre x, W)	block.rs:169-194 (forward) — the H_res x and H_postᵀ F terms; H_pre x is the caller’s	partial, deliberately. See §3.
Eq. 5 (HC) H = α tanh(θx̃ᵀ) + b	—	not used, and deliberately: mHC’s Eq. 7 and its kernel section (Eq. 14-16) both drop tanh for a linear x̃'φ. The prose says “we follow the original HC formulation”; the equations do not have tanh in them. The crate follows the printed equations and now says so at the top of lib.rs. Two pieces of external support, both from 2409.19606v3: the base paper’s “W/O tanh” ablation rows are neutral-to-better (Tab. 1: DHC×4 2.781 with tanh, 2.779 without; Tab. 2 same at n=2: 2.802 → 2.792), so dropping tanh is not a loss the base paper’s own measurements can see. Consequence still named: φ_* are unbounded linear maps where HC’s are tanh-squashed, which is why the Sinkhorn input is unbounded and why sinkhorn.rs must run in the log domain at all.
Eq. 6 the Birkhoff polytope H 1 = 1, 1ᵀH = 1ᵀ, H ⪰ 0	sinkhorn.rs:22-71 (log-domain alternating normalization)	✅
Eq. 7 x̃' = RMSNorm(vec x), H̃ = α(x̃'φ) + b, mat(·) reshape	block.rs:76-85 (x_norm), :94-116 (post/res), :125-154 (all three)	✅ including the R×n² reshape. x_norm is h / sqrt(mean(h²) + 1e-20) = Eq. 15’s 1/r = √(nC)/‖x⃗‖ with the paper’s Layer Norm ε = 1e-20 (App. A.1 Table 5). That epsilon is the paper’s, not the crate author’s invention.
Eq. 8 H_pre = σ(H̃_pre), H_post = 2σ(H̃_post), H_res = SK(H̃_res)	block.rs:113-116, :156-160	✅ three for three, factor 2 included
Eq. 9 M⁽⁰⁾ = exp(H̃_res), M⁽ᵗ⁾ = T_r(T_c(M⁽ᵗ⁻¹⁾)), t_max = 20	sinkhorn.rs:47-70; sinkhorn_cuda.rs:69-119	✅ both forms, SINKHORN_ITERS = 20 (sinkhorn.rs:6)
App. A.1 Table 5: n = 4, α init = 0.01, t_max = 20, ε = 1e-20	ALPHA_INIT (block.rs:6), SINKHORN_ITERS, x_norm’s 1e-20; n becomes our mhc_streams	✅ all four
static-bias init b_pre = b_res = 10, b_post = 0	block.rs:8 (STATIC_INIT), :55-57	OURS, transcription. The paper states no bias init. 10 is chosen so σ(10) ≈ 1 and SK(exp(diag 10)) ≈ I. Named in the crate header as of this commit. The idea is HC’s: 2409.19606 Eq. 14 initialises the static matrices to B = 1, A_m = e_{k mod n}, A_r = I precisely so that HC at step 0 is the Pre-Norm residual connection. mHC’s A_r = I is the same move through the Sinkhorn.
dynamic-projection init	block.rs:51-61 Normal(0, 0.01)	OURS, and it differs from HC. 2409.19606 §2.2-2.3: “To stabilize the training process, we introduce normalization before the linear transformation …” and §2.3 initialises W_β, W_m, W_r to zero, i.e. the dynamic half is switched off at step 0. burn-mhc draws them from N(0, 0.01), so its dynamic half is on. With α = 0.01 the contribution is O(10⁻²) so it is a small deviation — and the measured gradients (§5.3b) are consistent with it being small — but it is a transcription difference from the base paper, not from mHC (which says nothing about it), and it is named here rather than left to be found in a gradient dump.

2.2 Three structural differences from the paper, named

(§1.2 has the full delta table against the base paper 2409.19606; this is the crate-internal view of the two that change the operator’s shape.)

• Replicated streams → disjoint slices. §3 of the paper builds x_l by repeating the C-dim state n times: x_l = (x_{l,0}ᵀ … x_{l,n-1}ᵀ)ᵀ. The crate instead slices the D-dim hidden into n disjoint chunks of D/n (block.rs:180, :190). Same tensor shapes, different function: the paper’s H_res mixes n copies of one vector (a pure re-weighting), ours mixes n quarters of a real hidden state (a genuine cross-stream exchange). Ours is the more expressive operator and is the only one that can compose with a body that runs at full width D. It also means the “identity at init” claim is about the operator, not about the stream layout.
• The body is not per-stream. In the paper F is one function mapping C → C. Ours is one function mapping D → D, so the crate reads its output as n per-stream outputs and weights each by H_post_j. The [1,n]-shaped H_post broadcasting over [b,t,n,C] is what makes that work.
• Sum → concat at the readout. The base paper (2409.19606 §2.1) collapses the n streams back to one d-vector by summing them, and compensates in §4 by scaling every output module’s std by √n. We concatenate: the disjoint slices reassemble to exactly D, so out_proj is untouched and needs no rescaling. This is a consequence of slicing rather than replicating, not an independent choice, and it is one fewer difference between the two arms in row 7b.

---

3. PLACEMENT, decided explicitly

3.1 Where in the chain

ReZero (default)  |  use_gr -> Gated Residual  |  use_attnres -> AttnRes  |  use_mhc -> mHC

All four replace one statement — loop_block.rs’s residual write, h = h_ctx + y·scale — so any two are refused together in config::validate, now by a single three-arm check (validation.rs:85-115) so a future fourth arm cannot leave a pair unguarded. The branch is loop_block.rs:571-604; the forward_full_state assert at :319-324 re-checks it for a hand-built block that bypassed the config path.

3.2 Per-iteration connections, ONE shared block — and why

LoopBlock::mhc: Option<MhcBlock> is a single block applied at every iteration’s write, not Vec<MhcBlock> of length max_iter. Three reasons, in descending order of how much they would cost if I got it wrong:

1. The manifold’s argument is about the COMPOSITE. 2512.24880 Eq. 4 and §4.1 are entirely about ∏ᵢ H_res^i across depth. A shared block makes the loop’s composite H_res^T — a power of a doubly stochastic matrix, still in the Birkhoff polytope, because that set is closed under multiplication. So the constraint’s central claim survives our depth T = 2-4 exactly, and mhc_res_leaves_the_identity_when_the_bias_moves_and_stays_on_the_manifold asserts the closure on the real 4×4 composite.

2. Weight sharing is the loop’s defining property. Per-slot residual matrices would make the residual operator the only un-shared thing in a block that shares its controller, attention, memory, experts, norm and readout. iter_embed and residual_scale are shared for the same reason.

3. The connection count has to survive T small — and n is cheap in parameters and irrelevant to launches, which is the whole latency story. mhc_streams defaults to 2, changed from 4 on the external review of 2026-10-01 (“начать с n=2, одной вариацией за запуск”; and at 9.2M 2605.09165/2510.26692 answered “недостаточно свидетельств”), because K = 2 is exactly what the φ study ran (§0) and exactly what 2409.19606 Tab. 1 supports (§1.3). The crate’s n is a constructor argument — MhcBlock::new(n_branches, d_model, device) — so this was a one-flag change and not a fork; --set mhc_streams=4 is the follow-up row. The three papers that put the knee at 4 are still the right reason not to climb:

   • the base paper measures the ladder (2409.19606v3 §4.1, Tab. 1, OLMo-1B, 500 B tokens, V2 eval loss): baseline 2.811, DHC×1 2.819 (worse than the baseline — n = 1 is the one rung the family loses at), DHC×2 2.802, DHC×4 2.781, DHC×8 2.778. So n = 4 buys 0.021 and n = 8 buys 0.003 more.
   • mHC itself ships n = 4 (2512.24880 App. A.1 Table 5) — which is why this is a deliberate departure from mHC’s own default and not an oversight.
   • 2607.14530 (xHC) finds the family’s gains collapse past N = 4, with the two causes named: “insufficient write-back information for an expanding number of streams” and residual-mixing generation that “cost scales cubically with N”.

   At n = 2 the block costs 2·D·n + D·n² + 2n + n² + 3 = 6 155 parameters on small — measured on the instantiated block at d_model = 768, 0.067 % of small’s 9 197 390, and 0.067 % of nothing that matters. (An earlier draft of this file said 3·D·n + D·n² + n² + 5n = 2 306 “at D = 384”; the formula was wrong by the phi_pre term and the width was wrong because small is d_model = 768, not 384 — §3.2’s own numbers, caught by the same gate that now asserts the count against the built block.) So parameters were never the binding constraint at any n. What binds is the mixing, and n = 1 is the published proof that the mixing is the whole point: at n = 1 this arm is a worse ReZero with a Sinkhorn attached. That is also why MhcBlock::new refuses n = 0 loudly rather than clamping it to 1 (see §4.9), and why config::validate refuses mhc_streams = 0 for the same reason.

   The cost note is NOT an argument for a larger n, and I got this wrong before the review. I first wrote in schema.rs that “n is a launch-count knob”. It is not. sinkhorn_knopp issues a fixed SINKHORN_ITERS × 2 directions × 7 ops = 280 forward launches per call (sinkhorn.rs:48-70), plus a 40-step hand-derived reverse unroll in the backward, and that count does not depend on n at all — n scales the element count of each tiny [b, t, n, n] tensor and the parameter count, never the number of kernels the host enqueues. So a ladder in n will not buy the latency back; only the fused kernel would, and it is deliberately off (§6). The consequence for the experiment is in §6 and in row 7b: step time is priced as a separate verdict from BPB, and a smaller n is not the lever that would make it go away.

The loop boundary, not between sublayers. This is the paper’s placement, not mine: 2604.21106v3 §5.2 puts the lanes across loops, and cites 2604.21254 (“loop-level hyperconnections beat layer-level”). Our iterations are the loops.

3.3 H_pre is not wired, and the cost is 11 %

Eq. 3 is x' = H_res x + H_postᵀ F(H_pre x, W), and H_pre x is [1,n]·[n,C] = [1,C] — it feeds the block’s own function F, collapsing the n × C stream to C. Our body is built at width D; running it on a C = D/4-dimensional input means building a second, quarter-width block, which is a different model and a different experiment. The paper’s own Tab. 1 says what that costs:

arms enabled	absolute loss gap
none	0.0
H_res only	−0.022
H_res + H_pre	−0.025
all three	−0.027

H_res is 81 % of the total; H_pre adds 11 % and H_post 7 %. We wire H_res + H_post = 89 % of the measured effect and leave the 11 % to a decision that needs a quarter-width block.

The cost of that choice is concrete and asserted rather than absorbed: phi_pre ([D, n]) plus its gate alpha_pre and static bias b_pre are allocated and never read — 1 540 + 1 + 2 = 1 543 parameters on small at n = 2, 0.017 % of the model, and they are not in the graph at all, so grad() returns None for them rather than zeros. mhc_gradients_reach_every_wired_parameter asserts exactly that for all three, with a message that fires if something ever starts using H_pre.

3.4 The base state is h_ctx, not h

ReZero writes h = h_ctx + y·s where h_ctx = h + e_k. The mHC write uses h_ctx as x_l for the same reason: the A/B against ReZero must differ in the residual operator and in nothing else. Using h would move the iteration embedding from the write to the read side and confound the one comparison the arm exists to make. mhc_and_rezero_each_move_only_their_own_write asserts both directions of that separation, including that residual_scale is inert under mHC.

---

4. Defect classes hunted in burn-mhc, proactively

The two classes the paper-faithful AttnRes wiring found an hour before this (511daa5, AGENTS §2.2): a derivative factor that vanishes under a norm and an absolute-tolerance gate on a value that moved scale. Plus the burn-rmsnorm class: a test target that compiles to nothing.

4.1 FOUND — a SILENT 1e-7 floor in the fused kernel, untested, in the same

class as the defect the log-domain rewrite exists to fix

sinkhorn_cuda.rs:89-97 (row) and :105-111 (column) clamp the normalization sum to 1e-7 in the linear domain, with no counter and no escape:

let mut rn = rs;
if rn < F::new(1e-7_f32) { rn = F::new(1e-7_f32); }

The tensor path (sinkhorn.rs:47-70) runs the same normalization in the log domain and has no floor at all. Its own module doc (sinkhorn.rs:13-21) says why: entries above exp(88) overflow f32, and bounded inputs whose within-row range exceeds f32’s exponent span flush to zero, “after which a zero column sum makes the column pass emit inf and the next pass NaN”. So the kernel reinstates in linear space the exact failure the tensor path was rewritten to remove: a row whose true sum is 1e-20 is divided by 1e-7 and comes out 1e13 times too large, where the tensor path returns the 1e-20 and converges.

• Not reachable from MhcBlock at init (b_res = 10I puts the diagonal at e^10), so it is not a training hazard at step 0.
• No test covers it. sinkhorn_large_logits_stay_finite (lib.rs:93-125) exercises the non-autodiff path only — the inventory doc already named the autodiff fallback’s untested overflow; this is the kernel floor and it is a third instance.
• ADR-0011 mark: it carries SILENT, and SILENT is a defect. The fix is COUNTED (a counter next to FWD/BWD) or the row is dropped in favour of the log-domain formulation the tensor path already has.
• Not fixed in this commit: the kernel cannot be run or measured on this lane’s card, and a fix I cannot test is a fix I cannot claim. Named at the line (sinkhorn_cuda.rs:89-99) so the next reader meets it there, and it is the first precondition of the CUDA follow-up in §7.

4.2 FOUND — an absolute tolerance on a quantity whose scale is a free

parameter

ad_tests::sinkhorn_fused_backward_matches_tensor (sinkhorn_cuda.rs:711-732 after this commit) asserts maxdiff < 1e-3 absolute on dL/dlogits. That quantity’s scale is set by n (it multiplies exp(logits) and the 1/n a doubly stochastic output carries) and by iters (2·iters normalizations each dividing by a sum ≈ 1). The same kernel is exactly as correct at n = 16 as at n = 8 and fails the same number. This is the merge_state_writeback_matches_host_reference defect from 84844a6 — a gate whose verdict depends on the scale of the number it is looking at. The fix is a relative tolerance, and the gate I added for the fused node (§4.4) states it as a relative requirement rather than guessing an absolute one. Not changed here: it is a CUDA-only test I cannot run, and changing an unrunnable gate’s tolerance to a number I also cannot measure would be the same sin.

Also in this class, smaller: n1_identity_at_init (lib.rs:128-139) asserts mse < 1e-3 on h + branch with h, branch ~ N(0,1) — an absolute bound whose meaning moves with the input distribution. It is a loose gate on a claim (“starts as the identity residual”) that is true by construction, so it is informational. Noted, not changed.

4.3 FOUND — the gradient gate this crate never had, which is 8fa5d4c

seam_tests::balanced_checkpointing_reaches_the_seam_and_the_legacy_entry_does_not (sinkhorn_cuda.rs:525-562) asserts reach::<BalancedCheckpointing>(&x) .is_some(). Reachability, not trainedness. That is the exact gap 8fa5d4c fell into: burn-gdn2’s hand-rolled node came back UnTracked under BalancedCheckpointing, the arms ran thousands of forwards, the loss curve was healthy, and the attention arm sat at initialisation for the whole history of that project. A is_some() cannot tell a tracked node from a leaf.

Fixed as far as a CPU-only lane can: the_tensor_path_carries_a_gradient_under_ both_strategies asserts a non-zero, finite gradient on the tensor path under both strategies, and deliberately uses a non-constant fixture (the existing test’s ones input makes the sinkhorn’s output constant, so dL/dlogits is legitimately zero and any such gate is green for the wrong reason — the constant-fixture cousin of the defect 511daa5 documents).

What that does not close, stated plainly: the fused node lives behind cfg(all(cuda, autodiff)) and cannot be reached without a card. fused_node_carries_a_gradient is written as an #[ignore]d, CUDA-gated test carrying its four required clauses as documentation. Running it is owed, and it is a one-line command.

4.4 FOUND — a gate that could not run in the configuration it covers

fd_tests::sinkhorn_grad_matches_finite_difference is behind #[cfg(all(test, feature = "cuda"))], but nothing in it touches CUDA — it compares an analytic gradient against central differences, both on whatever Device::default() resolves to. It is a CPU test wearing a GPU gate, so on a CPU-only build the crate’s only independent check of a hand-derived backward (there are three such backwards in this repo: gdn2, attnres, this one) does not run. The honest reading of 34c5631: BURN_DEVICE appearing nowhere under a crate is a symptom, and so is a cfg wider than the code needs.

Not moved in this commit. Moving it means choosing a backend explicitly (Device::ndarray() instead of Device::default()), which changes what the existing CUDA-gated run measures, and I cannot run the CUDA half to confirm the change is behaviour-preserving. The gate is cargo test -p burn-mhc --features cuda,autodiff today; the fix is one cfg and one Device:: and it belongs to whoever next has the card. Recorded here instead of silently leaving it.

4.5 CHECKED, CLEAN — the derivative factors under the norm

The AttnRes defect was a hand-written adjoint missing a factor that vanished at the L2 norm, so every test stayed green. In this crate:

• The RMSNorm’s derivative is burn’s, not ours. x_norm (block.rs:76-85) is plain tensor arithmetic, so ∂(h/‖h‖)/∂h is autocomputed. There is no hand-written factor to drop, and the 1e-20 is the paper’s ε, not a fudge that could hide one.
• The hand-written adjoint IS the Sinkhorn’s, and it is correct. For m_post = m_pre / s, s = Σⱼ m_preⱼ: ∂m_postᵢ/∂m_preⱼ = δᵢⱼ/s − m_preᵢ/s², so ∂L/∂m_preⱼ = dⱼ/s − (Σᵢ dᵢ m_preᵢ)/s². Both sinkhorn_backward_tensor (sinkhorn_cuda.rs:480-506) and the fused sinkhorn_backward_kernel (:168-196) compute exactly that, recover m_pre = m_post·s rather than storing it, and chain to the logits with the · exp(logits) Jacobian of M⁽⁰⁾ = exp(H̃_res). I re-derived all three pieces and found no missing factor, no sign error and no wrong axis (the tensor version’s if j % 2 == 0 { 3 } else { 2 } correctly alternates row/ column with the unroll index, and it is unrolled in reverse).
• No norm in the Sinkhorn path, so the “vanishes under a norm” shape has nowhere to hide.

4.6 CHECKED — the two references are not comparing our own formula to itself

The 511daa5 lesson proper: every reference in fused_attnres.rs re-derived our own d^-0.5, so “kernel matches tensor” compared two transcriptions of the same wrong equation. Here the comparison is sinkhorn_kernel vs sinkhorn_knopp — two implementations of Eq. 9, one linear-domain fused and one log-domain, so agreement is real evidence about the two implementations and silent about whether either matches the paper. That is why the Eq. 3 fixture (§5.1) is hand-computed from the paper’s equations and not from the crate’s own tensor path, and why it is pinned to literals. No external reference exists for the numerics: DeepSeek’s own third-party/tilelang_ops/ref_mhc.py inside DeepGEMM is the closest thing and it cannot run here (SM90/SM100 only, per docs/research/2026-09-27-adopt-vs-port.md).

4.7 CHECKED, CLEAN — reachability of the fused path

sinkhorn.rs:23-45 probes sinkhorn_autodiff_s under both NoCheckpointing and BalancedCheckpointing, and the downcast compares the inner backend’s TypeId, which is the question the right way round (6e9624c fixed the first half of this). The ENTRY counter is incremented after the downcast, so it counts arrivals rather than interest (the f737710 fix). So the seam would engage on the trainer’s backend if burn-mhc/cuda were enabled — which, per §7, it deliberately is not yet.

4.8 A correction owed to an existing document

docs/archive/research/2026-09-27-fused-inventory-attention.md:206 records arXiv 2512.24880 … **VERIFIED**. The paper is real and the ID is right (so the inventory was not wrong about the citation), but “VERIFIED” with no reference named is the ADR-0020 shape, and the crate’s own README has since retracted the numerics to “no external reference exists — all three compare this crate against its own tensor path”. Those two statements disagree. The README is the correct one. Not edited here: that file is a dated research artifact and the correction belongs in its own lane.

4.9 A SECOND WRITER IN THIS WORKTREE — reviewed, kept, reconciled

At 01:01 on 2026-10-01, while this lane was queued on the build lock, an edit appeared in wt/mhc-model that the first pass did not make: burn-mhc/src/block.rs gained assert!(n_branches >= 1, …) in place of let n = n_branches.max(1);, with a doc comment citing “2409.19606 Tab. 1; n = 1 is measured there as below the Pre-Norm baseline”. A vendor/dormouse-fused/target/ build directory appeared in the same worktree at the same time.

Reviewed, and the change is correct and better than what it replaced. The silent max(1) is the ADR-0019 class: --set mhc_streams=0 would have trained n = 1, and 2409.19606v3 §4.1 Tab. 1 measures DHC×1 at V2 loss 2.819 against the baseline’s 2.811 — worse than the residual connection it replaced, with a healthy-looking loss curve. The citation was verified against the paper before being accepted, because a comment that cites a table wrongly is the same defect as a retracted claim. The first pass’s own config::validate already refused mhc_streams == 0 with a message; the assert! is the library-boundary half of the same rule and belongs there. Kept, and the 2026-10-01 pass re-verified it green (mhc_counts_every_iteration_and_is_refused_alongside_the_others now asserts the n = 0 refusal at both layers, and mhc_res_is_doubly_stochastic… reads n from the block rather than hardcoding a literal, so the gate follows mhc_streams).

What was not done: revert it, or commit it silently. The owner should know that two writers were in wt/mhc-model at once — the worktree discipline in ADR-0022 exists to prevent exactly the interleaving that made this ambiguous, and a find -newermt was the only thing that caught it. The file and the code are now reconciled (this pass owns the citation, the n default, the §5.1a and §5.1b corrections and the commit), but the collision is on the record rather than explained away, because a findings file that hides a concurrent writer is the exact instrument this lane is arguing against.

---

5. Gates, red → green, with the magnitudes

5.1 (a) FORM FIDELITY — three gates, because “the readout moved” is the

gate that let a d^-0.5 scale through

eq3_residual_is_the_papers_equation_on_a_hand_computed_fixture (burn-mhc/src/lib.rs). n = 2, C = 2, D = 4, so the whole residual is four numbers a human can check. α_* = 0 kills the dynamic part, leaving H_post = 2σ(b_post) and H_res = SK(b_res) exactly.

b_res = [[0, −ln4], [−ln4, 0]] → M = exp(b_res) = [[1, ¼], [¼, 1]]. Sinkhorn’s alternating normalization preserves the cross ratio M₀₀M₁₁/(M₀₁M₁₀) = 16, and a 2×2 doubly stochastic [[p,1−p],[q,1−q]] has cross ratio p(1−q)/((1−p)q), which is 16 exactly at p = q = 4/5. 20 iterations reach [[0.8,0.2],[0.2,0.8]] to the last bit of f32 (verified numerically before writing the test, so the expected value is the exact literal 0.8, not a tolerance band). With b_post = [−0.5, 1.5] → H_post = [0.75508134, 1.63514895], Eq. 3 gives, for stream j and column c, out[j,c] = Σᵢ H_res[j,i]·x[i,c] + H_post[j]·y[j,c] with x = [[1,2],[3,4]] and y = [[10,20],[30,40]]. All four numbers, in the row-major order reshape([b,t,n,C]) returns:

	c = 0	c = 1
j = 0	0.8·1 + 0.2·3 + 0.75508134·10 = 8.95081338	0.8·2 + 0.2·4 + 0.75508134·20 = 17.50162675
j = 1	0.2·1 + 0.8·3 + 1.63514895·30 = 51.65446857	0.2·2 + 0.8·4 + 1.63514895·40 = 69.00595810

5.1a A GATE THAT WAS NEVER RUN, WITH TWO OF ITS FOUR EXPECTED VALUES WRONG

This section is the most important one in the file, and it is about me. eq3_residual_is_the_papers_equation_on_a_hand_computed_fixture had a red→green table entry (this file’s §5.4, rows 1-2) claiming it was “demonstrated by applying the perturbation, recording the red reading, and restoring”. It was never executed. On this revision it was run for the first time and it failed — and the failure was in the test, not the operator:

1. Tensor::<1>::from_floats([b_post], …) on a [f32; 2] builds dims [1, 2] for a rank-1 tensor and dies in from_data with “Given dimensions differ from the tensor rank”. The gate had never got past its own fixture. Fixed with TensorData::new(b_post.to_vec(), [2]).
2. With that fixed, out[0] and out[1] passed and out[2] failed: crate 51.654465, expected 18.951490. The crate was right. The expected values I had written used y’s stream-0 element for stream 1 — the exact stream-indexing slip the gate exists to catch, made in the gate. Correct literals are 51.65446857 and 69.00595810, and I re-derived all four independently (hand arithmetic, then a separate script) before touching the test.

The separations, recomputed from the same source as the literals, on out[0,0]:

reading of Eq. 3	out[0,0]	gap	relative
Eq. 3 as printed	8.950813	—	—
H_post = σ(·), Eq. 8’s factor 2 dropped	5.175407	−3.775406	42.2 %
H_res = b_res, Eq. 9’s Sinkhorn dropped	3.391930	−5.558883	62.1 %
H_post indexed off by one stream	17.751490	+8.800677	98.3 %

42 %, 62 % and 98 % on an output of order 10 — two orders of magnitude above any tolerance, and a “did the output change” gate is green under all four. The test also asserts the projection itself (H_res = [[0.8,0.2],[0.2,0.8]]), not only its effect, because the projection is the mechanism.

One row of the original table is deleted rather than fixed. It claimed H_postᵀ read as H_post gives 10.150813, a 13.4 % separation. It does not: H_post is [1, n], so transposing it is a no-op, and recomputing gives 8.950813 — a 0.0 % separation. That row was a plausible-sounding number that does not exist, and it is in this file’s own record; the replacement row (the 98.3 % off-by-one) is a real one, and it is the mistake the test actually made.

What this fixture provably cannot catch, now stated instead of implied: a transposed H_res. The b_res is symmetric, so H_resᵀ = H_res and the wrong reading is the right answer. Catching it needs a non-symmetric b_res, which is a second fixture, not a tweak of this one. Named so nobody reads the remaining 42/62/98 % coverages as complete.

The generalisable rule, in the AttnRes/511daa5 family: a gate whose expected values are literals is a gate that must have been observed to go red. A red→green table row is a claim about a run. If the run is not in evidence, the row is a story — and this one was in a file whose whole subject is the cost of stories that look like measurements.

mhc_res_is_doubly_stochastic_in_both_directions (model level). The operator’s defining property, in both directions, as the paper’s own Amax Gain Magnitude metric: forward max|rowsum − 1| and backward max|colsum − 1|. Different bounds on purpose — the last normalization of Eq. 9 is a row pass, so the row sums are exact to fp and the column sums are the unconverged direction, which is what the paper’s Fig. 7(a) measures. Non-negativity (Eq. 6) too. This gate cannot be written for ReZero (a scalar has no matrix) or AttnRes (a softmax over sources has no stream mixing), which is the point: it tests the operator CLASS, not the output.

mhc_res_leaves_the_identity_when_the_bias_moves_and_stays_on_the_manifold (the one that matters). Two ends of the trajectory, because the middle is where it fails:

• at the paper’s init b_res = 10I, H_res is the identity to 4.5e-5 — and that is the point of the constraint (§4.1, “restores the identity mapping property”), not a defect, so it is asserted, against the predicted (n−1)·e^{−10} = 4.54e-5 rather than a tolerance band;
• perturbed off the init, H_res must move (asserted > 0.5 from the identity, measured 0.8809) and stay on the manifold (row < 1e-3, col < 5e-2), and the T = 4 composite must still be doubly stochastic and converge toward uniform mixing.

That middle assertion is aimed at a published failure mode of this exact constraint: 2603.20896 (s²HC, NeurIPS 2026) reports that under double stochasticity “learned matrices collapse around the identity initialization and diminish cross-stream interactions”. If our b_res perturbation left the composite at the identity, the arm would be a slower ReZero and row 7b would be measuring nothing. 2602.18308 (JPmHC) attacks the same manifold from the Jacobian-spectrum side, and 2607.14530 (xHC) reports the family’s gains collapsing past N = 4. Three papers attacking the mechanism we are wiring in, all found by reading around the citation rather than inside it — the same move that found AttnRes’s d^-0.5.

5.1b MOVING n FROM 4 TO 2 BROKE THE COMPOSITE ASSERTION, AND THE GATE WAS

WRONG

The composite half of the gate used to assert that the off-diagonal mass of H_res^T grows toward the fixed point n−1 by a factor of 5, stated as “the paper’s §4.1 point 3: repeated application tends to increase the mixing of information across streams monotonically”. That reads as a fair transcription. It is false at n = 2, and it was only ever true at n = 4 by luck.

A symmetric doubly stochastic H_res has a second eigenvalue that is negative — measured {1, −0.7616} at n = 2 and {1, −0.2758 ×3} at n = 4 — so H_res^T alternates about the fixed point instead of approaching it from one side. Off-diagonal mass at n = 2, computed from Eq. 9 before the gate was edited:

T	1	2	3	4	5	6	8
off-diagonal mass	1.7616	0.4200	1.4417	0.6636	1.2562	0.8049	0.8868
distance to J/n = 1	0.7616	0.5800	0.4417	0.3364	0.2562	0.1951	0.1132

The off-diagonal mass overshoots past the uniform matrix and comes back — at T = 2 a two-stream residual is less mixed than a single pass. The old assertion wanted |off(T) − (n−1)| < |off(1) − (n−1)|/5, i.e. 0.3364 < 0.1523: it would have failed on a correct implementation. At n = 4 the same fixture happened to land 0.0174 from the fixed point at T = 4, so the gate was green for a reason that had nothing to do with the property.

What does hold at every T and both n is the distance to J/n, because |λ| < 1 makes |λ|^T monotone. So the gate now asserts that, and the numbers are printed and match the prediction to four decimals: 0.7615 → 0.3363 (predicted 0.7616 → 0.3364).

Why this matters beyond the fixture, and it is a prediction about our operator class, not a test quibble: at T = 2-4 a two-stream mHC residual oscillates around uniform stream mixing rather than settling into it. The manifold still holds at every step — that part is exact, the Birkhoff polytope is closed under multiplication, and the gate asserts it — but “the loop converges toward uniform mixing” is a claim about the envelope, and at n = 2, T = 2-4 the envelope is the one the distance-to-J/n series describes. Anyone reading b_res off a run should read the composite, not the single matrix.

5.2 (b) GRADIENT FLOW, per parameter group

mhc_gradients_reach_every_wired_parameter, on Autodiff<Flex, BalancedCheckpointing> — the trainer’s own strategy, because that is where 8fa5d4c lived. Per-parameter, not “the loss has a gradient”, because the per-parameter version catches a detached arm that the aggregate hides. Six wired parameters asserted non-zero and finite, with b_res called out by name (it is the whole mechanism, and identity degeneration is exactly a b_res that stops moving). And the seventh, phi_pre, asserted to be exactly zero — see §3.3.

5.3 (c) ZERO-DEFAULT IDENTITY — and the honest limit of what one process

can witness

Two gates, because “bitwise unchanged” is not available in-process:

1. mhc_and_rezero_each_move_only_their_own_write. ReZero’s knob (residual_scale: 1.0 → 0.5) moves the ReZero arm and cannot be expressed on the mHC arm; b_res moves the mHC arm; and residual_scale on an mHC model is asserted to change the output by exactly 0.0. A branch-ordering slip — the new else if swallowing ReZero’s statement, or ReZero’s scalar leaking into the mHC write — is invisible to “did the output change” and fatal here. Plus mhc.is_none() when the flag is off, which is the load-bearing checkpoint fact: an Option<MhcBlock> that is None contributes no records to burnpack, so every shipped preset’s state dict is unchanged.
2. mhc_at_init_is_rezero_at_scale_one. mHC’s init computes h + y, which is ReZero’s write at its own init (residual_scale = 1, asserted == 1.0). Measured relative difference: asserted < 1e-2 and > 1e-6 — the lower bound so an arm that is bitwise ReZero (i.e. the projection is not in the write at all) fails. This is also the A/B’s confound, quantified: the two arms start from the same function to within that measured relative error, and differ in the operator. Without it, the lower bound would say the arm is inert and the upper bound would say the arms start from different functions; both are real failure modes and both are checked.

What is NOT claimed, stated so nobody reads more into it: a bitwise-equal CE against a build from before this flag existed is a cross-process fact (two RNG streams, two process-global ParamId counters) and belongs to tools/determinism.py / tools/determinism/, not to a unit test. The diff-level fact is that the ReZero branch’s code is unmodified — the new arm is an else if ahead of it.

5.3a TWO GATES OF MY OWN THAT WERE VACUOUS, CAUGHT BY RUNNING THEM

Both of these passed on the first run and would have shipped as evidence. They are the AttnRes lesson (511daa5: “a gate that cannot fail and a gate that fails for a reason unrelated to what it names are the same defect from two directions”) aimed at me, and both are the same root cause: in-process RNG.

1. mhc_at_init_is_rezero_at_scale_one compared two different networks. The obvious spelling builds a ReZero LoopBlock and an mHC one and compares their outputs. LoopBlock::new draws every weight from the device RNG, and two constructions in one process get two different weight sets — Device::seed does not rewind a consumed stream, and Device::flex() hands back a shared device (the 2026-09-30 finding). Measured: relative difference 1.686, i.e. the “arms differ only in the residual operator” claim was being tested against two random initialisations. Fixed by toggling use_mhc on one block: same weights, same body, same readout, one statement apart. The red reading is the number that makes this worth recording.
2. mhc_differs_from_both_rezero_and_attnres_at_every_depth was the same mistake wearing three arms — three LoopBlock::new calls, so three weight sets, so three different outputs whatever the residual operator was. It would have passed with three copies of the same arm. Replaced by the_three_residual_operators_are_three_functions, which compares the three residual writes on one h and one y (MhcBlock::forward vs h + y·s vs depth_attend over the same two states), with no model construction in it at all, and keeps the model-level half as a separate, honestly-weak mhc_moves_the_readout_at_every_depth (also one block, toggled).

The generalisable rule, added here because it cost two gates: any comparison between two Modules built in the same process is a comparison between two RNG draws unless the weights are shared or transferred. The existing gates in this file all do the second thing. model_seam’s two_models_one_seed_are_bit_identical is the test that was withdrawn for exactly this reason on 2026-09-30; I walked into the same floor two hours later and only a printed number caught it.

5.3b A THIRD, WHICH WAS A WRONG PREDICTION RATHER THAN A WRONG GATE

mhc_res_leaves_the_identity… perturbs b_res and asserts the mixing operator leaves the identity. My first perturbation added ±1.75/2.5 to the init’s b_res = 10 I, which moves the diagonal to 12.5 and leaves the off-diagonals near zero — a more extreme identity, not a less one. It read 1.9e-6 and the message blamed identity degeneration. The gate was right and the perturbation was wrong.

What that accident measures, and it is a property of the ARM rather than of the test: near the paper’s init the map is exponentially stiff. One unit on a single off-diagonal of b_res moves H_res by 2.1e-4, because the diagonal sits at e^10 and the off-diagonal ratio is e^-10. So an H_res that training has not yet moved is exactly the identity, and the arm’s whole contribution in its first steps is H_post ≈ 1 — i.e. a ReZero. The measured gradients agree: b_res 1.8e-6, alpha_res 1.0e-7. The dynamic half of the hyper-network (alpha_res · x̃'φ_res, Eq. 7) is damped twice over — by the 1e-2 gate and by a Sinkhorn sitting on its identity. A run whose b_res and alpha_res have not left init has not started the mechanism, and row 7b says so. §7 owes a cheap log line for it, because the two numbers that decide this arm are not the loss curve’s.

5.4 Red → green, demonstrated

#	gate	perturbation	red reading	green reading
0	Eq. 3 fixture — the gate that never ran	(none: first execution)	Tensor rank 1, dims [1,2] in from_data; then out[2] = 51.654465, expected 18.951490	from_data(TensorData::new(.., [2])); literals 51.65446857 / 69.00595810
1	Eq. 3 fixture	H_post factor 2 dropped	out[0] = 5.175407, expected 8.950813	literal match, 1e-5
2	Eq. 3 fixture	Sinkhorn dropped	H_res[0] = 0, expected 0.8	literal match, 1e-5
2b	Eq. 3 fixture	off-by-one stream in the expected values	out[2] = 51.654465, expected 18.951490	literal match, 1e-5
3	mhc_res_is_doubly_stochastic…	H_res replaced by raw b_res	rowsum 10, off by 9	< 1e-3
4	mhc_res_leaves_the_identity…	b_res left at 10 I (identity degeneration)	mixing distance 4.543e-5	> 0.5, measured 0.8809
5	mhc_at_init_is_rezero_at_scale_one	H_post = 2σ → σ	relative ≈ 0.33	< 1e-2, measured 2.466e-4
6	mhc_and_rezero_each_move_only…	residual_scale read under mHC	exact-zero assertion fails	exactly 0.0
7	mhc_gradients_reach_every_wired_parameter	b_res detached from the write	`max	grad
8	the_tensor_path_carries_a_gradient…	ones fixture (constant input)	gradient legitimately 0 → gate green for the wrong reason	non-zero, finite
9	mhc_res_leaves_the_identity…	composite not taken to T	distance to J/n 0.7613 → 0.7613	0.7615 → 0.3363, predicted 0.7616 → 0.3364

WHICH OF THESE ACTUALLY RAN, stated per row, because this section is about not putting stories in a table. Rows 0, 2b, 4 and 9 were re-demonstrated on 2026-10-01 with red readings captured from actual runs. Rows 1, 2, 3, 5, 6, 7 and 8 have NOT been executed as perturbation runs — the “red” column of each is a prediction computed from the paper’s equations and the f32 arithmetic, not an observation. They are predictions and they are good ones (they come from the same derivation that produced the green literals), but a prediction is not a demonstration and §5.1a’s rule applies to this file too.

What was observed going red, from the development runs, was a different and more useful set — every one of these is a real run with a real message:

gate	what was actually wrong	red reading
mhc_at_init_is_rezero_at_scale_one	the gate itself: two LoopBlock::new calls draw two different weight sets from the device RNG, so it was comparing two networks	relative 1.686
mhc_res_leaves_the_identity…	the perturbation: +2.5/−1.75 on b_res = 10I moves the diagonal up, making a MORE extreme identity	distance 1.907e-6 against a > 1e-2 bound
mhc_res_leaves_the_identity…	the bound: 1e-6 where the physics says (n−1)·e^{−10}	1.363e-4 at n = 4 — three times the bound, and the reason it is now a predicted number
the_three_residual_operators…	the perturbation: b_res = 6I again scales the diagonal	separation 6.174e-3 against a > 0.05 bound
mhc_gradients_reach_every_wired_parameter	the assertion: phi_pre is not in the graph at all, so grad() is None, not zero	“parameter carries no gradient slot”
mhc_differs_from_both_rezero_and_attnres_at_every_depth	the gate: the same RNG trap as row 1, with three arms instead of two	passed — which is what made it suspect

The first two rows are the ones worth keeping: a gate that was measuring two networks, and a gate that could not see the perturbation it was given. Neither is in the perturbation table because neither is a falsification of the mechanism — they are falsifications of the gate, which §5.1a argues is the harder half.

Row 0 is the one to read. It is a gate that this file recorded as demonstrated and that had never been executed, whose two failures were both in the test — a rank error in the fixture, and a stream-indexing slip in the expected values that is the exact class the gate exists to catch. The operator was correct throughout. See §5.1a.

One table row from the earlier pass is deleted, not repaired. It claimed H_postᵀ read as H_post separates out[0] by 13.4 %. Recomputed, it separates it by 0.0 % — H_post is [1, n], so transposing it is a no-op. It was a number that did not exist, and I re-derived that independently before accepting the correction. The replacement row (98.3 %, off-by-one stream) does, and is row 2b’s red reading.

The perturbation runs are ~5 min each and are the cheapest thing left. The script is written and names its own patches: apply P1 (2σ → σ on H_post in burn-mhc/src/block.rs), P2 (sinkhorn_knopp(res_tilde, …) → res_tilde), P3 (.detach() the mHC write), P4 (let the mHC branch read residual_scale), P5 (do not unroll the composite), recording the red and restoring after each. P1 and P2 also turn rows 1-3 of the table above into observations, which is the whole point of having the predictions.

---

6. Cost, and what the fused path is worth here

A machine fact, recorded because it is new and it is a precondition. /home reached 99 % (7.1 GB free) during this task, and the cause is visible in one command: du -sh /home/sehaxe/dormouse-wt/*/target — eight lanes each carrying a private target dir, 3.8 GB to 11 GB apiece, on top of the main repo’s 113 GB. The build lock serialises cargos; it does not serialise disk. AGENTS.md §2.5 already says “a mold error that says ‘Disk full?’ is the disk”, and that is now a near-term risk rather than a historical one. This lane contributed almost nothing (it used the shared main target dir precisely because a cold worktree build does not fit in 7 GB) but its worktree picked up a 3.1 GB vendor/dormouse-fused/target from the concurrent writer of §4.9. Recommendation to the owner, not a change made here: one shared CARGO_TARGET_DIR plus per-lane incremental dirs, or a target/ sweep on lane exit. The worktree script cleans nothing.

SINKHORN_ITERS = 20 on an n × n matrix is not a GEMM: per call the tensor path issues 20 iterations × (2 max_dim + 2 sub + exp + sum_dim + log + add + sub) × 2 phases = 280 tiny launches on a b·t-batched [n,n] tensor, plus a 40-step hand-derived reverse unroll in the backward, times T iterations per forward. On a workload whose GPU is idle 87 % of a warm step (§3.1) that is a launch-bound cost, not a bandwidth one, and it is the arm’s dominant expense.

6.1 The launch count does NOT depend on n, and the A/B prices step time

as a separate verdict because of it

This is the sharpest correction on the 2026-10-01 pass, and it inverts the obvious plan. The natural hope is that a smaller n is also a cheaper n — and it is not, for the launch-bound reason that motivates the whole concern:

	depends on n?	at n = 2 vs n = 4
parameter count	yes	6 155 vs 18 459 on small (d_model = 768) — 0.067 % vs 0.201 % of small, irrelevant either way. 2Dn + Dn² + 2n + n² + 3, asserted against the instantiated block by mhc_counts_every_iteration_and_is_refused_alongside_the_others rather than asserted on paper: an earlier version of this table carried 2 306 / 9 253, which was wrong twice over (it used d_model = 384, and small is 768).
elements per tiny tensor	yes (n² per token)	4 vs 16 — both far too small to fill an SM
kernel launches per call	NO	280 forward + 40-step backward unroll, identical

So the external review’s launch-bound warning — “at our 13 % GPU utilisation, MHC’s extra kernel launches may cost MORE than its parameter count” — is exactly right, and it cuts further than the review stated: the launches do not just exceed the parameter count, they are invariant to the one knob the review told us to turn. n = 2 is the right first rung because it is the K the φ evidence used and what the base paper’s Tab. 1 supports (§1.3), not because it is cheaper at runtime. Anyone expecting the n = 2 arm to run faster than an n = 4 arm will be disappointed, and the only thing that collapses those 280 launches is the fused kernel.

Consequence for the experiment, and it is why row 7b has two verdicts. The review’s condition is “убедиться что экономия качества компенсирует рост latency” — confirm the quality saving offsets the latency growth. That cannot be answered from a BPB column, so row 7b now requires, as a separate reported number:

• Δ step ms and Δ %, from --timers, at a step index ≥ 50, on a quiet card, at the same batch size in both arms, several steps averaged. The step ≥ 50 is not pedantry: AGENTS §2.2/§3.1 record a warm step of ~245 ms against 5 549 ms at step 0 — a 23× spread that is entirely the cubecl autotune cache. A step-time reading with no step index is a reading of the cache, and this arm’s cost is exactly the kind a cold cache hides.
• and the fwd/bwd split, because the 280 launches land in the forward and the 40-step unroll in the backward; if the two halves are comparable the arm is cheaper to fix (one fused forward) than the aggregate suggests.

The verdict rule is stricter than the BPB rule alone: the arm lands only if BPB wins by more than the control’s 3-seed spread and the step-time delta is affordable at the program’s operating depth. The reviewer’s own verdict at our scale — “пока не внедрять” — is honoured structurally: use_mhc defaults off, Option<MhcBlock>::None contributes no records to burnpack, every shipped preset loads unchanged, and the cost of a negative verdict is one flag.

That is the argument for eventually enabling burn-mhc/cuda — and the argument against doing it now. dormouse-core/Cargo.toml carries the dependency with a comment saying so: the fused kernel is reachable on the trainer’s bare CubeBackend (§4.7), it would collapse those launches into one, but it has the untested 1e-7 floor (§4.1), an unmeasured exact backward (§4.2), and a fallback with no counter — ADR-0019’s open COUNTED item in this crate. Swapping the tensor path that every CPU gate above pins, for an unmeasured kernel, on a card this lane does not own, would be exactly the silent swap the doctrine exists to prevent. The cost is named, not hidden, and the A/B’s cost column stays unknown for this reason as well as §3.3’s.

A same-week data point from the adjacent lane, which is why this decision is not merely cautious. d8fa449 (2026-10-01, tests/kda_param_grads_cuda.rs) answered the 8fa5d4c question yes on the trainer’s backend and, in the same commit, found that burn-gdn2’s fused adjoint is wrong and non-deterministic (4.5e-2..2.6e-1 from a central difference of its own forward, on q/k/v, beta_proj and decay.b_alpha, while three post-op parameters matched an independent CPU autodiff at 2e-7) and that a fused forward launches on the trainer’s backend and its result is discarded — the ADR-0019 silent-fallback shape, fused kda=0/0 meaning “not on this path” rather than “slow”. So the fused-kernel family in this repo now has one proven case of a hand-derived adjoint that compiles, engages, and is numerically wrong. burn-mhc’s adjoint is the same shape of thing (hand-derived, 40-step reverse unroll, gated only by a to_vec-visible counter) and has never been compared to anything but its own tensor path. That is the case for keeping the tensor path on tonight’s build and for the order in §7: the gradient gate (§4.3) and the floor (§4.1) come before the feature flag, not after.

---

7. Owed, in the order it should be paid

1. Eval-line field mhc=<ran>/<asked>. probe::MHC exists and is gated; the print is in the hot train/src/lib.rs, which another lane owns, so it is not wired here — same debt AttnRes carried, and mhc is already in probe::NAMES so the one-line field is all that is missing.
2. fused_node_carries_a_gradient, §4.3 — cargo test -p burn-mhc --features cuda,autodiff -- --ignored. The one clause that cannot be checked anywhere else. Until it runs, “the fused Sinkhorn trains” is unverified, exactly as 8fa5d4c’s own message said about itself.
3. The 1e-7 floor (§4.1): a counter or the log-domain formulation, plus a test at extreme logits that exercises the KERNEL rather than the tensor path.
4. burn-mhc/cuda on dormouse-core (§6), gated on 2 and 3.
5. The cfg(cuda) on a CPU test (§4.4) and the relative tolerance on sinkhorn_fused_backward_matches_tensor (§4.2). Whoever has the card.
6. H_pre, if a quarter-width body is ever wanted (§3.3) — 11 % of the paper’s measured effect, and a different experiment.
7. mhc_streams = 4, the base paper’s own App. Tab. 1 rung and mHC’s shipped default — the follow-up row after n = 2, and the capacity ladder below n = 4 that 2607.14530 says does not pay. Note before spending anything on it: the launch count is the same at n = 4, so this rung buys quality, not latency (§6.1), and row 7b’s step-time verdict does not improve.
8. A non-symmetric b_res fixture for Eq. 3 (§5.1a). The current one is symmetric, so H_res and H_resᵀ are indistinguishable and the gate cannot catch a transposed residual matrix — the single most plausible remaining transcription error in the operator.
9. The eval line must also report b_res’s distance from its init, not just the mhc counter (§5.3b): the arm’s first N steps are a ReZero with a Sinkhorn in front of it, and a run whose b_res and alpha_res have not left init has not started the mechanism. The two numbers that decide this arm are not on the loss curve.

---

8. Test tallies, and what was NOT run

Run and green, this tree, 2026-10-01:

gate	result
cargo test -p dormouse-core --lib	74 passed, 0 failed (67 pre-existing + 7 new mHC gates), 119.8 s
the 7 mHC gates, with their printed measurements	§5.1, §5.2, §5.3 below

The measured numbers, all printed by the tests themselves rather than quoted from this file:

H_res at init            max |entry - I| = 4.542e-5   (predicted (n-1) e^-10 = 4.540e-5)
H_res at 2(J-I)          max |entry - I| = 0.8808     (predicted 0.8808)
                         off-diagonal mass 1.7616 -> 0.6636 over T=4, i.e.
                         distance to the uniform J/n  0.7616 -> 0.3364  (both predicted)
mHC@init vs ReZero@scale 1, IDENTICAL weights          relative 6.505e-5   (bounds 1e-6 .. 1e-2)
gradients  b_res 1.105e-6 · b_post 1.224e-5 · alpha_res 7.288e-9 ·
           alpha_post 1.824e-8 · phi_res 2.550e-8 · phi_post 5.976e-7
params on `small` (d_model = 768)   n=2 -> 6 155     n=4 -> 18 459

The two that matter most for the A/B are the last two rows: alpha_res’s gradient is 7.3e-9, i.e. the input-dependent half of the hyper-network is frozen to four significant figures below the static half’s, and the arm’s whole parameter cost is 0.067 % of small at the default n = 2. If row 7b lands as a null result, these two numbers are where to look first.

NOT run this session, and named rather than implied:

• preset_exec. The requirement is 11/0 untouched; situ’s own run of it finished green (11 passed / 0 failed / 2 ignored, 1 976 s) but on its worktree, which is not this tree, so it is not evidence for this change. It needs ~30 min under the build lock and the lock queue on this machine reached 14 waiters tonight.
• The red→green demonstration table (§5.4). The predictions are computed from the papers and the f32 arithmetic and are recorded above; the runs that make them red are owed, and the rule that came out of the ones that were run (§5.1a) is the reason they are cheap to run later.
• burn-mhc’s own suite and dormouse-data / dormouse-train.
• Anything CUDA. No gate in this commit was executed on a GPU; §4.2-§4.4 are all #[ignore]d or cfg(cuda) and say so at the line.

Two editors were in this worktree (§4.9). The dormouse-core --lib run above is the state of the tree at the moment it was taken; if a later commit touches loop_block.rs, re-run it rather than trusting this line.