Board: RK3588 10.3.0.236, kernel 6.1.115-vendor-rk35xx-fence, governors verified inside every run (DDR 2112 MHz / A76 2400 MHz, both performance), NPU guard held, 0 faults. Stack: llama.cpp-rockchip diverged-history @ 35e76dbd0 + local P0.1/P0.2, ork-driver 1.0.168. Model: Qwen3-1.7B-BF16.gguf for every arm. ork_bench … 128 8 128, 2 reps, arms interleaved. Quality: ork_ppl, 8 × 512-token windows, 4088 scored tokens, ubatch=512, real wikitext. Calibration is DISJOINT from evaluation: GPTQ calibrated on wiki_sub.txt bytes 20000-60000, evaluated on bytes 0-20000. Every arm uses the same split.

Two findings, and they are independent.

Arm T is 1.30× better PPL than i4a8 (26.36 vs 34.30) and 1.92× better than rotation-only native (50.73). Plain-DT_I8 promotion is marginally better again at 25.78 — a 2.3% edge over I8_ROT, which is right at this project's ~2% screen, so call it a tie that leans the way AGENTS.md predicts ("rotation may help int4 and HURT int8").

The promotions genuinely ran, verified by census rather than by size: 56 I8_ROT/I8 entries = attn_v + ffn_down × 28 layers, with weight blobs 0.657 → 0.848 GiB ≈ 5.16 bpw.

So native 4-bit, built with the landed ladder, is not merely "acceptable" — it beats the i4a8 storage tier the 27B currently ships.

And this is the finding that matters, because it was initially mis-attributed.

The first reading of arm T's profile — run 25033 ms, 64852 µs/matmul — was divided by the invocation count and read as "the DT_I8_ROT tier is ~11× slower per matmul". That was an inference from an average and it was wrong. Two experiments killed it:

• ORK_I4_PROMOTE=i8 changes the 56 promoted entries from DT_I8_ROT to plain DT_I8, leaving the rotated path entirely. Prefill: 9.10 / 9.11 against arm T's 9.08 / 9.08. No change. If the rotated int8 tier were the cost, this would have fixed it.
• ORK_DEBUG_RESET=1 names the caller, and settles it outright:

110 is the promoted-weight invocation count. Every promoted weight drags the datapath I8_CHAIN → I4_CHAIN and back, once each way.

The arithmetic closes against a price this repo already measured. src/npu/core/mode.c:148 records "MEASURED 2026-08-24 on the 27B mixed-tier run: mean 105 ms per ACT_RESET, 160 of them = 16.8 s of a 161.6 s scored run". Here: 220 × 105 ms = 23.1 s, against a measured run time of 25.0 s — within 8%, with the residual being the induced cold re-warms the comment says land later in the run path.

So the cost is the transition, not the tier. A homogeneous pack of either precision pays one reset for the whole forward; a mixed pack pays two per promoted weight.

It was delivered correctly (sudo env ORK_MIXED_NOTHRASH=1 …, not a bare prefix — sudo strips ORK_*) and measured 9.07 tok/s, unchanged. The reason is now visible in the caller field: that knob's keep-warm predicate governs XP_MC_MM, and the rows actually firing here are profiles 8 and 2 on the I8_CHAIN ↔ I4_CHAIN transition. Different rows, untouched by the knob.

That makes this a XSPEC[] Phase-2 question of the kind AGENTS.md describes — one row per direction, each a literal RC_* reset condition — rather than an intrinsic property of mixed precision.

This supersedes §3's retraction as a statement about the present. §3 remains correct about what fc06bc6 measured in August and about why my 1.7B probe proved nothing. The question "is the reset needed on the kernel we run today?" is separate, and is answerable without reference to August.

Built a hazard-fidelity pack rather than reusing the August artifact (which is pre-v7 and would be regenerated — see the locking note below): Qwen3.6-27B Q8 source, native W4A4, ORK_I4_PROMOTE=i4a8, scattered attn_output,ffn_down, no GPTQ / no SmoothQuant (a coherence test, not a ladder arm — its PPL is not a quality number). Census: DT_I4 × 80 + DT_I4_NATIVE × 320, i.e. the 80 promoted ride the int8 datapath while the rest are int4 — exactly what 42cc655 describes.

All three controls hold, which is what distinguishes this from the §3 attempt:

The answer:

On the current kernel the I4_CHAIN↔I8_CHAIN ACT_RESET is pure overhead. The suppressed arm also ran faster (190.1 s vs 227.9 s wall).

Two findings, kept apart. Is it needed now? — answered, fully, by the table above. Did the kernel rework cause the change? — not establishable: neither fc06bc6 nor 42cc655 records a uname -r (grepped both), and that window sits inside active kernel work and an unresolved wedge investigation. We do not know what we are comparing against, so no causal claim is made.

Scope: 2 windows (enough to separate 10.64 from 3.4M by ~300,000×, but a small sample for subtle corruption — the bit-identical NLL across all three arms is the stronger evidence); one fresh pack, not the August artifact.

⚠ The August artifacts are pre-v7 and were unprotected. q27_ra8 and q27_g512 carry 32-byte footers that this build reads as version=2 with a garbage n_entries, so it would declare them stale and regenerate over them. Both are now .locked (the lock is checked before the footer is parsed, which is exactly the population that needs it) and q27_ra8 is copied to ~/ork-archive/.

Incidental, and it dissolves the 1.4% budget gap: this correctly-built pack is 14.22 GiB resident (14.69 with Bf) against the old pack's 22.656 — because 320 of 400 weights stay int4 nibbles instead of inflating to int8. Comfortably inside the 22.35 GiB budget. The convert itself logged tracked wcache: 22.66 GiB over 400 weights vs budget 22.35 GiB, reproducing the original figure exactly.

Read this correction before the numbers below. The suppressed-reset run was clean, and that is the expected non-result at this configuration — not evidence the reset is removable.

ORK_XSPEC_NORESET was added by fc06bc6 (2026-08-24), which had already run this experiment on the real 27B mixed-tier workload:

baseline               PPL      10.4421      160 resets
ORK_XSPEC_NORESET=1    PPL 3442988.92          0 resets   <- garbage
ORK_XSPEC_NOCLEAR=1    PPL      10.4421      160 resets   <- fine

"So the ACT_RESET is required across I4_CHAIN<->I8_CHAIN; the clears are not."

42cc655 measured the complement the same morning — widening ORK_INT_DT to keep warm across 4→3 took PPL 10.4421 → 5,029,965.1161 — concluding "The I4_CHAIN->I8_CHAIN ACT_RESET is load-bearing; the datapath cannot cross that boundary warm."

Why the run below came out clean. fc06bc6 states its own positive control "stays COHERENT at K=512/1024, M=1, plain packs, single domain, so the probe is not exercising the hazard", and that closing the gap "needs closer fidelity to the 27B config (large K, M=64, DT_I4_ROT_A8 packs, multi-domain)." Arm T on the 1.7B is K=2048, single domain, DT_I8_ROT — it misses the known-hazard configuration in several dimensions at once, so a coherent result there says nothing.

The lesson that commit exists to teach: ORK_XSPEC_NORESET must never be trusted without a positive control that demonstrably corrupts in the same configuration. A clean NORESET run at small K / single domain is the expected non-result; reading it as "the reset is removable" is precisely the misreading fc06bc6 added its control to prevent. This page made that mistake and is correcting it.

What the numbers below DO establish: a ceiling. They price what mixed-tier prefill would cost if the transition were avoided — which is the real lever, and the direction 42cc655 names: "The way to cut this churn is to avoid the transition (the uniform G=512 pack takes the same model from 160 resets to 2), not to widen the macro." Promoting contiguous whole layers instead of a tensor type scattered across every layer would cross the boundary twice per forward instead of ~110 times, with every reset intact.

ORK_XSPEC_NORESET=1 suppresses only the ACT_RESET, leaving the warm/size clears and last_dt intact (mode.c:146). That separation matters: mode.c:54-57 records that an earlier 27B corruption came from a kw=1 which turned off reset and clears together, so which half was load-bearing had never been established. This probe establishes it.

• 6.17× faster, interleaved arms, 2 reps each.
• Output is bit-identical — same PPL to four decimals, same mean NLL to six, and every intermediate window matched (2: 25.2915, 3: 25.7440, 5: 23.3881, 6: 24.2838). The generated text matched too.
• Run-time delta 25033 − 1810 = 23223 ms against the 220 × 105 ms = 23100 ms predicted from mode.c:148's measured price. Within 0.5%.
• With resets gone the profile inverts to the normal shape: quant dominates (2249 ms, 54%), run 44%.
• The transitions are still logged 221 times in the suppressed arm — the log line precedes the suppression — so the mode changes still happen; only the device reset is skipped.

55.64 tok/s also exceeds the 46.84 homogeneous-native ceiling, which makes sense: the 56 promoted weights run on the faster int8 kernel once they stop paying 105 ms each to get there.

42cc655 prescribed it: "The way to cut this churn is to avoid the transition … not to widen the macro." Measured here at matched budget — arm T promotes 56 of 196 weights, and 56 is exactly 8 whole layers × 7 matmuls, so the two arms promote the same count, the same bytes and the same bpw, differing only in where the promoted weights sit:

5.57×, with every reset intact and no XSPEC row touched. The reset breakdown shows exactly why: scatter pays 110 × profile=2 (4→3) + 110 × profile=8 (3→4); contiguous pays 2 × and 1 × plus the one cold entry. The forward crosses the precision boundary twice instead of 110 times.

This reaches ~91% of the 55.64 tok/s ceiling that the (invalid) suppressed-reset run priced, which is the expected result once the transition count is the only thing left to pay for.

⚠ These are SPEED arms only — built without GPTQ or SmoothQuant, so their perplexity is not comparable to arm T's and is not quoted. The open question is quality: contiguous whole-layer promotion is a different selection from the Q4_K_M structural prior (attn_v + ffn_down), which this project measured to beat qerr ranking by 3.8% at matched bits. Spending the same bits on whole layers may cost accuracy. That needs a GPTQ arm and a PPL run to answer.

§4 measured arrangement with the reset toll in place. With ORK_XSPEC_NORESET=1 applied to both arms (interleaved, order swapped, 2 reps, governors verified in-run):

Quality is a tie (0.84%, inside the 2% screen) and resident is a tie (0.9%). Scatter wins 1.92×. Both arms reproduced their resets-on baselines bit-identically (PPL to four decimals, NLL to six), which is what makes the timings trustworthy — contiguous had never been run under NORESET before.

So the ordering depends entirely on whether the reset is paid:

Opposite recommendations. The arrangement question cannot be answered before the reset question.

With the toll removed from both, transition count cannot explain a 1.92×. The cause is the measured orki_i4_hcap table (src/npu/i4/chain.c:94):

K=2048 -> H = 8          K=6144 -> H = 3

H is rows-per-batched-int4-submit, and per the weight-DMA amortization principle the per-submit cost is dominated by re-streaming the weight and is independent of row count — so H is the amortization, and K=6144 is 2.67× worse.

ffn_down is K=6144, the largest-K tensor in the model. Scatter promotes all 28 of them off the int4 path; contiguous (layers 20-27) leaves 20 of them on it. That is the 1.92×.

(First guess was Wb = (131072/K) & ~63. Wrong — it clamps to 64 at both K values. Recorded so the discarded hypothesis does not get re-derived.)

Contiguous promotion is strictly dominated: tied quality, tied resident, 1.92× slower, and 2.3× the build time (GPTQ 54.4 vs 23.7 min — the K³ term on those same K=6144 weights).

And one fact now explains three observations that had been treated as unrelated:

attn_v+ffn_down is not a heuristic. ffn_down is the largest-K tensor, and promoting it off the int4 tier is right on all three axes simultaneously.

The XSPEC-reset counts under NORESET are still 221 (scatter) and 4 (contiguous), not ~0. That is expected and not a knob failure: mode.c:143 logs the transition, :146 suppresses the reset after. The counter measures transitions-that-would-reset. Evidence the knob worked is the timing and the unchanged output.

2-factor isolation, all scatter (attn_v,ffn_down), all under ORK_XSPEC_NORESET=1 so nothing hides behind the reset toll. All three packs have identical census and resident (I4_NATIVE × 140

• I8_ROT × 56, 1.230 GiB), so any difference is purely runtime. Interleaved, order swapped, 2 reps.

GPTQ is free and worth 12.0%. C→B is 28.4573 → 25.4084 in PPL for 0.6% prefill — a tie. This also settles the alternative that could not be excluded from arithmetic: a GPTQ-altered weight does not measurably change the doorbell-miss rate. GPTQ has no runtime footprint, as its pack-time-only nature predicts.

SmoothQuant costs 7.0% prefill AND 3.8% PPL. B→A is 59.11 → 54.95 and 25.4084 → 26.3637. Both deltas are outside the 2% screen. On this model it is worse on both axes at once — strictly dominated, and it should not be an unconditional default.

⚠ Model-dependent, not a refutation of the recorded result. AGENTS.md records SmoothQuant as the first technique that stacks with GPTQ (57.49 → 52.65) — measured on qwen3.5-0.8B. This is Qwen3-1.7B with a different promoted set. Both can be true; what it rules out is shipping SmoothQuant on by default without a per-model check.

Best configuration measured: B — scatter + GPTQ, no SmoothQuant, resets off: 59.11 tok/s at PPL 25.4084. Better quality and faster than the full ladder.

§4 reported the ladder costing ~1.9× prefill, inferred from the contiguous pair (no-ladder 50.83 vs full-ladder 26.60). That does not transfer to scatter, where the measured cost is 7%. The earlier figure came from two contiguous packs and was read as a property of the ladder rather than of that arrangement. A plausible reason — unverified — is the same H story as §5: contiguous keeps 20 K=6144 weights on the int4 path, where SmoothQuant's per-element divide is 3× the work it is at K=2048. Stated as a hypothesis, not a result.

An earlier revision of this page proposed giving XP_I4CHAIN (profile 8, { KWP_NONE, RC_ALWAYS }) a keep-warm predicate, on the strength of the clean run above. That proposal is withdrawn. 42cc655 already tried the equivalent — widening ORK_INT_DT so KWP_MC keeps warm across 4→3 — and measured PPL 5,029,965. RC_ALWAYS on that row is not unexamined conservatism; it is the conclusion of a measurement, and the absence of a keep-warm predicate is deliberate.

The lever is transition COUNT, not the reset. Same ceiling, every reset intact, correctness untouched.

• Every promotion scheme is mixed by construction. The ladder's promoted arms are all affected together. Their quality numbers stand; their shippability is what is in question.
• MoE packs are mixed by design (int4 experts, int8 attn/dense) and may have been paying this silently.

• The "~54×" native-W4A4 prefill claim in the source was stale by ~15×. Measured 3.62× (169.46 → 46.84), and its stated cause — "single-row / no-M-amortization" — was also wrong: the path batches and chains (281.7 programs/ioctl, 1157 of 1158 submits chained). ORK_I4_BCHAIN is no longer a knob; it is the default for M≥2. Corrected in ggml-ork.cpp.
• Native W4A4 buys nothing at decode. All arms measured 1.02 tok/s, because dense M==1 is declined by supports_op and decode never reaches the NPU. The doc's "decode/streaming regimes where int4's bandwidth wins" do not exist on this stack.
• ORK_GPTQ_CLIP_MIN=0.531 should not be set. The live default is 0.10 with a self-tuning grid (src/ork_gptq.c:127); passing 0.531 would raise the floor and regress the result.
• GPTQ already skips the int8 tier. Both builds logged finalize: 140 weights against 196 entries with 56 promoted — 140 = 196 − 56 exactly. There is no wasted GPTQ on promoted weights to reclaim.

GPTQ finalize 23.7 min of a 28.5 min total build on the 1.7B (140 weights, chunked 8-wide over 8 cores, 172 MiB Hessian/weight at K=2048). Extrapolating to a 27B by the dominant O(K³) term (K 2048 → 5120) and weight count (140 → 400) gives ~18 hours and ~1 GiB Hessian per weight — an extrapolation, not a measurement, and the largest open risk to making GPTQ a default.

ggml-ork.cpp:11393 returns true early — ahead of the N-cap, the sbits < 5.0 gate and the M >= min_m decode decline — whenever the source is a stub and the tensor is in the pack index:

if ((g_ork_ctx->source_is_stub || cpudec) && g_ork_ctx->persist_idx.count(op->src[0]->name))
    return true;

So the same pack benchmarked against a stub and against a full GGUF routes decode to different engines. Anyone comparing those two configurations will get different answers for reasons that have nothing to do with what they were measuring. (It is also what makes packing the LM head work: under a stub it runs on the NPU at decode, and it cannot fall back to a file hole, because the early-accept fires precisely when the source is a stub.)

Section 6 left the native-vs-i4a8 run gap attributed to "the int4 datapath" without decomposing it. The working hypothesis going in was that native int4 issues ~16× more weight streams, because orki_i4_hcap(K) caps rows-per-batched-submit at H=8 (K=2048) and H=3 (K=6144) where int8 amortizes far more rows per weight load. That hypothesis is half right, and the missing half carries the larger coefficient.

Board 10.3.0.236 · kernel 6.1.115-vendor-rk35xx-fencefix · governors verified · re-baselined

Submit counts are identical. Native int4 does not issue more ioctls — it packs 167× more chained programs into the same ones. The run gap on this kernel is 3.50× (§6's 3.62× was a different kernel; they are not comparable).

programs = Σ over weights of  ceil(M/H(K)) × ceil(N/Wb(K))

Computed over the model's per-shape census this predicts 329,910 against the measured 325,696 — agreement to 1.3%, so the program count is fully explained by tiling geometry with nothing left for a hidden per-call effect.

Decomposing the ~855× per-weight inflation:

So H is a real mechanism — a fixed-N control measured 512 vs 1376 programs, 2.69×, against the predicted H ratio 8/3 = 2.67× — but it is the smaller term. The original hypothesis omitted the N-tiling dimension entirely, which is exactly why its arithmetic would not close (~8× predicted against 3.16× observed).

Why 855× more programs costs only 3.50× of wall: per-program cost is 48× lower on int4 — 6.9 µs against int8's 328 µs. Chaining absorbs most of the inflation. This is consistent with the regime-dependent per-task cost recorded in NPU-Quirks: a task's DMA hides under a neighbour's compute.

K=1024 → 131072/1024 = 128, &~63 = 128
K=2048 → 64,  &~63 = 64
K=4096 → 32,  &~63 = 0 → clamped to 64
K=6144 → 21,  &~63 = 0 → clamped to 64

For every K ≥ 2048 Wb is 64 and cannot widen. That constant is the ×43 term and the better optimization target.

Note that at K=6144 a 64-column tile is 64 × 6144 × 0.5 B = 192 KB — 3× the 64 KB the formula is documented to contain. The clamp already violates its own stated bank-containment invariant at large K and works anyway, which matches OPS_REGISTRY's standing note that Wb is a modelling gap (Wb=256 fails at M=8 and works at M=64, yet the documented law has no M term).

ORK_I4_NOROT=1 on the native pack gives 325,696 programs, 1157 chained, 1155 ioctls — identical to the rotated run. The Hadamard rotation was previously excluded from the run gap by a structural argument; it is now excluded by measurement.

Worth recording: NOROT was slower (28.48 vs 47.89 tok/s) at identical geometry. That is value-driven, not geometric — garbage values hit the in-band int16 doorbell sentinel 0x7fff and cost misses and retries (see NPU-Quirks, the int4 doorbell seed being in-band for int16 output). Its output is garbage by construction and its correctness is not reported.

Read §8 before acting on anything below. These were arithmetic over the measured hcap table and the Wb expression, written down as predictions so that measurement could test them. It did: (a) the n(12−n) objective is wrong — the binding term is activation DMA, a function of Wb alone, and the knob that unlocks Wb is H, not the weight-bank count; (b) the DMA asymmetry is inverted (the weight streams once, the activation is re-streamed) and transposition is net negative; (c) the slicing table is arithmetically right but physically wrong — slicing below K=2048 is 4.5× slower. What survives: Wb genuinely can be raised, worth 2.21×, at K ≤ 2048 only. The text is kept unedited below as the record of what was predicted.

Flagged explicitly as arithmetic over the measured hcap table and the Wb expression, pending a board run:

(a) H and Wb compete for one 12-bank budget. 0x1040 = ((12-n)<<4)|n splits CBUF into n DATA banks and 12-n WEIGHT banks. H scales with the DATA side (H_max = n×16384/K, and the hcap table matches n = 1 exactly at K=1024/2048/4096); Wb's 131072 constant is 65536 B = 2 banks, i.e. the weight side. Since programs ∝ 1/(H·Wb) ∝ 1/(n(12-n)), and n(12-n) is 11 at the current n=1 but 36 at n=6, the shipped split sits near the worst point of that curve — a balanced split would be up to 3.27× better with no new capacity.

The obstacle: chain.c says "DO NOT RAISE IT", citing a 2026-08-21 sweep where raising Wb gave 1–2 second submits (the driver self-healing a straddled batch). The open question is whether that sweep moved the bank split alongside Wb; if it raised Wb alone, its failure is explained without refuting the capacity reading. Any retest needs both validity criteria — bit-exactness and uninflated submit time — because a self-heal returns a correct checksum at 1–2 s, so checksum-only probing cannot distinguish valid from healed.

(b) Transposing the matmul cannot help program count. Computing Cᵀ = Bᵀ·Aᵀ swaps which dimension is divided by H and which by Wb, and the product is invariant:

untransposed:  (M/H) × (N/Wb)  = M·N / (H·Wb)
transposed:    (N/H) × (M/Wb)  = M·N / (H·Wb)

At M=128, H=3, N=2048 the ceiling effects are 43×32 = 1376 vs 683×2 = 1366 — 0.7%. The open part is not the count but the DMA asymmetry: per program one operand streams as the "weight", and in prefill that is the large K×N matrix while the activation is small. Whether a transposed formulation could make the small matrix the streamed operand is unresolved.

(c) K-slicing, not padding, is the lever on tile shape. N is already forced to a multiple of 64, so there is no Wb remainder to recover; M-padding to a multiple of H recovers ~0.8%. But both H and Wb scale as 1/K, so per-slice programs scale as K² while slice count scales as 1/K, making the total linear in slice width:

8× fewer and fatter programs — a short-K tile fits more rows and more columns in the same bank. The slicing machinery already exists (pack.c slices at 8192), so this may be a threshold change rather than new code. The cost is host-side partial-sum accumulation, since CACC partial-sum accumulation is architecturally absent on this part; the crossover is unestimated.

The gap is not "int4 issues more submits" and not H alone. The int4 datapath decomposes each matmul into ~855× more chained programs — Wb=64 N-tiling at ×43 and H at ×19 — against a 48× lower per-program cost, netting 3.50×. Wb is where the remaining headroom is, and the 12-bank split is the mechanism that would move it.

Section 7 left three derivations open. All three are now measured, and two are refuted. Board 10.3.0.236 · kernel 6.1.115-vendor-rk35xx-fencefix · governors verified · matmul-level through ork_i4_mm_run (the production BCHAIN path) at real shapes.

The headline: Wb can be raised, worth up to 2.21× on a single matmul — but the mechanism is not the hypothesised one, the objective function proposed in §7(a) is wrong, and the lever dies entirely at K ≥ 4096.

Measured (git archaeology). The 2026-08-21 commit order was e730244 (23:57, "5.5× too small") → 4796f8c (00:07, "M=64-only, times out at production M") → 0c89c31 (00:13, "containment rule, DO NOT RAISE IT") → 0a7a6f1 (00:33, the 0x1040 crack) → 105438e (11:59, ORK_I4_DBNK added). ORK_I4_DBNK did not exist when the Wb sweep ran, and the sweep set only ORK_I4_WB. i4_bank_sweep.c still had exactly two modes — wb alone and split alone. Nothing had ever swept them jointly. So §7's suspicion that the knob was unpaired is confirmed.

But §7's reasoning for why that mattered was backwards. The int4 default 0x1040 is the captured value 0xb1 = DATA=1 / WEIGHT=11, and mregs defaults to 0x5f, which deliberately excludes bit 0x20 — so int4 never recomputes 0x1040 and the split is frozen at its most weight-biased setting. There was never any weight-bank capacity to be gained by pairing; pairing can only shrink the weight side. Leaving the split alone was the most favourable choice for a large Wb.

Measured controls (K=2048, N=2048, M=128, 2 passes each):

The split alone unlocks nothing. H is what unlocks Wb — and H in turn needs the split, because DATA=1 cannot hold H=16. So the pairing is real but runs through H.

The two failure modes are distinct, and that is diagnostic: capacity violations give WRONG (silent miscompute); H-gate violations give SELFHEAL (hardware stall, driver retry, correct checksum after ~4.7 s). 4796f8c's observation was correct; its attribution — "a correctness wall at the bank width" — was not.

Measured, K=2048 N=2048 M=128, 4 reps, both validity criteria applied (bit-exact against the default-config reference and not time-inflated). Baseline 3570–3654 µs.

Best valid point: n=2 / Wb=320 → 1707 µs = 2.09×, 56 programs against 512.

The WRONG boundary is predicted exactly by the weight-bank capacity rule Wb·K/2 ≤ (12−n)·32768 for every n=2…8. Only n=1 breaks it, because the separate H-gate bites first.

§7(a) proposed maximising H·Wb ∝ n(12−n), predicting n=6 and up to 3.27×. The measurement is not program-count-bound at the margin. n=4/Wb=256 has 32 programs and is slower (1807 µs) than n=3/Wb=256 with 48 (1718 µs) and n=2/Wb=320 with 56 (1707 µs). The best point has n(12−n) = 20, not 36, and 6/6 does not even reach it (n=6 caps at Wb=192, 1.82×).

The binding term is activation DMA = M·N·K/Wb, which depends on Wb alone (see §8.3). The correct objective is therefore: take the smallest n that clears the H-gate, leaving the most weight banks for the largest Wb. At K=2048 that is n=2.

Measured, K=2048 n=2, H forced, Wb=128/256: H=8 SELFHEAL, H=10 SELFHEAL, H=12 ok, H=14 SELFHEAL, H=16 ok, H=24/32/40/48/64 ok. At K=1536 the default H=11 already clears it. No H-threshold, byte-budget or mod-4 rule survives all points. Same shape as orki_i4_hcap and the fp16 K=128 envelope: scan, never bisect.

Measured, M=128, N=2048 unless noted:

At K ≤ 1536 no split change is needed at all — Wb raises at the shipped 0xb1, because H is already ≥ 11 there. The shipped (131072/K)&~63 under-uses by 5× at K=1024 and 7× at K=768. e730244's 11-bank reading was right at those K; 4796f8c's correction was right only at K=2048/H=8. Both were partial, which is why the question stayed open for seven weeks.

At K ≥ 4096 the lever is dead: 0x1040 ≠ 0xb1 gives rc-fail at every Wb including the default 64 (n=2/4/6/8 all rc-fail at K=4096 and K=5120), and Wb>64 at 0xb1 also rc-fails. 0x1040 carries the K-reduction schedule as well as the split, and the captured value is K-sensitive — only 0xb1 survives at large K. The boundary in (2048, 4096) is unmeasured.

M-independence (measured): 2.09× at M=128, 2.07× at M=256, 2.08× at M=512; n=1/Wb>64 self-heals at M=32…512 alike. The gate is H, not NG and not M.

The algebra is confirmed: ratio exactly 1.000 at K=1024/2048/4096/8192; 0.993 at K=5120. §7's 0.7% figure stands.

But §7(b) guessed the streaming direction wrong. Measured from the code (i4/chain.c:363-390, i4/run.c:121): the loop is N-tile outer / M-group inner, and ORK_I4_WREUSE defaults to 1, setting CBUF_CON0[13] for every g>0 — so each weight N-tile is loaded once and held CBUF-resident across all NG M-groups. Therefore:

• weight DMA = N·K/2 — the whole weight, exactly once (2 MiB at K=2048 N=2048);
• activation DMA = M·N·K/Wb — re-streamed every program (8.4 MiB at Wb=64, 1.68 MiB at Wb=320).

The big matrix already streams once; the small 128 KiB activation is the one re-streamed ~65×. That is also the mechanism behind §8.2 — Wb is the activation-re-stream divisor, which is why time tracks Wb and not H·Wb.

Transposition is net negative, on three counts: (1) program count invariant; (2) the activation term M·N·K/Wb is identical under transposition — only the once-streamed term changes, 2 MiB → 128 KiB; (3) decisive — nc is clamped to NC = ceil(dim/Wb), so transposed NC' = ceil(M/Wb) and at M=128 you would need Wb ≤ 42 for 3 cores, below the hard 64 minimum. A transposed prefill matmul could never use more than 2 cores. Plus Cᵀ must be transposed back (measured 244 µs) against a ~1.9 MiB saving worth ~410 µs — a wash before the core loss, clearly negative after.

Padding: confirmed negligible. N%64==0 is enforced at chain.c:445, so there is no Wb remainder. M-padding recovers ceil(M/H) slack only: 0.8% at M=128/H=3, worst case 3.1% (H=22 at K=768), 0–0.3% at the lever's configs.

§7(c)'s table reproduces exactly — and measured time goes the other way (M=128, N=2048):

8× fewer programs costs 4.5× more time. There is a ~1.5 ms per-call floor (submit + drain + A-stage) that does not shrink with K. Confirmed end-to-end through the real sliced path (K=4096 as 2×2048): unsliced 5327/5397 µs, sliced 7643 µs — 1.43× worse.

The premise also partly fails on paper: Wb floors to 64 for all K ≥ 2048, so 4096→2048 does not raise Wb at all and the program count is exactly neutral.

The one form that pays — the refined version. Slice down to 2048, not below, not to cut programs but to move a K out of the dead zone into the range where the Wb lever is valid:

Bit-exact verified — unsliced ref=10db0b1a8dc63164 (K=4096) and 1e842b943ba3c7e4 (K=5120) are identical to the sliced and lever-arm checksums.

Reachability: the K-slice→BCHAIN decomposition already exists (i4/run.c:195-240) and already accumulates C[e] += Ct[e]. The only gate is the pack-time slice width, #define ORK_I4_KS 10752 (src/npu/internal.h:567) — 10752, not 8192; 8192 is the int8 value. It is a compile-time macro used by the loaders, so lowering it changes the .orkpack tile layout and existing packs must be rebuilt. The accumulate loop is scalar with a per-call malloc — both easy wins if pursued.

1. It does not make native W4A4 competitive. On Qwen3-1.7B (hidden 2048, ffn 6144), q/k/v/o/gate/up are K=2048 and get ~2.1–2.2×; ffn_down is K=6144 and gets nothing without slicing, ~1.2× with it. Expect ~1.5–1.7× overall, leaving W4A4 still ~2× behind the inflate→W8A8 path. Worth taking for capacity (half the weight RAM, the documented reason native W4A4 exists), not for speed parity.
2. The dominant failure mode is SELFHEAL, which the geometry ladder never sees. The ladder in orki_i4_run_bchain_db protects against rc=-1 only; a self-heal returns a correct answer ~1300× slower. One point (n=2/H=8/Wb=128) was fast-and-correct on pass 1 and a 4.7 s self-heal on pass 2 — non-deterministic at the frontier. Any shipped point needs margin, and the self-heal path needs to feed the ladder.
3. The best point has zero margin — n=2/Wb=320 is exactly 10 banks (327680 B). Ship Wb=256 (2.01×, one bank of slack) over 320 (2.09×).
4. nc ≤ NC = ceil(N/Wb) starves multi-core on narrow N. At N=512, Wb=320 leaves 2 chunks → 2 cores. Any new rule needs Wb ≤ N/cores.
5. ORK_I4_DBNK has a latent bug. It computes H = n * orki_i4_hcap(K), and hcap is a ceil at non-dividing K, so it overshoots: at K=768, n=2 gives H=44 against a true ceiling of 42 — measured WRONG at every Wb tried. It needs min(n*hcap(K), n*16384/K) or a per-(K,n) table.
6. Per-part risk. chain.c:542 records that Wb and hcap were measured on one die and that a second RK3588 stalls at geometries the first runs. Shipping a wider default enlarges exposure to exactly that.

• §7(a) n(12−n) is refuted as the objective. The binding term is activation DMA M·N·K/Wb, a function of Wb alone. Correct rule: smallest n clearing the H-gate. n=6 is not optimal and cannot reach the best point.
• §7(a)'s "pairing would unlock weight capacity" is backwards. The shipped split is already maximally weight-biased; pairing only shrinks the weight side. The pair that matters is H.
• §7(b)'s DMA asymmetry is inverted. The weight streams once (CBUF-resident); the activation is the re-streamed operand. Transposition is net negative, decisively on the nc ≤ ceil(M/Wb) core clamp.
• §7(c) is arithmetically right and physically wrong. Slicing below 2048 is strictly negative (per-call floor ~1.5 ms); slicing to 2048 to reach the lever is the version that pays, 1.24×.
• ORK_I4_KS is 10752, not the 8192 named in §7(c) — 8192 is the int8 slice width.
• No end-to-end model number was taken. The gate for one is a dense W4A4 model where int4 matmuls dominate prefill; Qwen3-1.7B-W4A4 is the shape that would answer it.

§8 killed transposition on a core clamp. A follow-up tested whether re-optimising the bank split rescued it, on the argument that the two reuse bits are one capability attached to whichever operand sits in the weight role — so transposing would move the matrix you want resident into the role that the validated WEIGHT_REUSE[13] already covers, and the corrupting DATA_REUSE[12] would never be needed. The staleness cap tops out near 360 KiB, and only the small matrix fits under it untiled, which would give NC=1 and both operands read exactly once.

It is expressible, and it is slower. Board 10.3.0.236 · 6.1.115-vendor-rk35xx-fencefix · full-random int4 data · ORK_I4_NOLADDER=1 so a refused geometry reports rc!=0 instead of being silently re-tiled · new probe tools/re/i4_role_probe.c.

The clean comparison controls everything except role assignment — same split, same Wb, same 32 programs, same per-program geometry, same single core, same A/B position: 3106 µs untransposed against 3642 µs swapped. The swapped shape is ~15% slower while moving roughly half the bytes. Reproduced at K=1024 (2884 vs 2239 single-core). The prediction was that single-core transposed would beat a third of the untransposed 3-core time (~580 µs); it missed by 6.3×.

(a) The DMA arithmetic was overstated 2×. The activation is nibble-packed (need_af = NG·(2H)·(K/2)), so the term is NC·M·K/2, not NC·M·K. The transposed advantage over the untransposed best is 1.35×, not 1.73×.

(b) Decisive: these shapes are not activation-DMA-bound, so the saving buys nothing. At matched programs, split and cores, 2.13 MiB took more time than 4.0 MiB. Program count is not binding either at this working point — on the swapped shape 32 programs → 3642 µs, 43 → 3732, 64 → 3771, 512 → 4977: a 16× swing in program count costs 37%. Neither of the two terms the derivation optimises is the one setting the clock here.

(c) The roles are not symmetric in SOFTWARE, even where they look symmetric in hardware. The weight role is packed once; the activation role is nibble-packed on the host, every call, on the critical path (orki_i4_tile_Aslice in bch_db_worker):

The measured single-core deficit is 3678 − 3190 = 488 µs; the host-pack delta is 361 µs, ~74% of it. A real Cᵀ = Bᵀ·Aᵀ puts the big matrix into the per-call host-pack role and the small one into the pre-packed resident role — a 16× increase in per-call host work, structurally, and the exact opposite of what the idea was reaching for. On top of the ~244 µs Cᵀ transpose-back, and on top of ork_i4_mm_pack having to run per call rather than once, since the activation changes every call.

nc ≤ NC exists only because work is distributed by N-tile (lc0 = i·NC/nc). CBUF is per-core, so splitting the g axis instead is a pure work-distribution change — ~60–100 lines over 4 files, mechanical, no new hardware risk. But measured 3-core scaling at fixed geometry is 1.35×, which puts the transposed best at ~2700 µs (~2140 µs on an optimistic 1.70×) against 1738 µs today. Removing the clamp just makes the transposed route lose by less.

NC=1 is also not a general regime: it is reachable only because the transposed weight is tiny. K=2048 N=2048 M=128 at Wb=2048 (NC=1, 8 programs, reuse off) was REFUSED — doorbell miss.

K=2048 N=2048 M=128, n=2 (WEIGHT=10 banks → the law predicts Wb ≤ 320):

The WRONG boundary vanishes entirely with reuse off, confirming ork_regs.h:139 by removing the reader of the retained bytes rather than by inference. It does not pay — best-with-reuse 1767 µs beats best-without 2071 µs (6:512:3, 12 programs), and reuse is worth 1.31× at the default geometry — but Wb and the reuse flag are a pair that nothing enforces: orki_i4_bchain_try picks Wb with no reference to ORK_I4_WREUSE, so setting ORK_I4_WREUSE=0 silently forfeits ~1.6× of legal Wb headroom, and a Wb raise validated with reuse off would miscompute the moment reuse came back on.

1. The WRONG points were the FASTEST in the sweep (1714–1776 µs). A timing-only sweep would have crowned Wb=448 as the optimum. Combined with the known RNG trap — a patterned generator hides this corruption from a checksum — neither gate alone is close to sufficient, and the failure mode is not the slow self-heal the existing tooling is built to catch. It is fast and wrong.

2. Geometry validity is CORE-COUNT-dependent, and nothing in the code models that. 2:320 is the best known config at nc=3 (1738 µs) and fine at nc=2 (1924 µs) — but at nc=1 it took 259 930 µs, a 150× self-heal storm. Same K,N,M,H,Wb; only the per-core chain length changed (56 programs on one core against ≤24). orki_i4_bchain_try chooses Wb from K alone, so any future Wb table keyed on (K), or even on (K,N,M), is unsound. This is a live wedge risk independent of everything else here.

2:320:3 at 1738 µs against the shipped default 3536–3668 µs is a 2.03× bit-exact single-matmul win, already expressible with env knobs. The shipped Wb = (131072/K)&~63 is 5× below its legal ceiling at K=2048, and the default 0x1040 = 0xb1 (DATA=1) starves H. The blockers to making it a default are the two hazards above, not the speedup.

§8–9 reasoned about DMA volume. This section measures where the time goes, after a challenge that the DMA accounting was comparing non-fungible quantities. It was. Board 10.3.0.236 · 6.1.115-vendor-rk35xx-fencefix · whole-process MEM_SYNC census over both orki_bsync and orki_bsync_off · full-random data · new probe tools/re/i4_sync_prof.c.

24.0–24.9 MEM_SYNC ioctls per matmul call, 16 138 KB/call, against a per-buffer prediction of 15 684 KB — residual 454 KB. Weight sizes for the three shapes are 2048 / 6144 / 6144 KB; nothing of that order appears. Every weight orki_bsync is in i4/pack.c, at pack time. Nothing to win.

NPU execution is 72–73%. The host activation nibble-pack — the term §9 credited with deciding the transposition A/B — is 1.3–2.2%. Host regcmd synthesis is 11–14%, four to eight times the nibble-pack, and nobody had accounted it.

Neither operand's DMA sets the clock. Weight bandwidth is 0.55–0.85 GB/s against a ~30 GB/s DRAM ceiling — two orders of magnitude off. The NPU time is per-program dispatch. Decisive: at n=2/Wb=320 the program count drops ~7× and the NPU poll falls 2630 → 1357 µs (1.94×) while the same weight and activation bytes move.

So §9's conclusion survives and is strengthened, but its stated cause was wrong by an order of magnitude — the host asymmetry is real and is not a clock-setter at these M.

⚠ Unresolved tension, recorded rather than smoothed over. §9's transposed A/B had program counts matched at 32, so program count cannot explain that particular 536 µs delta, and §9's host-pack attribution (361 of 488 µs) remains the only account of it. The two findings are consistent if read precisely — the nibble-pack is small in absolute terms across the shapes that matter, while still dominating the delta of a matched-geometry swap at M=2048. The claim "program count decided that A/B" is not supported for the matched case. Transposition stays refuted either way; only the attribution is open.

Reframing the Wb win (K2048 N2048 M=128, shipped → n=2/Wb=320): wall 3625 → 1853, poll 2630 → 1357 (−1273), synth 401 → 43 (−358), de-tile 243 → 151. ~72% of the saving is NPU-side, ~20% is host regcmd synthesis — the latter unaccounted until now.

The premise that only a multi-shape run could expose the whole-buffer orki_bsync overage is wrong. The four scratch buffers are not merely high-water marked, they are pre-sized:

• ork_npu_init → core/sched.c:113-114: mcc 2 MiB per core, unconditionally
• orki_mc_ensure → core/sched.c:380: mrc 64 KiB, mtk 64 KiB, maf 256 KiB per core

Both deliberate, to dodge a real bcreate EINVAL under import pressure — so shrinking the allocation is the wrong lever; orki_bsync_live is the right one. Consequence: a single-shape process already overpays. Measured process-level overage: K2048 N2048 2.84×, K6144 N2048 2.36×, n=2/Wb=320 2.97× (mcc 3.20×, mtk 16×, mrc 4×). Only K2048 N6144 is ~1.0×, its live bytes exceeding the floors. Multi-shape high-water adds on top (maf 1.81× / mrc 1.39× / mcc 1.84×).

Prior sweeps missed this because they never looked at sync cost — not because the shape was single.

Same binary both arms (ORK_BCH_LIVE), arms alternated, order flipped:

Order-independent (the live arm won from first position too, 1675 vs 1786). MEM_SYNC traffic 21 764 → 9 864 KB/call. Bit-exact — fnv64(C) identical across every arm and both geometries under full-random weights and activations.

It compounds with the geometry change, because the cost is absolute while the call gets 1.9× shorter: 1.6–2.4% shipped, 3.4% with skew, 8.3% at the geometry worth shipping.

Sites — core/dyn.c:91 is the largest and was missed: mcc FROM_DEVICE in ork_dyn_done_i, live h->nout[i]*2, ~100 µs of the 152 µs wall delta at the optimised geometry. Plus i4/chain.c:371 (maf), :395 (mrc), :399 (mtk), :421 (mcc seed). core/dyn_ctl.c:393 is already orki_bsync_live. Same pattern unmeasured at chain.c:581/601/605/612/628 and :911/:983 (MoE expert chains).

_sw_t0 (core/domain.c:217) starts after orki_dom_drain(c) and ork_dom_flush_if_dirty(c), so it reports 0.97 µs/swap. Timing the whole of orki_dom_activate:

ndom=2, 20 switches:  410.7 us/swap = drain 300.7 + flush 0.3 + park/restore/first-touch 109.7
ndom=3, 29 switches:  396.7 us/swap = drain 288.8 + flush 0.3 + park/restore/first-touch 107.5

Steady-state ≈ 290–300 µs/swap, essentially all orki_dom_drain — three blocking sacrificial f16 matmuls, one per core, plus an Af bsync. ork_dom_flush_if_dirty is 0.3 µs, not a term. End-to-end cross-check agrees: +332 µs/swap (ndom=2), +197 µs/swap (ndom=3).

Anyone trusting the existing hook would conclude domain switching is free. It is ~300 µs.

Weight in the end-to-end estimate: at today's ~7.2 ms/matmul, a ~7-matmul layer crossing 1–2 domain boundaries pays 0.6–1.2% — noise, so it does not undermine §8's 1.5–1.7× extrapolation. But at the optimised geometry a matmul is ~1.7–1.9 ms and one swap is 16–18% of it. If a domain layout made swaps per-matmul rather than per-layer, domain switching would eat most of the geometry win. Worth knowing before laying out domains; not worth optimising now.

1. NPU poll, 72% — program count; largely banked by the Wb/n geometry change.
2. Host regcmd synthesis, 11–14% — falls out with the same change (9.3×).
3. De-tile, 3.4–8.1% — single-threaded over all cores on the calling thread, untouched. 858 µs on K2048 N6144. The largest unexamined item.
4. Cache maintenance, 2.6–5.2% — orki_bsync_live takes about half.

Not worth pursuing: weight zero-copy work (already zero-copy); attacking the nibble-pack (1.3–2.6%); tuning orki_dom_drain before knowing whether the target model swaps per-layer or per-matmul; the submit ioctl (0.15%).

Every number in §7–10 was matmul-level on synthetic shapes. This is the first model measurement of this line of work. Board 10.3.0.236 · 6.1.115-vendor-rk35xx-fencefix · governors verified before and after · Qwen3-1.7B native W4A4 (ORK_MIXED_DISPATCH=1 ORK_MIXED_W4A4=1 ORK_QUANT=4, pack 1.34 GiB / 193 resident weights) · ork_bench, FA off, warm · arms differ in exactly one file, libggml-ork.so, sha256-verified distinct.

Order-independent: BASE-first 45.54 vs BASE-second 45.38; OPT-first 51.13 vs OPT-second 50.85 — spread ≤0.6%, far under this board's 1.126× arm-position bias. The two ubatches agree to 0.7%.

Decode is exactly 1.00 tok/s on all eight runs, both arms, as the code predicts: dense M==1 is declined by supports_op, so decode never reaches the NPU and BCHAIN (which needs M≥2) cannot touch it. PPL matches to every digit, per-window too — bit-identity confirmed at the model level.

ORK_PROFILE=1, 2 reps alternated:

The entire saving sits in t_run, exactly where the three phases live; t_quant and t_deq are flat to noise, and the accounting closes (wall delta 5143 ms vs t_run delta 4616 ms).

On real model shapes the run call is 1.3499×, against the microbenchmark's 1.209× — the microbenchmark UNDER-stated the lever by 1.12×. So the gap between 1.209× and 1.128× is not decay of the optimization. It is Amdahl.

~61% of prefill wall was never addressable by these three commits. The int4 matmul call was 39% of BASE wall; the commits took 26% out of it = 10% of wall = 1.128×. That is the whole story.

Seven shapes, all at M = ubatch exactly — no M spread at all:

K=2048 is 79.1% of W4A4 run time but only 25.7% of prefill wall. K=6144 (ffn_down) is the sole K≥4096 shape at 6.8% of wall; there is no K<2048. Every NPU matmul in this model's prefill is int4 — 1544/1544, zero int8 — and output.weight never reaches the NPU.

Calculated from the measured census (the n=2/Wb=320 geometry was deliberately not run, as it carries the unresolved storm):

Realistic payoff is 1.05–1.09× model prefill, and it bids for a share of what is left: the 2.03× was sized on the pre-optimization tree, and these three commits have since taken 1.3499× out of that same call. Against that, it carries an intermittent self-heal storm (4.9 ms → 60/116/228 ms, bit-exact but 12–46× late) — a board-wedging failure mode — for a single-digit-percent model win.

K-slicing ffn_down 6144 → 3×2048 widens the lever's domain from 25.7% to 32.5% of wall, a ×1.26 multiplier on whatever the lever is worth, before paying host-side int32 partial-sum accumulation. It is not what makes the lever pay.

The W4A4 activation rotate + quant is 33.1% of prefill wall — as large as the entire NPU run — it is pure host code with no NPU risk, and it is completely unattributed: ORK_PROFILE's own ACTQ breakdown reports arith 0 ms / resolve 0 ms / misc 26824 ms. That is precisely the state regcmd synthesis was in before instrumenting it yielded 18.7× on its host term.

• cp -a of the fork tree carries llama-pinned's absolute CMake paths, so the first build silently compiled the original tree and produced two byte-identical .sos. Caught only by the sha256 check; a fresh cmake -S/-B fixed it. Always hash the artifact that differs between arms.
• ggml-ork.cpp's W4A4 header comment is now stale — it states a flat 3.6× vs W8A8 and NPU run 5840 us/matmul, both predating these commits, so it is ~1.35× pessimistic. Re-measure or date-stamp it before landing.

§9–11 treated the n=2/Wb=320 self-heal storm as a tail risk on a geometry lever. Root-causing it on .235 found something larger and unrelated to the lever. Board 10.3.0.235 only (uname -r = 6.1.115-vendor-rk35xx-fix, DMC performance@2400, A76 @2.4 GHz) — timings not comparable to .236.

The live sentinel is still ORK_DYN_SENT16 = (int16_t)0x7fff (internal.h:210), adjudicated by orki_dyn_collision_landed (core/dyn.c:70). Direct correlation on every storming call:

sat+ = 0    sat- = 0    max|C| = 183        <- not one element within 180x of the sentinel
"equal the int16 seed" in every stderr log:  0 occurrences

Positive control — forceable on demand, at a geometry this die runs cleanly (H=4, Wb=64), 8 calls each:

So the mechanism is real, reproducible on demand, and removing saturation removes it — the adjudicator cannot rescue a 100%-saturated surface by design (≤0.1% survivors). It is simply not reachable with real data: at K=2048 with codes in [−8,7] the accumulator is mean ≈512, σ≈970, so 32767 is ~33σ.

My hypothesis was wrong about this storm, and right about the mechanism existing.

.235's measured BCHAIN envelope at K=2048, ladder off, checksum-graded:

H_max = 4, where orki_i4_hcap(2048) ships 8 — 2× too high. Wb_max = 64, i.e. exactly the ladder's minimum rung, so this die has no Wb headroom at all, lever or no lever.

Consequences, and the second one is the serious one:

1. The storm on .235 is not intermittent — it is 100%, and 100% at the SHIPPED DEFAULT geometry (60/60 at the default, ~5.2 s/call against a 3.5 ms correct runtime, ~1000× late). .235 is in a different, deterministic regime from .236's ~1-in-6, so it cannot serve as a higher-rate instrument for .236's storm.
2. .235 has been silently paying ~5.2 s per int4 prefill matmul all along. The geometry ladder rescues every call to a bit-exact answer, so nothing ever reported an error. Any .235 int4 benchmark taken before now measured the ladder, not BCHAIN.

Chain of events on a failing call: all three cores submit NONBLOCK and are accepted (mc_rc=[0 0 0]), the hardware executes only a small prefix (~87.5% of the cells it should write are still at the seed, hw_elapse=0, int_status=0x0 — the registered over-ceiling signature), ork_dyn_end declares a miss, orki_mc_recover_resubmit runs 6× (ACT_RESET + full-surface re-seed + resubmit), recovery exhausts, and the geometry ladder retries at Wb=64 then H=4 and succeeds. ~9 soft resets per call.

• The EWMA miss-timeout key omits the geometry. miss_to comes from orki_op_predict_us(pkey) keyed on (dt, 2, N, b_M, S) — no H, no Wb — so a geometry that works and one that cannot share a key. With the ladder on, the successful H=4 retry trains it to ~4 ms and the window clamps to its 50 ms floor; with the ladder off nothing succeeds, est stays 0, and the 2 s CATCHALL applies (6×2 s = 17 s).
• The 50 ms miss floor is below one poll pass. Each recover round reports waited 102262us/50000us — one pass over three cores' surfaces costs ~50 ms, so miss detection cannot resolve finer than ~15× the op's correct runtime.
• The BCHAIN output scratch is 75% padding the NPU never writes (slot is 4*H*Wmax*64 int16, only Hg*NBc*64 written). orki_mc_recover_resubmit's esz==2 branch re-seeds all nout, so after any recover the padding is permanently 0x7fff, and orki_dyn_collision_landed scans the whole surface against a 0.1% cap — leaving the adjudicator structurally unable to pass on this path (92–94% survivors after a re-seed, ~75 pp of it padding). Latent; did not bite here.

The sentinel collision (above); ORK_I4_WREUSE (off = identical); chain length — at H=4, Wb=64, nc=1 the whole 1024-program chain runs on ONE core, 20/20 clean at 7615 µs, so "56 programs on one core storms" is not a chain-length law; job accounting (unpow=0, nojob=0, no "NOT dispatching"); and the DRAM clock — H=4 passes at both 2400 and 1968, H=5 fails at both, independently confirming the 2026-10-12 OPS_REGISTRY correction that the DMC term governs only the tiny discriminator shape.

H=4 × 60 calls: 60 ok / 0 fail (med 3556 µs); H=5 × 12: 0 ok / 12 fail; H=4, Wb=320 × 15: 0/15. No stochastic band exists on .235, so a one-shot checksum-graded probe per (K, shape class) at pack or init time classifies it deterministically. Which also means .236's intermittency at H=16 is either that die sitting exactly on its own ceiling, or a different mechanism — untested, since .235 could not be used as a proxy.

A fix would additionally have to address: the ladder can only lower to the shipped rungs, so it cannot express "this die needs H=4 though the table says 8" without paying the full failed attempt every call; pkey should include the geometry; the recover re-seed should cover only the written cell set; and the 50 ms floor should scale with surface size.

An over-ceiling H=5 run left core 1 pinned at 99–100% with zero render-node holders (blocked_slow 503, blocked_age 44.4 s); the 6× ACT_RESET recover did not clear it and it needed a reboot plus governor re-pinning. Expect the guard to block ~120 s on the fault count after a storming run. Also, pgrep -f <script> from an ssh one-liner matches the ssh command itself and killed the agent's own session — the liveness-checks-lie trap again. Target by pid.

§11 named the CPU-side W4A4 activation rotate+quant as the largest addressable term (33.1% of prefill wall) and "100% unattributed". Instrumenting it first — before optimising anything — found the cost in a different place than anyone had guessed, and a broken thread schedule nobody was looking at.

Board 10.3.0.236 · 6.1.115-vendor-rk35xx-fencefix · governors verified · Qwen3-1.7B native W4A4 · fork commit a26c9d5ad on diverged-history (sole file ggml-ork.cpp; no ork-driver src/ change, so no attest refresh).

t_actq has exactly two += sites and both are on the int8 path. On W4A4 it was never incremented at all — which is why ORK_PROFILE read arith 0 / resolve 0 / misc 26824. Not a mis-placed timer, a missing one. That is the third piece of lying instrumentation found in this investigation, after ORK_DOM_PROFILE's 420× undercount and the Wb sweep whose fast points were the wrong ones.

wall 26659ms over 1544 calls (440 reuse-hits, 1104 computed)
1.611 Gelem   8.06 GB   16.55 ns/elem   threads=8
cpu 127855ms = memcpy 850 (1%)  smooth-div 47 (0%)  rotate 16864 (13%)
               scale-search 101789 (80%)  quantise 8305 (6%)  | par-eff 4.78x of 8

The rotation was never the problem. It is 12–13% of the term, and ork_fwht_norm — scalar-looking C — already compiles to a 4-wide butterfly at -O3, running ~1 cycle per element-pass. Irreducible.

★ Two findings that matter more than the table:

1. This is a pure-ALU term, not a memory one. 1.611 Gelem = 8.06 GB over 26.7 s = 0.29 GB/s, with 8–24 KB rows staying L1-resident across ~20 passes. So "move fewer bytes" — the standard lever in this codebase — is the wrong one here. An explicitly fused single-pass form measured slower (12.3 vs 7.75 ns/elem) on register pressure.
2. The thread schedule was broken. par-eff 4.78 of 8. The loop is #pragma omp parallel for with the default static schedule across 4×A76 + 4×A55. Solving 4(1+k)/k = 4.78 gives k ≈ 5 — the little cores are ~5× slower on this code, and with an even row split every big core idled ~80% of the loop.

NEON clip search (micro 40.0 → 7.75 ns/elem, 5.21×, reduction-order only — FRINTN is exactly lrintf's round-to-nearest-even); NEON quantise (5.61 → 0.36 ns/elem, 15.8×, byte-identical); and schedule(dynamic), which was worth more than the vectorisation and is bit-exact by construction since rows are independent.

End-to-end, alternated BASE/OPT/OPT/BASE: prefill 50.28 → 76.62 tok/s = 1.524×. Activation prep 27077 → 4689 ms = 5.77×. Decode unchanged at 1.00 (this path is prefill-only). Chunk sweep: dynamic,1/4/8 tied, ,16/,32 degrade; static with OMP_NUM_THREADS=4 reaches 72.6, independently confirming the A55 straggler.

★ The A/B method is worth copying. Both arms are the same binary (sha256 cceec203…), selected at runtime by ORK_ACTQ_SCALAR=1 OMP_SCHEDULE=static. That eliminates the cp -a/pinned-CMake-path hazard from §11 rather than guarding against it — with one artifact, two arms cannot silently be the same code.

Activation digest 025ce13df13cc9ed over 805 306 368 quantised elements, identical BASE and OPT across three builds. ork_ppl 52.1747 both, mean NLL 3.954598 to six digits.

Stated plainly rather than rounded to zero: a 90 000-case kernel sweep found 2 mismatches, both on unrotated heavy-tailed synthetic rows — not 1-ulp wobbles but alphas two grid steps apart, because on such rows the MSE curve is flat to the last bit and a 1e-16 reduction-order difference moves the argmin. Every rotated row matched, and the production path always rotates. The conservative half was taken anyway: ork_w_scale_mse (the weight clip) is reverted to scalar — it runs once per weight at pack time, vectorising it bought nothing measurable, and a .orkpack is persisted and compared across builds, so it does not get to wobble.

• float32 accumulator in the clip search — micro says 10.2× instead of 5.21×, ≈1.03× end-to-end. It changes the formula, not just the reduction order, and the 2/90 000 result shows the argmin can flip. Not worth the exactness risk for 3%.
• Overlap with the NPU — the most promising structural idea going in, now largely removed by its own success: act-prep is 2.34 s against 13.6 s of NPU run per pass, so perfect double-buffering caps at ~1.10× for an architectural change to a synchronous ggml op.
• Activation-cache reuse — 440/1544 (28.5%) against a structural ideal of 3-in-7 (42.9%); the gap is last_abits being in the reuse key while the pack promotes some weights to 8-bit activations. Recovering it means caching both a 4-bit and an 8-bit quantisation of the same row, worth ≈1.8% of wall. The next item in this area if anyone returns to it.

The host side of W4A4 prefill is no longer the problem. The NPU run is, at ~51% — and the ~38% non-matmul remainder has never been broken down and is now the second-largest term in the model.