A running blog of NPU experiments for the open inference stack (ork-driver → ggml-ork → llama.cpp). Each entry is a short writeup with a benchmark table; click through to the full page for the hypothesis, the actual code (commit + symbol + env-knob), the method, and the verdict.

Ground rule: every entry records the implementing code so a "ruled-out" approach can be re-run later (bigger model, new SoC, kernel update). Experiments are kept runnable behind env flags rather than deleted. All perf numbers come from the standardized same-model, warm, governor-verified benchmark — never extrapolated. See Benchmark Standards & Methodology for the exact procedure and env knobs.

Reference model: Qwen3-1.7B-w8a8 on RK3588 (10.3.0.236), DDR 2112 MHz / A76 2.4 GHz, llama-bench -p 512 -n 128 -r 2, warm. Closed-runtime reference: librkllmrt v1.2.3 on the same .rkllm model — prefill 183.9 / decode 11.71 tok/s (rknpu_bench, fresh 2026-06-13). Every applicable table carries that rkllm row; librkllmrt has no env knobs, so it is a single fixed point, not an A/B arm. All tables are tok/s (pp512 = prefill, tg128 = decode) unless noted.

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

Native W4A4 (DT_I4_NATIVE) runs prefill 3.50× slower than i4a8 on the same model, and the cause is not what the source claimed. Submit counts are identical — 1155 ioctls on both paths. The int4 datapath packs 167× more chained programs into the same submits (1,946 vs 325,696), against a 48× lower per-program cost (6.9 µs vs 328 µs), which is why 855× more programs per weight nets only 3.50× of wall. A tiling model programs = Σ ceil(M/H(K)) × ceil(N/Wb(K)) predicts the measured count to 1.3%, leaving nothing for a hidden per-call effect.

Decomposing the inflation: Wb N-tiling at 64 columns contributes ×43.1, H contributes ×18.7 — so the hcap row-grouping ceiling is a real mechanism (a fixed-N control measured 2.69× against the predicted 2.67×) but the smaller term. Wb is pinned at 64 for every K ≥ 2048 and is where the headroom is. ORK_I4_NOROT=1 returns identical geometry, closing the Hadamard rotation as a suspect by measurement rather than by argument.

Follow-up, same day: Wb can be raised — 2.21× on a single matmul — but the pair that unlocks it is H, not the weight-bank count (the shipped 0x1040 = 0xb1 is already DATA=1/WEIGHT=11, i.e. maximally weight-biased, so pairing can only shrink the weight side). Best valid point n=2 / Wb=320 at K=2048, 1707 vs 3626 µs, 56 programs against 512. The lever works at K ≤ 2048 only — at K ≤ 1536 no split change is needed at all (the shipped (131072/K)&~63 under-uses by 5× at K=1024 and 7× at K=768), and at K ≥ 4096 it is dead, because 0x1040 also carries the K-reduction schedule and only the captured 0xb1 survives there. Three predictions were written down first and measured against: maximising n(12−n) is the wrong objective (the binding term is activation DMA M·N·K/Wb, a function of Wb alone); transposition is net negative (the weight already streams once CBUF-resident, and nc ≤ ceil(M/Wb) would cap a transposed prefill at 2 cores); and K-slicing below 2048 is 4.5× slower, though slicing to 2048 to reach the lever pays 1.24×. Expect ~1.5–1.7× overall on a real model — worth taking for capacity, not speed parity.

Transposition, tested and refuted (§9). The follow-up idea — that the two reuse bits are one capability attached to whichever operand holds the weight role, so transposing would put the small matrix under the validated WEIGHT_REUSE[13] and never need the corrupting DATA_REUSE[12] — is expressible on this path and measurably slower: at matched split, Wb, programs and cores, 3106 µs untransposed against 3642 µs swapped, while moving half the bytes. Three reasons, the third new: the activation is nibble-packed so the DMA advantage was overstated 2×; these shapes are not DMA- or program-count-bound (a 16× swing in program count costs 37%); and decisively, the roles are not symmetric in software — the weight role is packed once, the activation role is nibble-packed on the host every call, so transposing moves the big matrix into that role for a 16× rise in per-call host work (23.9 → 385.0 µs), which is 74% of the measured deficit. Confirmed in the same run: the Wb WRONG boundary vanishes entirely with ORK_I4_WREUSE=0, so the ceiling is reuse staleness, not containment — proven by removing the reader of the retained bytes. ⚠ Two live hazards fell out: the WRONG points were the fastest in the sweep (a timing-only sweep would crown an incorrect Wb, and the failure is fast-and-wrong, not the slow self-heal the tooling catches), and geometry validity is core-count-dependent — 2:320 is best at nc=3 but self-heal-storms to 259 930 µs at nc=1, so any Wb rule keyed on K alone is unsound.

Where the time actually goes (§10), after the DMA accounting was challenged. It was comparing non-fungible quantities — weight bytes are a pure NPU-side read, activation bytes additionally carry a host nibble-pack and host cache maintenance. Measured attribution at M=128/nc=3: NPU execution 72–73%, host regcmd synthesis 11–14% (four to eight times the nibble-pack, and previously unaccounted), de-tile 3.4–8.1%, cache maintenance 2.6–4.2%, nibble-pack 1.3–2.2%, submit 0.15%. Neither operand's DMA sets the clock — weight bandwidth is 0.55–0.85 GB/s against a ~30 GB/s ceiling, and at n=2/Wb=320 the NPU poll falls 2630 → 1357 µs while the same bytes move. Two pieces of instrumentation were found to lie: the grow-only scratch also has fixed pre-allocation floors (mcc 2 MiB/core), so even a single-shape process over-syncs 2.4–3.0× — switching i4/chain.c and core/dyn.c:91 to the existing orki_bsync_live is bit-exact and worth 1.6–2.4% shipped, 8.3% at the optimised geometry, compounding because the cost is absolute while the call gets 1.9× shorter; and ORK_DOM_PROFILE undercounts domain swaps ~420× (its timer starts after the drain barrier), reporting 0.97 µs against a real ~300 µs/swap — noise per-layer today, but 16–18% of a matmul at the optimised geometry.

Same page: the reset toll on mixed-precision packs (~105 ms per weight, 221 resets), the fix being arrangement not policy (contiguous promotion 221 → 4 resets, 5.57×), scatter beating contiguous 1.92× once resets are free, and the ladder's real cost — GPTQ is free and worth 12%, SmoothQuant costs 7% speed and 3.8% PPL on this model. Three retractions are recorded inline, including a NORESET result that did not exercise the hazard it appeared to clear. → full writeup

Board 10.3.0.236 · kernel 6.1.115-vendor-rk35xx-fence · governors verified · AFTER arm ran FIRST

orki_bsync() syncs b->size, and the per-core colsplit scratch (maf/mrc/mtk/mcc) only ever GROWS — so every call paid DMA cache maintenance over the widest shape the process had ever run. Measured with a new ORK_CS_PROF=1 phase profiler: 1.47 GB/call, 24.6 ms of a 30.4 ms call, against a 3.5 ms submit (59.6 GB/s — a DRAM-rate figure, so the cost is bytes, not ioctls). This is the fixed ~24 ms/call that (I) below left open, and it refutes all three suspects named there: pool dispatch+join is 0.024 ms/call (the pool is already persistent — there is no per-call pthread_create), the Sk==1 copy-back is 1.838 ms, and the "A gather" cost is its trailing whole-buffer orki_bsync, not the memcpy.

Three cores now BEAT one by 1.21x where they lost by 1.49x, so the multi-core default stays and (I)'s ORK_NPU_MC=1 advice is retired. Byte-identical output: test_matmul's static fp16 golden back to the pristine value, test_f16colsplit bit-exact, the training gradient table and loss unchanged to the digit, --break rms still failing, make test ALL PASS. The same fix on the SINGLE-CORE path is not yet viable: it is worth another 5% on the 1-core arm but it changed the fp16 cold-run output, caught by that same golden. → full writeup

Board 10.3.0.235 (ROCK 5B+) · kernel 6.1.115-vendor-rk35xx-npu = armbian + all 13 PR #390 commits · ork-driver 1.0.165

Closes the first open limit of entry (III). rknpu_job_timeout_clean() reaps the RUNNING job and drains the rest of that core's todo_list into schedule_work(&job->cleanup_work) — but rknpu_job_free() never looked at job->dom_held, so every queued job it frees took a domain reference with it. The trigger is ork-driver's own ork_npu_reap_stuck(), whose NONBLOCK fp16 dummy queues (rather than runs) when the stuck core's job is not yet past its timeout, and which the BLOCKING drain behind it does not retire.

Per-callsite counters (rknpu.dom_dbg) plus the fix behind rknpu.dom_fix, one kernel, two arms, full make test each:

[leaked] == [queued reaped] == the reclaimed refcount. The fix is three lines in rknpu_job_free(); ORKD_DOM_API passes in both arms, and in the second it passes with the reclaim backstop disabled. Branch rknpu-domain-ref-leak in oRKLLM/rk3588-kernel, off rknpu-iommu-rework; keep #patchB5 as the backstop.

Also settled here: TEST_SLICE_RESCUE / the four failing i4 shapes reproduce identically on the stock armbian kernel, so they are a .235 board property, not PR #390 and not our kernel work. → full writeup

Board 10.3.0.235 (ROCK 5B+, NEW) · kernel 6.1.115-vendor-rk35xx-npu = armbian + exactly PR #390's ten commits

First time the clean upstreamable series has been booted and exercised rather than only compiled — .236 could never show this, because it runs #59 with ~49 extra hunks that silently include the fix.

The mechanism is in PR #390's own comment: a switch proceeds only when iommu_domain_refcount reads exactly zero, and the PR clamps the counter at zero but cannot recover a count already too high — so one leaked reference wedges every later switch and, via rknpu_gem_object_create(), every later allocation. #patchB5 (Change 19) reaps the cores on timeout and, if nothing is in flight, forces the leaked count to zero and retries once. 98-line patch attached, no debug counters. Two limits kept explicit: B5 treats the symptom — the leak itself is unexplained and still open; and a single isolated pass proves nothing, because orkd_dom_api passes on the unpatched kernel too when run alone on a fresh boot (the wedge is created by the preceding tests). → full writeup

Board 10.3.0.236 · kernel 6.1.115-vendor-rk35xx-fence · ork-driver 1.0.160 → 1.0.165 · fork e32561c79

ork_bench on qwen3-1.7b-f16.gguf (P=128/G=32/ubatch=128, 3 reps, warm, governors verified, new arm run first i.e. disadvantaged):

+0.4% / −0.2% — inside the ~2% screen and far below the arm-position bias band. A tie. The reason is dispatch, not the fix: ORK_PROFILE shows run_i8: 501 calls and nothing else, because ggml-ork routes even an f16 GGUF through the int8 datapath, so ork_f16_mm_* is never reached. The 2.3–5.0× measured on fp16 shapes in isolation is real but lives in the training path, not here. The bump is therefore safe and neutral. Caveat recorded: ork-driver has no fp16 profile counter, so "zero fp16 calls" is an absent line rather than a measurement — hence the timing A/B was run anyway. → full writeup

Board 10.3.0.236 · kernel 6.1.115-vendor-rk35xx-fence · ork-driver PR #25

orki_f16_sched's window is [128,2048) and the SoC's fp16 K-slice was 2048, so every full fp16 K-slice landed on sched=0, whose M-cap is one CBUF bank of activations — 8 rows. A submit re-streams its whole weight slice regardless of row count, so a tied head's dX issued 15616 submits × 4 MB = 64 GB of weight DMA (64 GB / 7.9 s = 8.3 GB/s, exactly the single-core DRAM weight-stream rate). Slicing at 1024 gives 65.9 → 330.5 GFLOP/s (5.01×) on that shape, 2.3–2.5× at ordinary K, with three Sk=1 controls flat at 1.00–1.01×. Full Qwen3.5-0.8B training step 112.626 → 30.133 s (3.74×), 9.1 → 34.0 tok/s, worst gradient tensor improving 3.351e-03 → 2.451e-03. Also recorded there, and not acted on: the fp16 multi-core colsplit is a net loss on that workload (one core 1.49× faster than three; a fixed ~24 ms/call on the 3-core path, located but not root-caused). → full writeup

Board 10.3.0.236 · kernel 6.1.115-vendor-rk35xx-fence · ork-driver PR #15

The int4 NONBLOCK doorbell infers completion from "the output no longer equals the sentinel we seeded". Sound for int32 (INT32_MAX is unreachable); not sound for int4, whose output is int16 against a seed of 0x7fff — a value a saturating W4A4 accumulator reaches easily (K=3072 with codes in [-8,7] runs to 196,608). One legitimate output element equal to 32767 made the op look permanently stuck, the recovery resubmitted a round that had already landed six times, and the decode failed.

The ORK_MC_DIAG census had the answer all along and mislabelled it:

sentinel 1/81920 (0.00%) | first=1 last=1 runs=1 => ONE-RUN(stalled write)

One element out of 81,920, at index 1 — classified a stalled write because the classifier keyed on the run count. A stalled write leaves a run that reaches the end; a single short run near the front is a collision wearing a stall's shape.

Four signatures said collision and none said dropped round: deterministic (always the same ffn_down, after exactly 55 of them), gone when the calibration text changed, gone under ORK_I4_BLOCKING, and blocked_slow/unpow/nojob all flat — the kernel was never involved. Per-call instrumentation confirmed everything reaching the NPU was well-formed (scales finite, codes in [-8,7]). Three theories died first: starved WEIGHT_BANK, a kernel blocked-dispatch stall, degenerate SmoothQuant scales (all 140 vectors checked clean offline, worst ratio 28.5).

Fix: promote the census from a printout to a decision — after the spin times out, re-read through the kernel's DMA sync and accept only if survivors are ≤0.1% of the surface, loudly. It had to go in the i4batch branch as well as esz==2; the BCHAIN prefill path returns before the latter, so the first cut changed nothing. Validated bit-identical to the blocking path (PPL 112.0663 both ways).

This is latent for any W4A4 workload — the seed is in-band for every int16 output. The sound long-term fix is a completion signal outside the output surface, which needs a regcmd change to the captured int4 template.

Page: Exp-2026-10-08 — fp16 Wide-K Head Hang Board 10.3.0.236 · kernel 6.1.115-vendor-rk35xx-fence · base origin/main eb3e66f

ork_f16_mm_run at M=1024 K=248320 N=1024 (a tied output head's dX) committed a job that never raised a completion interrupt; core 0 stuck at 99%, board power-cycled. The suspicion was that r103's starved-WEIGHT_BANK mechanism was not limited to narrow N once K is sliced. It was not that at all.

orki_run computed its M-chunk twice — once to size the Af/Cc staging buffers, once to emit the programs — and the two formulas had drifted. The emitter moved onto the measured fp16 envelope (orki_f16_mcap_n); the sizing pass kept the legacy int8 mg_max*64. Where fp16 is the taller of the two (Kp=512: 352 vs 320; Kp=1024: 176 vs 128) the NPU was told to write 131072 bytes past the end of Cc. Walking off a mapped IOVA region on this hardware is not a fault — it is a hang. K=248320 at the SoC's KS=2048 leaves a 512 tail, and that is the slice that died: the reported nout=360448 is exactly 352 × 1024, which identified it before any board run.

Found entirely offline. The suite missed it because fp16 M>1 at N>=32 goes multi-core → colsplit, which reads the same cap for planning and sizing; the single-core run_loop is reached only when colsplit declines, i.e. only at output-head widths.

Fix: one function, orki_run_loop_chunk, called by both passes; an extent guard (orki_check_extent) at all four emit sites so future drift refuses instead of wedging; the int8 full-K branch's Cc sized explicitly (it was covered by luck); ORK_F16_HISCHED removed (its only reader was the sizing pass, so it could only create disagreement). New test test_f16_wide_k in make test. Lesson: a cap enforced in two places is a cap enforced in one.

Page: Exp-2026-10-07 — W4A4 vs GGUF 4-bit, Part II Model qwen3-0.6b-f16 (dense) · target = its own Q4_K_M @ 5.09 BPW · 64 × 256-tok windows · ork-driver 1.0.154

ORK_QUANT=4 alone does NOT give you W4A4. Without ORK_MIXED_W4A4 you get int4 weights inflated to int8 through mul_mat_i8 — effectively W4A8 — with no rotation and no GPTQ registration, and the run writes a plain RTN pack while exiting 0. That flag explains every anomaly in Part I.

★ 2026-10-08, MIXED-RADIX ROTATION — 1.7% on seen text, 0.33% held-out. A null where it counts, and the clearest demonstration of why the held-out eval had to come first.

b = K & -K, so the FWHT was applied block-diagonally. At K=3072 (= 2¹⁰·3) that left ffn_down — the largest weight, carrying the top qerr — rotated only within three independent 1024-channel groups, unable to spread an outlier beyond its own third. Every other shape here is a power of two and already got a full-K rotation. Fixed by Kronecker: R = Q₃ ⊗ H_b, i.e. FWHT each block then a 3-point orthogonal ACROSS blocks at each position, O(K·m). ork_w4a4_rot_k; ORK_I4_NOKRON=1 reverts. Six call sites (3 weight-side, 3 activation-side) had to change in lockstep — a basis change silently corrupts if the two operands disagree.

Verified before measuring: dot products preserved to 6.1e-05 over 3072 terms with planted outliers (so ⟨Ra,RW⟩ = ⟨a,W⟩ holds), and energy escaping block 0 went 0% → 66.7%, the 2/3 a uniform 3-way mix gives. An orthogonal transform that left the blocks independent would have passed the first check and changed nothing, which is why both were needed.

A five-fold divergence between seen and held-out. The mechanism is clean: the rotation sets the basis GPTQ's Hessian is accumulated in, so better mixing gives a better-conditioned Hessian and a better fit TO THE CALIBRATION TEXT — and that extra fit is most of what fails to transfer. The transform does real work; the benefit is largely not generalisable. Treat a seen-text gain as unproven from here on.

Also recorded, since it is awkward: the best held-out number remains lad_d03 at 56.13, built with the FIXED-16 clip grid. Self-tuning cost ~1% held-out (56.13 → 56.72) and Kronecker recovered a third of it (→ 56.53). All three are inside the ~2% screen and mutually indistinguishable, so self-tuning is not shown to be harmful — but it shipped on a "tie" measured on seen text, and on held-out it sits on the wrong side of that tie.

★★ 2026-10-08, HELD-OUT EVAL — the real gap is 1.27×, and every ladder number before this was measured with our method calibrated on the evaluation text. Each arm calibrated GPTQ on ppl_screen.txt and then scored perplexity on the same text, while Q4_K_M never saw it. Four cells so the contamination is MEASURED rather than asserted — ppl_screen.txt is ~31733 tokens, both calibration (≤3072 rows) and eval (64×256) live in the first half, so the tail is genuinely unseen and same-distribution. 56 × 256-token windows throughout:

The contamination is real but modest. Ours degrades more than the baseline, the direction fitting a Hessian on the eval text predicts — but the gap widens only 1.229× → 1.273×. Most of the seen→held-out jump is that the tail is simply harder, which a two-cell design would have blamed entirely on contamination. (Prediction before the run was ~1.35×; that over-feared it.)

Use 56.13 vs 44.10 — 1.27×, about 12 points — as the headline. Relative orderings among the ladder arms stand; their absolute values carry this asterisk. Any future quantisation work here should calibrate and evaluate on disjoint text by default.

★★ 2026-10-08, THE CLIP-GRID WIN SURVIVES HELD-OUT — the one genuine quality result of the day. Same recipe, same binary, clip depth the only variable, both texts at 56 × 256-token windows:

2.4% held-out, right at the screen's resolution but on the correct side — unlike the mixed-radix rotation, which collapsed from 1.7% seen to 0.33% held-out. So the clip fix is real and the Kronecker rotation is not, and only the held-out eval distinguishes them. That is the day's methodological result: a seen-text gain is a hypothesis, not a finding.

The same build also gave the counter its strongest confirmation. On the real 140-weight pack at K=3072 the old 8-point grid pins 100% of rows against its floor — every single row wanted a tighter scale than the search could reach, and the quantiser never said so. Offline the pinned fraction falls with depth: 76.6% at floor 0.781, 71.9% at 0.656, 50.0% at 0.531, 40.6% at 0.406, and 0.0% only for the self-tuning search. Even 16 — briefly the default — left half the rows pinned. That argues for the self-tuning grid better than its PPL tie did.

★ 2026-10-08, THE PLATEAU BROKE: the GPTQ clip grid was stopping above its optimum. R 52.65 → 50.42. The MSE clip search tried 8 fractions, 1.000 down to 0.781. That floor sits ABOVE the optimum at per-row granularity — ng==1 is every production shape here, one scale spanning all K, pinned by the single largest weight in thousands. Widening to 16 points (floor 0.531):

Read this as a PLATEAU, not a peak. On the first three points (0.781 / 0.531 / 0.406) it looked like a located interior optimum, and that was written up as one. Filling in 0.594 and 0.469 showed the curve is non-monotonic and the whole 0.594-0.406 range spans 1.4% — INSIDE the ~2% screen resolution. So 0.531 is the best of five mutually indistinguishable samples, a reasonable default, but it is not a measured optimum. A cautionary instance of the project's own rule: three points drew a peak that two more points dissolved.

What does survive, solidly: widening the floor from 0.781 is worth ~4.4% (52.65 to 50.4-51.1), which is well outside the screen. Free: no extra bits, no runtime cost, pack-time only. Shipped as the default in ork-driver 1.0.159.

Invisible at the granularity we test. test_gptq runs group=32, where the optimum is already inside the old grid and widening is byte-for-byte identical. The defect bound exactly where we ship and nowhere we test — worth remembering when a knob "has a test".

And it reproduces the day's other lesson in miniature. The layer-local H-weighted error keeps FALLING past floor 0.406 while PPL turns over at 0.531, so the local objective would have picked the wrong point even here. Found via desk research into K-grouping, which flagged that a free per-channel clip search beats group-128 outright and that our grid stopped above the per-channel optimum — see K-Grouping-Cost-On-RK3588.

2026-10-08, PROMOTION at a matched bit budget — the hand prior beats qerr ranking by 3.8%. Three arms on the corrected clip grid:

t16x vs t16r is a MATCHED-BUDGET comparison — same 34 tensors' worth of bits, different tensors — and the structural prior wins by 3.8%. We already knew qerr was flat and a weak ranking signal; this is stronger: spending the same bits by qerr's ordering is actively worse than the hand prior. It also clears attn_output of suspicion — t16x contains five attn_output promotions and is the best arm, so the problem is qerr's ordering as a whole, not that tensor type.

Promotion's returns have flattened: the 6 extra tensors (+0.06 bpw) bought 0.8%, inside the screen. The headroom to Q4_K_M's 5.09 bpw is effectively spent.

A PROCESS NOTE WORTH MORE THAN THE RESULT. t16x was first reported as 58.42 — a 25% REGRESSION — and a whole mechanism hunt was built on it (attn_output promotion toxic, qerr-ranked promotion actively harmful, a bisect proposed). It was a misread: the monitoring grep was PPL = [0-9.]+ | tail -1 against a log still being written, which matches the per-window running PPL = lines, and it caught window ~30 where the running value was 57.48. Always anchor on ^\[ork_ppl\] PPL = . The same hazard had already been flagged in this session over an RHT window-51 misread, and was then repeated.

2026-10-08, ACTIVATION clip grid — a NULL, and a predicted lever that was not one. After the WEIGHT grid's floor turned out mis-set, the activation grid looked like the same defect on the side where the error actually lives (0.386 activation vs 0.169 weight at true W4A4). It is not. Against t16 = 46.55, runtime knobs so no repack:

All inside 0.7%. The reasoning error worth keeping: unswept does not imply truncated. The activation floor was already 0.5625 — inside the plateau the weight side turned out to want — so it had been set adequately; the weight grid's 0.781 had not. The two differed in VALUE, not merely in whether anyone had checked. And "this half dominates the error budget" does not imply "clipping is the lever for it": 4-bit activation noise is largely irreducible rather than outlier-driven.

2026-10-08, SIX MORE ARMS — the ladder had PLATEAUED at R=52.65 (4.0 bpw) / T=47.82 (4.89 bpw). Everything tried this day is a null or worse against those two, which is itself the finding:

Three things worth carrying forward. (1) α=0.5 is optimal — both neighbours lose, so SmoothQuant's one hyperparameter is already at its best and that knob is closed. (2) GPTQ is not replaceable but is not the only route: at a matched 4.89 bpw it is worth ÷1.078 (T 47.82 vs SP 51.55), yet S at 4.0 bpw already beats GPTQ-alone's 57.49 with no Hessian and no calibration pass. (3) Promotion is mildly COMPLEMENTARY to GPTQ, not substitutive — ÷1.069 without it, ÷1.101 with it.

The layer-local objective is a poor predictor of perplexity, now measured twice. Back-to-front ordering (Babai's nearest plane, arXiv 2507.18553) measured 1.3% better on GPTQ's own H-weighted objective and 0.45% — a null — end to end. Note also that Babai's bound holds under the no-clipping condition and we run the MSE clip grid on by default, so the guarantee never applied to our configuration in the first place. Discount any future proposal justified only by a local reconstruction metric; this is the same shape as the independent finding in arXiv 2609.01587, where none of three local-damage hypotheses predicted which layers actually recover.

The int4 doorbell value-collision fix is load-bearing, not theoretical. Arm a75 logged 128 adjudicated collisions and completed; a25 and SP logged 0. Before the fix, a75 would have hard-failed exactly as arm S did.

ARM RSG, 2026-10-08 — RHT + SmoothQuant + GPTQ = 52.86, a NULL against R's 52.65. Third independent confirmation that RHT and GPTQ are substitutes: RHT is worth ÷1.32 when GPTQ is absent (Q 72.63 → S 55.11) and nothing at all when it is present (GPTQ alone 57.49 → RHT+GPTQ 58.59; R 52.65 → RSG 52.86, both inside the ~2% screen resolution). The pattern now holds across three different bases, so treat "add RHT on top of GPTQ" as closed rather than unmeasured. Zero doorbell value collisions on this pack, incidentally — the arm-S collision was pack-specific, as expected.

ARM S, 2026-10-08 — the transform stack BEATS GPTQ alone, with no calibration pass: 55.11 vs 57.49. RHT + SmoothQuant and no GPTQ at all scores 55.11 over the full 64 windows. That answers the standing question directly: the techniques GPTQ substitutes for do stack together and collectively beat it — while costing no Hessian, no Cholesky, and no O(M·K²) calibration. They are not pure substitutes either, since SmoothQuant + GPTQ (52.65) still beats both, so the honest reading is that GPTQ and the transform stack attack overlapping but not identical error, and the cheap half is worth more than it looked. Measuring this arm needed the int4 doorbell value-collision fix below; it had been failing at window 2 and its window-1 reading of 49.76 was, as usual, not the trend.

Rotation is not one lever among several, it is the lever — ÷11.3 against ÷1.3–1.5 for everything else.

Transform refinements and GPTQ are SUBSTITUTES. RHT is −2.6% without GPTQ and +1.9% with it; act-order is +3.6% with it. Both attack weight-side conditioning, and GPTQ's error feedback already does that. WUSH's own numbers show the same shape (+2.8 over Hadamard with RTN, +0.7 with GPTQ), so a data-aware transform on top of a GPTQ stack is a much weaker bet than it looks. The headroom is on the activation half, which only clipping touches.

Screen resolution: PPL is deterministic per pack, but this 16320-token screen cannot resolve <2% — RHT reads −1.9% at one rung and +1.2% at the next. Rank on the ≥20% effects only.

act-order is a measured negative (+3.6%), and the reason is structural: its benefit comes from interacting with group boundaries, and per-row scales have none — while the disruption to GPTQ's propagation order still costs. The property that made it free for us (no permutation to store) is the one that removes the upside. Don't re-try it while scales are per-row.

Promotion: the ranking signal beats the budget. Q4_K_M's own bump list (attn_v + ffn_down, every other layer) scores 49.39 on 28 MiB; our qerr ranking scores 53.95 on 34 MiB. qerr measures a weight's own rounding error, not how much it reaches the output — and its distribution is flat anyway (top decile = 12.8% of total error ≈ its proportional share), independently reproducing arXiv 2609.16391.

Also ruled out with reasons: NF4 is permanently impossible on native W4A4 (the only LUT is in SDP, after accumulation); grouped scales run on the NPU but cost ~24×; and W4A8 is 8 bits RESIDENT (Bi4 is a host-side store), so it does nothing for the IOVA-capacity tier though it remains right for smaller models.

Page: Exp-2026-10-07 — W4A4 vs GGUF 4-bit Harness: ork_ppl · wikitext-2 test (first 128 KiB) · 256-token windows, ubatch=256, KV cleared Knobs: ORK_QUANT=4 · ORK_GPTQ=1 · ORK_OFF=1 (CPU baseline) · ORK_OP_STATS=1 (claim + handler tallies) Kernel 6.1.115-vendor-rk35xx-fence, board .236, ork-driver 1.0.154, governors verified.

The 4-bit goal is parity with GGUF 4-bit, reached by stacking techniques — but that target had never been measured on the same harness as the W4A4 numbers, so "how far from parity" was not a quantity. It is now. All arms share corpus, window count, ubatch and binary (2 × 256-tok windows):

int8 already beats the GGUF 4-bit target. W4A4 sits at 3.7× it, and the W4A4/int8 ratio here (3.66×) is 2.6× worse than Exp-2026-08-23 recorded (1.40×).

GPTQ could not be measured: it registers ZERO weights on this model. Phase 1 hooks only native-W4A4 weights; ggml_backend_ork_gptq_max_k() returns 0 after the discovery decode, so ork_bench reports window 0 registered no weights — skipping and silently produces a plain RTN int4 pack. A GPTQ run is therefore indistinguishable from a non-GPTQ run — worth knowing before anyone attributes a PPL to it.

It is not model-specific. The obvious suspect was the model family (Gated DeltaNet, n_ff=0), and it was tested and refuted: qwen3-0.6b-f16, a plain dense transformer, registers zero weights too and also reports n_ff=0. Packing itself works and takes the int4 tier, so weights pack through a path that is not the native-W4A4 branch holding the phase-1 hook — likely the i4a8 tier rather than DT_I4. Moving the ladder to a plain transformer does not route around this; the non-GPTQ rungs are unaffected. The GDN model remains the leading suspect for the ratio anomaly only.

Two instrument errors are recorded on the page, both of which looked exactly like results: the rknpu counters are gated and default OFF after a reboot (zeros mean nothing until you write on), and a grep matching the ROUTE-STATS header but not its rows produced a phantom "claimed but never executed" routing bug. Also: build the pack as its own step — coupling build and measurement under one timeout yielded no number at all in 45 minutes.

Probes: tools/wr_diag.c (mechanism) · tools/wr_sweep.c (validated sweep) · tools/dr_time.c (correctness‖speed, self-voiding) · tools/wreuse_probe.c (timing — see the caveat) Knobs: ORK_MTILE_WR (bit 13) · ORK_MTILE_DR (bit 12) · ORK_WR_WBANK · ORK_WR_GRAINS · ORK_WR_GEOM / ORK_WR_DUMP · ORK_WR_NC (core count) · ORK_WR_FULLRAND (input distribution) Kernel 6.1.115-vendor-rk35xx-fence, board .236. This entry replaces several written mid-investigation whose conclusions later measurement overturned; the corrections are kept as method record.

reuse correct  ⟺  K × segw ≤ WEIGHT_BANK banks × 32 KB   AND   mc == mcap
correct_fraction = min(1, capacity / (K × segw))

The mc == mcap clause was found late (2026-09-06) and is not optional. 0x1040 is derived from mc, so a SHORT final M-tile gets a different bank split, its boundaries do not line up with the loader's, and it reads stale bytes even when the capacity half passes. Measured: K=1024 (mcap 320) with M=512 → tiles 320+192 → silent mismatch; with the clause it declines and is bit-exact. Every all-full-tile shape was already bit-exact (K=2048 mcap=128 M=512 → 128×4; K=3072 mcap=64 → 64×8), which is what isolated it.

Do not PREDICT the bank counts. 12 − mcap·K/32768 is wrong below K=2048 — K=1024 emits 0xa2 (WEIGHT_BANK=10, DATA_BANK=2), not the 2/10 that formula gives. Read WEIGHT_BANK back out of the emitted word. A first implementation of the law gate predicted them, had its readback guard conditioned on the forced-segmentation path only, and therefore shipped the K=1024 mismatch above with no protection at all.

where segw = N / nc is the columns one CORE owns, and capacity = WEIGHT_BANK × 32 KB = 128 KB.

Confirmed on a lever that is not shape. segw = N/nc, so forcing the core count (ORK_WR_NC) must move the corruption at a FIXED shape — and it does, to the same 128 KB constant. This matters because every shape-only correlation this investigation produced turned out to be worthless (see §2):

And confirmed against the input distribution, which is the check §2 shows is mandatory: widening the test data from wr_diag's 2-bit PRNG to the full int8 range moved the numbers toward the law (82.9 → 83.3%, exact), i.e. the narrow distribution was mildly flattering via value collisions but the law is data-independent. DATA_REUSE was not so lucky.

Independently corroborated by the vendor-equivalent open stack. Mesa's rocket driver (src/gallium/drivers/rocket/rkt_task.c) enables weight reuse on exactly this predicate: if (weights_banks_required + 1 < CBUF_BANKS) operation->reuse_weights_cbuf = true; — our capacity law, plus one spare bank, derived independently from the same silicon. Its calc_weights_banks() even carries banks++; /* Why do we need an extra bank? The calc above might be wrong on this HW */. It also documents a constraint we had not known: "Submit all tasks to the same core, so weights can be reused" — CBUF is per-core, which is precisely why nc belongs in the law.

Mechanism. The loader tile streams the whole K×segw weight past the CBUF weight banks, which retain only the last capacity bytes; the reuse tile reads the banks — correct for those columns, stale beyond. Corruption is deterministic (0/524288 differ across runs), confined to reuse tiles (loader 0.0% wrong), uniform across core segments, values are plausible dot products (no zeros — it is a real dot product, against the wrong bytes), and the correct columns form a contiguous run at the end measuring exactly 2 × 32-col groups = 128 KB = WEIGHT_BANK. Priming with a different matmul changes nothing.

This is the vendor's conv rule and it transfers. #39 observed a real out-of-envelope case, not a hardware limit.

Our pack layout is not at fault. The correct region being contiguous shows the column-major 32-wide grouping (woff = ((n/32)·KT + k/32)·1024 + …) lands coherently in the banks. Layout only becomes observable when the bit is set — without it every M-tile re-streams from DRAM, which is why make test passes.

The fix works. Widening WEIGHT_BANK 4 → 11 (0x1040 → 0xb1) turns a 49.8%-wrong shape into 0.0%. 0x1040 couples the split to the K-reduction schedule — mc × K = DATA_BANK × 32 KB — so DATA=1 forces mcap 128 → 16; both must move together or the activations overflow.

But it is worth ~2–4%, under the noise floor. M=144…256, 2 runs/cell: gains 3.4 / 2.5 / 5.1 / 2.2 / 2.0% against 6.7–7.0% within-config spread. A single-run "5.0%" at M=256 was 2.0% on repeat. Marginal cost of the extra programs is ~1.3 µs (the per-task hardware floor — chaining already amortises the submit), not the ~10 µs assumed. The real cap: weight streaming is only ~8% of runtime, and weight share falls as 1/M while tile count rises with M, so the lever and the regime never overlap. M=1 decode is the worst case, not the best — one M-tile, nothing to reuse.

BUILT AND MEASURED 2026-09-06 — it does not pay. See "N-tiling measured" below; the framing in this subsection was half wrong and the correction is there. Kept as written because the reasoning that led to the build is the part worth auditing.

…but that 2–4% measures the WRONG CONFIGURATION, and Mesa says so. The number above comes from widening WEIGHT_BANK 4 → 11 to make an overflowing shape fit. Because 0x1040 couples the split to the K-reduction schedule (mc × K = DATA_BANK × 32 KB), that collapses mcap 128 → 16 and quadruples the program count. We measured the cost of the bank widening and charged it to weight reuse.

Mesa never does this. Its rule is "full weights, partial input" — enable reuse only when the weights already fit (weights_banks_required + 1 < CBUF_BANKS), then split the INPUT across tasks; where they do not fit it sets reuse_weights_cbuf = false and takes available_input_banks = 7. So the shipping strategy on this silicon is the one design we identified and never tested:

keep the default 4/8 split and N-tile so segw ≤ capacity/K — same weight saving, ~4 programs instead of 16, no mcap collapse.

Blocked only by colsplit giving each core ONE segment (it subdivides only at N > NMAX). Status is therefore "untested in the regime that works", not "marginal". The ~8%-of-runtime weight-stream share still caps the upside for prefill-M matmuls, and M=1 decode still has one M-tile and nothing to reuse — but neither of those is what the 2–4% measured.

The gate was K == 4096 && Sn == 1 — a placeholder that "fails closed on a shape that happens to keep make test green". A whitelist can only be wrong about shapes nobody has run yet, so it is now the law evaluated on the geometry already being emitted: engage where the weight ALREADY fits, adding no tasks and narrowing nothing (unlike the forcing function below, which measured 0.39–0.42×).

tools/law_bench.c, both arms checked bit-exact against each other, reuse arm FIRST so the position bias works against it:

A real win, in a shape class nobody had tested — found by the law, not by intuition. It needs many FULL M-tiles: at M=256 (2 tiles) there is one reuse program per core and no measurable gain.

But eligibility is N/nc ≲ 64, i.e. N ≲ 192, so it declines on every production shape we hold: dense FFN 7168 (read from the LFM2.5 GGUF), attention 2048, MoE experts ~768. The only eligible tensor in a routed MoE is the router (N = expert_count, 32–128) — the smallest matmul in the model, a few hundred KB against ~100 MB of experts. So the change is inert today at zero cost. Its value is structural: a future model, draft head or adapter that lands in the eligible region engages automatically and correctly, instead of waiting for someone to widen a whitelist. make test ALL PASSED with it default-ON.

Built (ORK_WR_NTILE, segcap = WEIGHT_BANK×32KB/K columns, cap carried on the chain handle so begin_colsplit's emission and end()'s boundary-scatter cut identically) and measured (tools/ntile_bench.c, which checks the N-tiled checksum against the untiled one on every row):

make test also FAILS with the knob on — test_matmul {512,2048,256} reports run_i8 failed (the submit itself, plus 2 recovery-path ioctl failures), not a wrong answer. The completed rows are bit-exact, so the segmentation and the scatter agree; the code is right and the DESIGN is wrong.

The reading error. Mesa's weights_banks_required + 1 < CBUF_BANKS is a test of whether the weights already fit, not an instruction to make them fit. When they do not, rkt_split_tasks sets reuse_weights_cbuf = false, takes available_input_banks = 7, and moves on. It never tiles down to force a fit. A conv kernel (3×3×C×F) passes that test naturally with no narrowing; our matmul weight at K=N=2048 is 4 MB against a 384 KB CBUF, ~10× over, so forcing the fit demands segw = 64 columns — far below the output width the NPU computes efficiently. The per-task floor and the narrow-output inefficiency swamp the weight saving several times over.

So the correction to this entry's earlier claim. "Untested in the regime that works" was wrong: for LLM-shaped matmuls there is no such regime at the default bank split. Weight reuse requires the weight to fit CBUF already, and at K ≥ 2048 with any useful N it never does. That restores the original "marginal" verdict on a sounder footing than the retraction it replaced — not the mcap collapse, but the weight being an order of magnitude larger than the buffer the feature reuses. The lever is real, the capacity law is exact, and neither applies to our shapes.

It corrupts at every shape tested. There are no clean K. The earlier table below — eight shapes split into {2048, 2560, 4096} clean and {512, 1024, 1536, 3072, 3584} corrupt — was produced by wr_diag, whose PRNG was

static int rnd(void){ sd = sd*1103515245u + 12345u; return (int)((sd>>16)&3); }

Bit k of an LCG mod 2³² has period 2^(k+1), so bits 16–17 repeat every 262144 values. The activation bytes held by one M-tile is mcap × K:

At exactly those K, consecutive M-tiles held bit-identical activations, so reusing the previous tile's data was correct BY COINCIDENCE — no fetch had been declined. Confirmed by intervention: change ONLY the input distribution (ORK_WR_FULLRAND=1), leaving every register identical, and

(K=2560 is the one row the period account does not explain; it also submit-times-out under reuse, errno=110, and is excluded from timing.)

What this retires. The eight hypotheses below were each correctly refuted — but they were competing to explain a pattern that does not exist in the hardware. That is why no register field could separate the two sets, and why the WEIGHT_BANK power-of-two correlation held 8/8 across shapes and died instantly under intervention. The recommendation this entry previously carried — capture a vendor conv with bit 12 set — would have explained nothing. Retained as method record, not as open questions:

Hypothesis 1 was dismissed because failure looked binary (0% or ~99%, never partial). That inference does not hold, because the two operands corrupt asymmetrically through C[m,n] = Σ_k A[m,k]·B[k,n]:

• a stale weight column ruins only that column — columns are independent, so damage is proportional;
• a stale activation at any single k ruins the whole row for every n, because every output in the row sums over all k.

So partial data staleness CASCADES to ~100% of outputs and the error count saturates — it cannot distinguish "one bank stale" from "everything stale". The magnitude can, and measured mean |diff|/|ref| = 0.134 at K=3072: outputs ~13% off, not unrelated — i.e. a small fraction of the k-terms were stale. Residency was never refuted; the wrong statistic was being read.

ORK_WR_GEOM shows that within a core, wbase is constant across all four M-tiles while m0 varies. Consecutive tasks share the weight and differ in the activation. Our loop order (segment-outer, M-tile-inner) is therefore weight-reuse-correct and data-reuse-invalid by construction: setting bit 12 asks the hardware to compute tile N with tile N−1's rows. It is always wrong, and it only ever looked right when the two tiles happened to contain identical bytes.

The speedup is real and is real because the fetch is genuinely skipped — tools/dr_time.c, cold checksum, 4 warm iterations, self-voiding positive control:

DATA_REUSE is usable only under the transposed loop — M-tile-outer / segment-inner, so consecutive tasks share the activation and differ in weights. That is the decode-shaped schedule (one token's activation against many weight tiles), and it is the open follow-up.

Corroboration for "invalid by construction" from the one open-source production user of this silicon. src/gallium/drivers/rocket/rkt_regcmd.c has the only CBUF_CON0 emit site in the driver:

uint32_t con0 = CNA_CBUF_CON0_WEIGHT_BANK(task->weights_banks) |
                CNA_CBUF_CON0_DATA_BANK(task->input_banks);
if (task_num > 0 && operation->reuse_weights_cbuf)
   con0 |= CNA_CBUF_CON0_WEIGHT_REUSE(1);

CNA_CBUF_CON0_DATA_REUSE() is defined in rocket_registers.h and called from nowhere — 0 hits across rkt_ml.c, rkt_task.c, rkt_regcmd.c, rkt_coefs.c, rkt_device.c, intercept.c and every header. Only bit 13 is ever set, and only on task_num > 0 (our m0 > 0: skip the loader).

The reason is structural, not a Mesa limitation: rkt_split_tasks splits a conv by input height, so its consecutive tasks carry different input slices against the same weights — the identical loop shape to our segment-outer / M-tile-inner colsplit. Data reuse is invalid for that ordering in both stacks.

Built (ORK_WR_MOUTER=1 transpose + bit 12, =2 transpose only so the reorder and the bit stay separable; ORK_WR_SEGW forces a segment width, since without it a core owns too few segments to reuse across). The segments are cut once into a table and walked in either order — the scratch address is (segbase + m0*segw) and the activation is (adma + m0*K), neither of which depends on the walk order, so the SEGMENT-MAJOR layout end() expects is preserved and the copy-back is untouched. It is a pure reordering.

The bit that corrupts 100% of outputs in the default order is bit-exact in the transposed one, at every shape and every M (tools/mouter_bench.c, segw=64, three arms against a common reference):

That is the direct confirmation of "invalid by construction": DATA_REUSE was never broken silicon and never shape-dependent — it was the wrong bit for a segment-outer loop. Against the reordered baseline the bit does real work: 2.00× at M=256, 1.60× at M=64, removing the per-segment activation re-fetch (≈5.6 MB at M=256, K=2048).

It still does not pay, for three separate reasons.

1. Decode — the whole premise — gains exactly nothing. M=1: 608 µs with the bit, 608 without. The activation is K bytes (2 KB at K=2048); skipping its re-fetch is invisible against a K×N weight stream. Decode is weight-bound, not activation-bound. Verified that reuse actually FIRED there rather than inferring it: ORK_WR_SEGW=64 at M=1 emits 99 programs against 9 unforced (3 runs × 3 cores × 11 segments, segw=64), so there were 10 reuse opportunities per core and they bought nothing. The first attempt at this check read 9 both ways and looked like the knob was dead — that was a STALE wr_diag (make wr_diag relinked nothing; strings showed zero occurrences of the knob while mouter_bench, built the same minute, had it).
2. Creating segments to reuse across costs more than the reuse recovers — 0.80–0.85× of base, the same narrow-segment penalty that killed the N-tiling above.
3. Natural segmentation is far too coarse to help. nmax = 8192, so on the widest real projection (K=3584, N=18944) the three cores get 1, 2 and 2 segments — at most ONE reuse program per M-tile, on two of three cores. Measured 1.01–1.035× on those shapes, but the control shows the same 1.028% while holding a single segment per core, where reuse cannot fire at all. That band is arm-ordering drift, not the bit. (Segment counts read off ORK_WR_GEOM, not inferred.)

Asked directly: is there an operand size big enough that reuse pays? The structure says no, and it is not about size. Per core the unsegmented baseline already fetches the activation once per M-tile (M×K), which is the floor. Segmenting into nsg pieces inflates that to nsg × M×K; DATA_REUSE returns it to M×K. Reuse's best possible outcome is to EQUAL the baseline — it repairs damage segmentation caused rather than creating a saving.

Searched anyway, across 14 shapes chosen to favour it (small K ⇒ large mcap, wide M, narrow segw, since the lever's strength is the ratio mcap/segw):

Zero wins. Every row is below 1.000 and they approach it ASYMPTOTICALLY FROM BELOW as segw widens toward the unsegmented width — the predicted curve, whose limit is the baseline, reached only when there is nothing left to reuse across. Reuse is doing real work (it recovers 57% / 74% / 82% of the segmentation cost at K=512 / 1024 / 2048, improving with K as activations grow) but repair caps at "undamaged".

A bigger model makes it worse, not better. mcap shrinks as K grows (mc×K = DATA_BANK×32KB: 704 at K=512, 128 at K=2048, 64 at K=4096), so large-K models get the smallest M-tiles and the weakest activation-to-weight ratio. The only free segmentation is nmax-forced (1/2/2 segments on wide-N, ~0.5%), and nsg there is set by N/nmax, which model size does not move.

Unmeasured: K=2048/M=512 at segw=128 and 224 failed to run, so the largest-K asymptote is missing. Out of scope: attention BMM inverts the premise — both operands are activations and the "weight" is not resident — but that is a different code path from ork_i8_mm_run and untested here.

Sweeping the forced segment width at M=1, K=2048, N=2048 (nc=3, one M-tile, so tasks/core = segments/core) turns the task count into an independent variable:

~9–13 µs per extra task, and that UNDER-states it (each base is arm 0 and carries the position bias above). This does not contradict the ~1.3 µs marginal figure recorded elsewhere in this log — it bounds its regime. At large M a task's weight DMA overlaps its neighbour's compute and the fixed cost hides; at M=1 the task is compute-trivial, nothing hides it, and it dominates. That single number explains every negative result in this entry: N-tiling, forced segmentation and the transposed loop all pay ~10 µs per task they add, and no reuse of a ≤512 KB operand can repay that.

Reuse off vs on at MATCHED task count is the comparison that matters, and it is null: Δ = −16.0, +1.7, +6.1, +23.8 µs — sign flips, inside a ±25 µs floor (the base column wanders 478→527 for identical work). The harness's own null is the calibration: at segw=672 reuse CANNOT fire (one segment) and the arms land 0.7 µs apart, so a real null reads ~0 and these rows do not.

An earlier version of this section reported "608 vs 608" at M=1 and called it suspicious consistency. It was %.0f rounding over 4 iterations; at 20 iterations the same pair is 570.7 vs 576.8. The bit was genuinely firing throughout — 90 readbacks with b12=1, exactly 10 per core per run.

Any arm-ordered A/B inherits a position bias, and it needed sizing rather than assuming. tools/aa_probe.c runs the identical configuration in all three arms — no knob is ever set:

Later arms run 1.01–1.126× "faster" doing identical work: the first arm absorbs mode entry, the first weight DMA and page/cache warm, and the rest inherit it. This is not a lever — production has no first arm, so it is already on the warm side after the first call; there is nothing to implement. What it does buy is a scale for reading everything else: the 1.01–1.035× on wide-N is INSIDE the bias band and is therefore null (confirmed by control, not by arithmetic), while the 2.00× reuse-vs-reorder is arm1-vs-arm2 where the bias is ~1.013×, and the 0.80–0.85× penalty is one the bias can only understate.

Verdict: the mechanism is now fully explained and provably correct, and no shape in this workload earns its cost. Both reuse bits end in the same place for the same underlying reason — CBUF is ~384 KB and our operands are not sized for it: the weight is ~10× too big to stay resident, and the activation is far too small for its re-fetch to matter.

• drivers/accel/rocket/rocket_registers.h (mainline) confirms our bit map exactly: WEIGHT_REUSE 13, DATA_REUSE 12, FC_DATA_BANK [10:8], WEIGHT_BANK [7:4], DATA_BANK [3:0]. CBUF_CON1 is DATA_ENTRIES[13:0]; FC_CON1/FC_CON2 are DATA_OFFSET/WEIGHT_OFFSET [16:0]; FC_CON0 is FC_SKIP_DATA[31:16] + FC_SKIP_EN[0].
• NVDLA programming guide documents the semantics: "WEIGHT/DATA_REUSE: re-use the weight/data from previous hardware-layer. If this flag is set, CDMA fetch will be fully(partially) skipped", gated on CDMA_HEIGHT_N <= (CSC_HEIGHT − D_RELEASE)_{N−1}, with a companion WEIGHT/DATA_SKIP_RELEASE that keeps the producing layer's slices from being released. We never program the release side.
• NVDLA CBUF is a circular buffer: "Each buffer acts as circular buffers… if the address reaches the max, it wraps to zero" — capacity and wrap phase both matter, which is the shape of the weight-side law.

Ten wrong conclusions were reported during this work. Every one came from the same shape: a favourable result on one configuration, explained by a plausible mechanism, stated before a validated detector could disagree.

1. An unvalidated detector. wreuse_probe reported bit-exact for shapes that were corrupt — its timing loop re-ran the identical submit 20× so the checked output came from a warm iteration with the weight still resident; test_matmul runs each shape once. Four rounds of findings rested on it and are void (a "narrow envelope K=4096", a K-sweep, an N-scaling curve, a vendor-capacity test). → A probe never shown to detect a known failure carries no information when it reports success. wr_sweep therefore embeds its own positive control.
2. Stale binaries, four times. ccache not rebuilding after a knob moved; three source files never synced to the board; make wreuse_probe rebuilding only the probe while test_matmul kept old objects; wr_sweep unrebuilt so a "DATA_REUSE sweep" actually measured WEIGHT_REUSE. → Rebuild the specific binary and prove the knob fired (count applications, read the register back).
3. Lost source state. colsplit.c silently reverted mid-session to a pre-commit version; the ORK_WR_BANKS knob vanished, and a "refutation" was published from runs where the knob did not exist. → Check git status against HEAD before trusting a knobbed result.
4. Closing on a black box. "It corrupts" was recorded as final. Four questions — deterministic? which rows? which columns? what do the values look like? — turned it into an exact law in one run.
5. A shell bug inverting an A/B. env VAR= … sets an empty string, which getenv reports as set, so a run labelled "base first" was actually reuse-first.
6. Correlation read as causation. A CI failure was attributed to this work and "fixed" with a commit message saying so; it had been red for a day from an unrelated commit (393f28b, two if(a) x; if(b) y; lines in prof.c).
7. A noise excursion reported as a result. 5.0% at M=256 became 2.0% on repeat.
8. Instrumentation placed where it could not observe. The regcmd dump ran before the reuse bit was applied, so a raw word diff showed zero differences between the two bits.
9. The gate mistaken for the feature. dr_time was first run with ORK_MTILE_WR_ALL=1 alone — which only lifts the SHAPE gate; the feature is ORK_MTILE_WR. So wrm=0, both arms emitted the identical baseline program, and it printed six CLEAN rows and a 1.03× "speedup" that was warm-up drift between arm 1 and arm 2. Caught only because three of those rows are known-corrupt shapes and cannot be clean. → This is why the probe now carries its control INSIDE it and returns VOID instead of a table.
10. The test data itself was the variable — the largest error here, and the one no amount of register discipline would have caught. Eight shapes were classified clean/corrupt, eight hypotheses were raised and refuted against that classification, and a wiki entry recommended a vendor capture to explain it. The classification was an artifact of a 2-bit LCG whose period (262144) exactly matched mcap × K at the two "clean" K. Every register was correct; the inputs were degenerate. → When a pattern resists every explanation, suspect the harness before the silicon. Vary the one thing a hardware account says must be irrelevant — here, the input distribution — and see if it moves.

Full page: Exp-2026-09-05 Heterogeneous Dispatch Ceiling

Retraction first: the "37.4 GB/s aggregate, ~1.5×" figure from earlier today was wrong. It added the NPU's useful weight bytes to the CPU's read+write traffic — two different currencies — and the sum exceeded the board's theoretical LPDDR peak (33.8 GB/s), which is proof it over-counted.

Re-measured in one currency (bytes read) and checked against the DDR controller's own load counter. The check passes independently in every configuration — 27.40/0.81, 30.17/0.90, 20.88/0.63, 22.66/0.70 all imply the same ~33 GB/s at 100% load:

Two engines: CPU+NPU reach 29.30 GB/s at 88% — a 1.41× gain in the decode-shaped regime but only 1.10× when the CPU already streams well. The gain is entirely a function of how badly each engine saturates the bus alone, and decode is latency-bound, so it leaves 40% idle for the NPU to fill.

Three engines: running Mali concurrently, CPU+NPU fell 29.30 → 16.92 GB/s while DMC went 88% → 85%. Total bus traffic did not rise; Mali took ~12 GB/s straight out of the other two. The ceiling had already shown itself in the two-engine data — CPU2+NPU 30.08 vs CPU3+NPU 30.12, i.e. a whole extra core of demand bought 0.04 GB/s and only changed the split.

A third MEMORY-bound engine cannot beat two on bandwidth. For decode this sets a hard wall of ~30 GB/s ÷ bytes per token: a 1.7B Q8_0 tops out near 17 tok/s (measured 12.8 = 73% of the wall, little left), while a 7B Q8_0 tops out near 4.3 (measured 0.89 = 20%, lots left).

But a COMPUTE-bound third engine is genuinely additive — the homogeneous test above is the special case, not the rule. Mali running a memory-free kernel (verified DMC 0%) beside the CPU+NPU pair already holding the bus at 88%:

Both keep essentially everything, versus the same GPU doing memory work cutting CPU+NPU to 58%. And additivity scales with how memory-free the compute work actually is: the NPU made compute-bound via M=256 (97.9 GMAC/s at DMC 16%) retains 96% against one CPU core but only 61% against three, because at M=256 it still moves A and C. Engines are additive when they contend for different resources and zero-sum when they contend for the same one — dispatch on resource profile, not engine count.

Full page: Exp-2026-09-05 Fixed Per-Task Cost

The SRAM investigation left a bigger lever behind than the one it closed. Decomposing a chained int8 matmul as time = a + b·N, a CPU DRAM load inflates the fixed part 2.23× while the weight stream moves only 1.22×. The kernel's #patch74 hardware-time counters say where that lands:

113% of the growth is host-side; the NPU's own fixed per-task cost is ~1.27 µs and does not inflate at all. Which explains, after the fact, why SRAM placement couldn't help — the NPU-side fixed cost was never the problem, so on-chip placement had nothing to protect.

Removing it works. Host cost (wall − kernel hardware time), single-core both arms:

80% removed quiet, 72% under load — an absolute saving 1.82× larger under contention, worth ~1.30× at decode-sized shapes and tapering to 1.10× at 640 KiB as the fixed cost amortises. So every quiet-system benchmark of ork_pc_* understates it, and it understates it most in exactly the CPU∥NPU partition case. The spread against the recorded 6.03× (regime A) and 1.47× (FFN layer) is not inconsistency — it is one fixed cost divided by different amounts of work per task.

Method note worth keeping: SPIN (same cores busy, zero memory traffic) rather than an idle baseline. The NPU submit path is host-CPU-bound, so an idle baseline mixes scheduling into the contention number.

Full page: Exp-2026-09-05 SRAM Aggregate Bandwidth

The standing partition thesis is that CPU∥NPU decode only pays if the NPU's share reads a different memory — otherwise both engines just divide one ~28 GB/s DRAM budget. syssram was the only candidate. Five probes across four weeks, and the answer is directional, which is why it read as a contradiction:

So SRAM traffic does not spend DRAM's budget when the CPU issues it, and does when the NPU does. Every previous unqualified "SRAM is/isn't a separate port" has been wrong in one direction or the other.

The methodological content is the point. The 2026-09-05 morning probe (sram_bw_probe) reported flat and was void: a 128 KiB weight and one submit per iteration meant ~93% of each iteration was submit overhead, so weight placement governed ~7% of the runtime. That flatness was then explained by inventing a mechanism ("same LPDDR controller") that contradicted a measurement already on disk. The replacement probe fixes it structurally: it chains 16 tasks over a 640 KiB weight, fits time = a + b·N and refuses to print a verdict below a 70% streaming fraction, and adds a SPIN contender (same cores, zero memory traffic) so the scheduling term is divided out rather than assumed away — SPIN retention came back 99.8–100.1%, proving the affinity split leaves no scheduling artifact.

With that in place the null replicates at three weight sizes with the sign flipping (+0.3 / −0.5 / +0.3 pp), and the residual contention localises: fitting retention(N) gives the weight stream slowing only 1.22× while the fixed, DRAM-resident per-task remainder (A/C/descriptors) slows 2.19× — a fit that predicts the unseen midpoint to 0.1 pp. Most of the NPU's loss under CPU load was never the weight stream, and that remainder cannot move to SRAM (task descriptors in SRAM is a hard wedge).

Rules earned here: never explain a null before proving the experiment could have shown a positive; a contention probe must verify its own contender (at -O3 the saturator's memcpy was deleted outright and "reached" 7,056,858 GB/s while generating no traffic); and a cross-project finding needs an anchor in both project indexes — this record lived in the llama.cpp session history, which is why it was re-derived.

Full page: NPU on-chip SRAM

The 956 KB syssram block (DTB Change 1) has been a standing "next lever" since 2026-06-26. It is now closed, and the closing is more instructive than the verdict — this page existed but was never linked from this log, which is why the same ground was re-covered three times in one session.

Placement is not the doubt — the driver prints 956 KB of the regcmd landed in on-chip SRAM, and a weight alloc drops the pool by exactly its size. It works. It just buys nothing.

Do not read the flat weight row as "SRAM is not a separate port." It is one — measured, 2026-08-11. A 2026-09-05 revision of the topic page asserted the opposite ("same LPDDR controller"), invented to explain the flat row; it was corrected once the original probes (test_sram_bw, test_sram_npubw, test_sram_pipe) were found. That gap is now closed — see the entry below: NPU-read ‖ CPU-write was re-measured properly and is a null, so "separate port" is directional (true for CPU access to syssram, false for NPU weight reads).

Three wrong explanations, each killed by the next measurement: (1) bandwidth (1.2–3 GB/s vs DRAM's measured 10.5 — and note the 1.2 is the uncached mapping, where SRAM and DRAM are identical, not a property of SRAM) — true but irrelevant, a descriptor ring is not bandwidth-bound and syssram is designed for descriptor rings; (2) "read once per submit" — false, ork_dyn_halt rewrites a not-yet-executed program's 0x0014 after submit and the sequencer honours it, proving reads are per-task at walk time; (3) "ping-pong hides the fetch, so it should pay with ping-pong off" — tested and falsified, a 2×2 on a real 6-program single-core chain is flat in all four cells, and ping-pong itself is flat too, so the fetch was never on the critical path in either mode. What survives is magnitude: ~1 KB against 115 µs.

Two false-positive "wins", both killed by replication (2026-08-26 and 2026-09-04). The second one's tell: the single-shot delta included kernel hardware time improving 15%, which regcmd placement cannot plausibly cause. A physically implausible component of a delta is a noise signature, not a mechanism. These probe timings carry ~10% run-to-run spread — a single A/B here is worthless; use ≥3 trials or A-B-A.

What SRAM is actually good for here: memory that costs no DRAM/IOVA (the one deliberate use in-tree — ork_pc_compile parks its regcmd pool there), and low latency for small always-available objects. We have no latency bottleneck: the M=1 per-call cost is CPU work (synth, bsync, descriptors). Verdict: ❌ do not spend more on it without a mechanism predicting a win in latency of a small hot object.

Full page: Exp-2026-09-03-KV-Cache-Quantization

Hypothesis was that attention is memory-bound (2.0 FLOP/byte before reuse vs ~15 machine balance), so halving KV bytes should roughly halve the attention term. Refuted — it went 4.2× the other way. The cost is dequant inside the flash-attention inner loop; FA tiles, so the no-reuse traffic bound was never approached and there was no bandwidth wall to hit.

llama-bench -p 400,2000 -n 0 -r 3 -t 4 -b 512 -ub 512 -ngl 99 -dev ORK -fa on, int4 .orkpack, GGML_DISABLE_VULKAN=1. Slope = per-token cost of one extra token of context.

Two penalties stack: quantising at all (3.6×), plus mismatching the pair (~1.7× more — both mixed pairs are worse than both matched pairs). At 4000 tokens f16→q8_0 is −68%. Not the GPU: the q8_0 arm never touches Vulkan and pays the same as TurboQuant on the Mali.

TurboQuant (measured through oRKLLM — turbo needs ggml_vk_set_mode, llama-bench can't reach it; types verified live via [CPAR-DBG]): (q8_0,turbo3) 91.60/41.10/24.48 vs matched (turbo3,turbo3) 69.57/40.80/24.96 at 402/1975/3943. The mixed-pair penalty does not transfer to turbo — turbo-on-K is 24% worse short, a wash long. oRKLLM's never lead with turbo K policy is right on speed as well as quality. (A 2-point fit projected an 11% matched win at 4000; measured 2%. Projection recorded as wrong.)

Side result — llama.cpp-rockchip#6 closed as a false lead. The reported ~2.6× "oRKLLM worker vs bare process" gap was this setting: the served path ran (q8_0,turbo3) while every probe ran f16/f16, and the exclusion table varied n_ctx but never the KV dtypes. With KV matched, served beats the CLI (244.66/136.30/85.59 vs 227.07/130.19/84.56) and the slopes are identical. There is no serving overhead. Read [CPAR-DBG] before comparing two processes.

Correction carried on the full page: the ~365 µs "submit floor" cited in ggml-ork's supports_op (and Exp-2026-06-12-Submit-Floor-Profiling) is the blocking mcworker path. The chained/doorbell figures were already on record in this wiki: Exp-2026-06-21-Mixture-of-Experts-PC-Chaining-and-Scheduler-Programming measured 58.1 → 5.2 µs/matmul (11.21×) going from 8 separate submits to 1 chained submit with "zero CPU intervention", and Exp-2026-07-16-Dynamic-Submit-API-and-Doorbell puts NONBLOCK per-op progress overhead at "a few µs/op" (its ~167 µs is the NONBLOCK submit, not a per-op cost). So ORK_MINM and the M=1-decode-is-CPU policy are both calibrated against a cost the doorbell does not pay — worth re-deriving.

Open: decode unmeasured (bandwidth-bound, could be a turbo win); ORK_ATTN unmeasured and is the largest remaining prefill lever (attention is ~70% of prefill at 4000 tokens even on f16, ~56 GFLOP/s).

Filed: llama.cpp-rockchip#7, oRKLLM#1.

Kernels: #56 (armbian base + only the five upstream fixes) vs #57 (the full redesign). Same board, same workload, same session. Outcome: #388 (5 fixes only) was closed and replaced by #390, which carries those five plus the rework and closes #387.

The redesign (Kernel-Modifications changes 18–23) had never been tested against the alternative: keep the vendor architecture and apply only the five reference-counting fixes. Both arms were built on the armbian tree — the rknpu driver is 0–2 differing lines from vendor develop-6.1, so "vendor architecture" is what the armbian tree already contains, and this avoids booting an untested vendor-tree kernel on a board whose recovery needs SSH.

The probes pass on both arms — only the real workload discriminates. Worth remembering: a probe that reproduces a defect is not the same as a probe that can clear one.

Why Arm A dies (persistent NVMe log, previous boot):

rk_iommu fdab9000.iommu: Error during raw reset. MMU_DTE_ADDR is not functioning
iommu driver failed to attach a compatible domain
WARNING: CPU: 5 PID: 7555 at drivers/iommu/iommu.c:2116 __iommu_attach_group+0x9c/0xac
Comm: ork_bench  Hardware name: Radxa ROCK 5B (DT)

The vendor switch is iommu_detach_device() + iommu_attach_device(), and rk_iommu_attach_device() → rk_iommu_enable() → rk_iommu_force_reset(). Under concurrent load the force-reset fails its DTE_ADDR readback, the attach then fails, and the core leaves the device attached to nothing. This is the mechanism the very first diagnostic patch (#patch30) suspected and that the light switch removes by never force-resetting on a switch.

So the redesign's purpose is now empirical, not theoretical: it exists to take rk_iommu_force_reset() out of the domain-switch path. The five fixes are necessary but not sufficient — they remove the refcount wedge (swtmo 12 → 0, no more mismatch domain), and the force-reset failure is what remains.

Consequence for packaging: the five fixes live entirely in drivers/rknpu/, which builds as a module, so they can ship as a DKMS package with no kernel replacement. The redesign touches drivers/iommu/rockchip-iommu.c, which is built-in — it requires shipping a kernel. That is a real cost, and this A/B is what justifies paying it.

Kernel: 6.1.115-vendor-rk35xx-fence builds #45 → #52 (Kernel-Modifications changes 18–21). ork-driver: 657193c 9c506a9 05571c5 32c9158 5e2b7da cad6f2b 10a9324 2aebfd1. Knobs: ORK_PIN_DOM ORK_DOM_BASE ORK_DOM_WASH ORK_DOM_RAMP ORK_WARM_REPS ORK_SUBTRACE · kernel params: dom_reclaim switch_reset dbg_healthy_ctrl dbg_false_ok. Probes: tools/dom_scale_probe.c (new), orkd_dom_api.

The multi-domain path went from wedging the board to failing safely and loudly — and two classes of silent wrong output were found and fixed on the way. The underlying dispatch stall is untouched and has now survived fifteen distinct interventions.

What shipped.

1. Ch.18 — light domain switch. dma-iommu resolves every mapping through iommu_group->default_domain, so the DMA API is single-domain by construction; the vendor made domain N reachable by overwriting the core's default via a pasted-in copy of struct iommu_group. Replaced with rk_iommu_switch_domain() = stall → write DTE_ADDR → ZAP_CACHE → unstall, no force_reset, no core involvement, plus a driver-owned IOVA allocator per domain. mismatch domain → 0, switch timeouts 12 → 0, no perf cost (217.0 → 217.2 tok/s). ZAP alone suffices — force_reset was not doing necessary work (verified via switch_reset=0/1).
2. Ch.19 — reclaim-and-retry. A stalled job pinned the domain refcount forever, so every later switch burned 6 s and every allocation failed -EINVAL. patch54 already reaped but then returned the error anyway. Now reaps on the first timeout and retries. Unkillable D-state → clean exit.
3. Ch.20 — soft reset must reprogram the LIVE domain. ★ A regression from ch.18: rknpu_soft_reset() re-established the MMU via iommu_get_domain_for_dev(), which after ch.18 returns domain 0 — so a reset silently re-attached the wrong page table and the next job was committed and never completed. This was the "first int8 op after an ACT_RESET is silently dropped" bug. Fixed with rk_iommu_reprogram().
4. Ch.21 — audit of every remaining iommu_get_domain_for_dev() caller. Four more: the per-core reset (the twin ch.20 missed — the path the watchdog drives), the post-map IOTLB flush, and the SRAM/NBUF cache iommu_map/iommu_unmap. All targeted domain 0.
5. ork-driver: two silent-success classes. orki_submit1_db() polled a sentinel and returned 0 on timeout, so a dropped op was indistinguishable from a completed one. Auditing all six poll sites found a worse one on a default path: the fp16 K-split verify in ork_dyn_end() — written to stop the accumulate reading a mid-drain partial — summed the partials anyway on timeout, doing the exact thing it guards against.

The reproducer (replaces multi-minute bench runs): sudo -E ORK_MM_TIMEOUT=3000 ./orkd_dom_api failed 5/5 in ~3 min with no model load; ORK_PIN_DOM=0 passed. The short timeout is what makes it deterministic — at the 60 s default the same fault merely looks "intermittent" because each occurrence costs a minute.

Ruled out by measurement (each A/B'd, most repeated): domain count (the failing bench uses two, same as the passing probe — capping is a no-op) · domain 0's involvement (ORK_DOM_BASE=1 runs in {1,2,3} and still fails) · switch/submit race (kernel logs job-vs-device domain: 115 of 121 agree) · cycle count (2000 switch+submit cycles, bit-exact) · resident bytes (600 cycles at 1024 MiB across 2 domains, bit-exact — the bench's own profile) · warmup depth (2/3/4) · ping-pong (0x5 vs 0x1) · sacrificial ops at every position, count and core mask (ORK_DOM_WASH, ORK_DOM_RAMP, post-ACT_RESET drain) · drain on/off · submit timeout 3 s→60 s (identical, wall time scaling with the timeout) · switch_reset full-vs-light · task_ctrl readback · status bit 0x2000 · deferred GEM reclaim (queue objects whose destroy bailed, drain them when the domain becomes reachable — 4 runs, deferred=20 reclaimed=0 every time: reclaim needs a switch into the very domain that is unreachable, so it can never fire in the failure it was written for; kept for the visibility, it recovers nothing).

★★★ ROOT CAUSE (same day, d1cb00a): dma-buf IMPORTS are mapped into domain 0. The bench named it in its own failure line — weight[... imported=1]. Giving dom_scale_probe an imported mode (pack → ork_w_dump → ork_i8_mm_load_import) makes the weight SOURCE the only variable:

It fails on the FIRST cycle, so scale was never involved — the bench only looked scale-dependent because every one of its weights is imported. Mechanism: rknpu_gem_prime_import_sg_table() takes its IOVA from sg_dma_address(sgt->sgl), and that sg was mapped by dma_buf_map_attachment() → the DMA API → iommu_get_dma_domain() → default_domain = domain 0 — while the object is recorded as domain N. A submit in domain N then resolves an IOVA that exists only in domain 0's page table. This is the THIRD consequence of ch.18, and the one a iommu_get_domain_for_dev grep structurally cannot find: the DMA API does it implicitly. Predicted-then-confirmed with zero code: imported + ORK_PIN_DOM=0 (live == default) PASSES; imported + multi-domain FAILS.

★★★ FIXED same day (kernel ch.22, build #54; ork-driver 4f8db7f). Two halves, because the sg is mapped in two places that both got the domain wrong. (a) dma_buf_map_attachment() runs inside DRM core's PRIME_FD_TO_HANDLE, before any rknpu hook, and maps through the DMA API → domain 0. The .gem_prime_import wrapper is the interception point: scope an override of the core's default domain to that one call (and to the matching unmap at destroy), under domain_lock, restored immediately — not the vendor's permanent overwrite, whose exposure was that any core path reattaching "the default" used a domain the driver might free. (b) PRIME_FD_TO_HANDLE carries no domain at all, and the later MEM_CREATE does not re-map an existing handle — so a new switch-only action RKNPU_ACT_SET_DOMAIN is issued before the PRIME ioctl in both import paths, putting the mapping where the weight will be used.

The multi-domain bench, failing since the start of this investigation, now passes:

Multi-domain now matches single-domain performance exactly, output coherent, make test green with 0 sentinel misses. The "dispatch stall" was never silicon — it was a page-table mismatch, and the register archaeology that chased it (task counter, task_ctrl, status bits, interrupt masks) was chasing a symptom of software the whole time.

Seventh lesson, the expensive one: vm_kbuild.sh installed the PREVIOUS image after a failed build, so the board booted a stale kernel while the log said "installed" — a compile error in this very fix was briefly scored as "the fix does not work". It now aborts on a failed build. Never let a build failure fall through to a deploy step; a stale binary reads as a clean negative.

Mainline cross-reference (accel/rocket). rocket_core_reset() does no IOMMU work at all — only reset-line assert/deassert; rknpu entangles an IOMMU detach/attach with the hardware reset, and that extra step is exactly what broke in ch.20. Rocket also has no stall detection: no task-counter monitoring, no missing-interrupt logic, just a 500 ms drm_sched timeout → core reset. So mainline offers no evidence that this stall is known silicon behaviour.

Reusable lessons. (1) A matched control kills most "discoveries". Four single-sample leads died this way (task_ctrl readback, CNA/DPU S_STATUS, PC_OP_EN, INT_RAW_STATUS). The task_ctrl one looked decisive — write 0x7001, read back 0x1001, trigger bits gone — until a healthy job read back 0x1002 from 0x7002. It is self-clearing. Sample the control at the same signature; dbg_healthy_ctrl exists because a hard-coded 0x7002 gate logged nothing against a 0x7001 stall. (2) Verify the intervention actually ran. ORK_DOM_RAMP=3 "had no effect" — and the kernel commit counter (15 → 42, +27 = 3 ops × 3 cores × 3 switches) is what makes that a negative rather than a no-op. (3) One trial per boot, until failures stop being destructive. An early switch_escalate A/B was uninterpretable because each failure left the next run crashing early (rc=139, 0 stalls) — data that looks like a result. Ch.19 is what made same-boot A/B valid again. (4) A dmesg NO TASK PROGRESS line is a 120 ms watchdog DETECTION, not a failure. The residency runs show 1–2 with every output bit-exact. (5) Guard everything, including make test. The board hard-faulted and auto-rebooted when a make test was launched while an A/B still held the NPU; mt3.sh did not go through npu_guard.sh. The NVMe log caught the signature (job dom 2 / dev dom 0, domain_put UNDERFLOW). (6) Silent success is worse than a wedge. Two independent poll sites returned 0 on timeout. For a model run that is garbage output instead of an error; both now fail loudly.

Probe: tools/re/i4_widek_stall_probe.c (new). Kernel: 6.1.115-vendor-rk35xx-fence (changes 4+5, CONFIG_ROCKCHIP_RKNPU_FENCE=y). Full writeup: NPU-Quirks · Kernel-Modifications.

The doorbell's long-standing "~1/4000 dispatch drop" costs ~23% of 27B W4A4 prefill (time = 23.01 s + 3.61 s × misses, R² 0.967). A kernel watchdog sampling the PC completed-task counter at 1 kHz settled it: on every stall the counter reads 0x0 with the job >100 ms old, on all three cores — the submit returns rc=0 and the hardware never executes a single task. Not a stall, not a partial write.

Reusable lessons. (1) The contiguous-tail sentinel pattern is not a partial write — cores get contiguous op ranges, so a core that never runs leaves its slice at the sentinel. (2) At a ~2% rate, n=400 gives ~10 events (Poisson ±32%) — direction only; use ≥4000 and always run the positive control in the same pass. Two wrong conclusions here came from n=60 runs. (3) A probe can faithfully reproduce a defect and still over-weight it relative to production — check against a model run before believing a transfer. (4) Reading NPU MMIO from a kernel timer without holding power_lock hard-hangs the board; pm_runtime_get_if_in_use() on the parent is not sufficient on RK3588 (per-core genpds).

Ruled out by measurement: miss-window width (flat 300 ms → 30 s), read visibility (a kernel FROM_DEVICE re-count recovers nothing), the missing dsb between dc civac and the load (a real ARM bug, fixed, but zero effect on the rate), kernel-side kill (no dmesg fault), int16 sentinel collision, forbidden shape, and a post-switch dummy op.

Commits: ork-driver main 3b1eb96 (int8+fp16 SDP), 6033334 + bfefc64 (int4). Probes: tools/re/sdp_mcap_probe.c, tools/re/i4_hcap_probe.c. Full writeup: Exp 2026-08-20 SDP and Int4 Envelopes.

Finishes the sweep started on the fp16 matmul path, so every precision now runs on a measured bound.

The SDP 8192 was a register FIELD WIDTH (RK_DPU_DATA_CUBE_WIDTH mask 0x1fff) read as a capability, copy-pasted to 21 sites. A field width bounds what you can ENCODE, never what the hardware COMPLETES — and here the guard was admitting a shape that faults the NPU rather than refusing it. fp16 SDP frays rather than cliffs (8188 OK, 8189 BAD, 8190 OK, 8191 BAD), so the usable bound is the CONTIGUOUS max.

int4 was a throughput bound, not a correctness one — BCHAIN already tiles M. But if(H>16)H=16 was invented (2× at K=512, 4× at K=256), and 16384/K needed to be CEIL: plain integer division is right only when K divides 16384, losing a row at K=768 (22 not 21) and K=1536 (11 not 10) — K=768 being a real down-proj shape. Six K bracketed at integer resolution. Checksums-across-processes was the trick: H only changes M-tiling, so a valid H must reproduce a known-good H bit-for-bit, needing no model of int4 packing.

★ The recurring error is LADDER RESOLUTION. The int4 sweep first used only powers of two — every K probed happens to divide 16384, so floor and ceil coincided and the data looked conclusive. Same error hid the fp16 352 ceiling behind a 320/384 probe. A ladder proves nothing between its rungs.

★ Derived probe list (in the writeup): Wb=(131072/K)&~63 one line below the H fix is the same class and UNPROBED — it discards 25% of the N-tile at K=768/1536; the int4 H>64 clamp is a boundary this campaign introduced; the fp16 CBUF>32768 clamp discards 43% of cbuf_elems. Also recorded as already-fine so nobody re-probes: fp16 sched=0 floor (bracketed at 4 non-dividing K), fp16 sched=1 (moot — all legal K divide evenly), int8 mg_max floor (bracketed at K=1536). And a live inconsistency: fp16 FLOORS while int4 CEILS on the same nominal 16384 budget, both measured at K=768 — so the "16384" is not the same quantity in both. Do not port either rounding across dtypes.

Commit: ork-driver main b351c14. Probe: tools/re/i16_mcap_probe.c.

Applied the fp16 playbook to the int16 activation path (orki_i16_act_lut, shared by ork_i16_npu_{exp,silu,gelu,rsqrt}). Two results.

The roadmap item does not reproduce. Tier 14 A #2 claims "ork_npu_exp_i16 has an unmapped row-count envelope, NRMSE 0.55–0.69 @ MR=6144/4096". The op is bit-exact at all 37 probed points — M=8..8176 at N=64 and N=8..2048 at M=512 — including M=4096 and M=6144 exactly. The knobs that item cites to isolate it (ORK_ATTN_TNORM_OFF, ORK_ATTN_MAX_CPU, ORK_ATTN_SM_SUMCPU) do not exist in ork-driver; they are fork-side, so the NRMSE was measured through the fork's attention composition, not this op. Re-filed against the fork.

But a real bug sat underneath it. The guard accepted M<=8192 while the hardware stops at 8176: M=8184 returns errno=110 (submit TIMEOUT) plus a self-healing reset — the guard was admitting a fault-generating shape rather than refusing it. 8192 was never measured; it is RK_DPU_DATA_CUBE_WIDTH's 13-bit field (mask 0x1fff ⇒ M-1 <= 8191) read as a capability — the same provenance as the fp16 bound fixed the same day. It is a row limit, not a size limit: 8176 is bit-exact at N=64 and N=128 (twice the bytes).

The guard is copy-pasted 21× across the SDP family, so the fix adds one ORK_SDP_MAXM constant and applies it to the 8 int16 sites. The f16 perchan/replay and i8 SDP sites are deliberately left at 8192 — same geometry patcher, likely the same ceiling, but unmeasured, and propagating an unmeasured bound across dtypes is exactly what put an int8 constant on the fp16 datapath. Probing them is the open follow-up.

Pattern worth naming: three bounds audited this session, three that were inferred rather than measured — fp16's mg_max*64 (from int8), fp16's sched=0 cap (an int8 element count at 2 B/elem), and int16's 8192 (a register field width). A field width is an upper bound on what you can encode, never evidence of what the hardware completes.

Commit: ork-driver main 6584e2a. Probes: tools/re/f16_mcap_probe.c, tools/re/f16_k128_probe.c, tools/re/i8_mcap_probe.c. Full writeup: Exp 2026-08-20 fp16 M Envelope and the CBUF Bank Split.

Tier 14 A asked for a refusal on out-of-envelope M; measuring the envelope first showed the envelope itself was wrong. 0x1040 is a CBUF bank split — v = (WEIGHT_BANK<<4)|DATA_BANK, DBNK+WBNK == 12, walked one bank data-ward per 64-row group, bank = 32 KB — so the fp16 rule is just banks × 32 KB / row-bytes: sched=0 → M ≤ 16384/K, sched=1 → M ≤ 180224/K. Measured 704@K256 / 352@K512 / 176@K1024 (vs the inherited 320/128), and 8@K2048 … 256@K64 for sched=0, each with ceiling+1 confirmed wrong. Two live miscomputes fixed: ork_f16_mm_run_stream's guard was the int8 byte budget (M*K>32768), 2× too loose at 2 B/elem → wrong output at non-pow2 K; and ork_f16_mm_run miscomputed at K=128 for any M ∈ [257,1472] (4*R gave a 1024-row tile against a real 256). Both now tile to one measured orki_f16_mcap(K,sched), which replaced five drifted copies of the old formula; only inexpressible cases (differing M across a batch) refuse. K=128 is broken by encoding — it needs half-bank steps (slope=(int)7.5=7), so its envelope is non-monotonic (256 ok, 320 bad, 384 ok, …) and only the 256-row prefix is usable. make test byte-identical.

Why it went unnoticed: silu_f16_check had validated fp16 at 320 and 384 — never 352 — so the boundary sat between the probe points. Two methodology rules came out of it: never bisect this predicate (non-monotonic), and never derive an fp16 tile from an int8 constant.

int8 checked and unchanged: predicted the same win, measured none — int8 is already correct at its truncated mg_max*64 (704/320/128/64 at K=512/1024/2048/4096). Why, RESOLVED: the split walks data-ward as M grows, so WEIGHT_BANK shrinks — and int8 keeps its weight CBUF-resident (dies once WBNK no longer fits it) while fp16 streams its weight (unaffected, runs to the full 11 data banks). Closed form for int8 at N=64: M_max = 64·floor(12/w − 1), w = K/512, exact at K=512/1024/1536/2048; K=1536 was the discriminating measurement (bit-exact at M=192 with WBNK=3, every element wrong at M=193 where the clamp drops WBNK to 1). They coincide only when the weight is exactly 1 bank, which is why int8 K=512 and fp16 K=256 both reach 704. Predicts an N-dependent int8 ceiling (untested). One corner unexplained: K=4096.

Board-ops: forcing v below the 0x1b floor (WEIGHT_BANK=0) hangs the submit; an external timeout then killed it mid-submit and hard-wedged the board (plug cycle recovered it, no SPI damage). Use ORK_MM_TIMEOUT so the driver reaps its own job.

Commit: ork-driver main 96efa7a (include/ork_native_cpu.h batched ork_cpu_gemm_i4/ork_cpu_gemm_nf4, probe examples/test_i4_gemm.c) + llama.cpp-rockchip diverged-history 4cf8661d0 (ggml-ork.cpp ORK_MOE_CPU route + run_cold batched path + footprint residence; submodule → 96efa7a).

Prefill is compute-bound on expert dequant (weights read ~once/batch ≈ 1 GB/s ≈ 5% of the bus), so a cheaper unpack on the CPU beats native ggml Q4_K. Store experts as NF4-int4 resident in the orkpack, run them with a batched (M>1) 4×4 NEON GEMM (NF4 unpack = one vqtbl1q_s8, no Hadamard), keep attention on the NPU (int8).

A/B — 35B Qwen3.6-A3B, RK3588, mixed.gguf base, -t4, r=3, governors=performance, same board:

Kernel bit-exact vs the M=1 gemv (I4 2.3× / NF4 1.85×, NF4 faster absolute). Correctness validated (ork_bench: "What color is the sky?" → "Blue"); make test ALL PASS; decode unchanged (native path, 4.4 t/s pre/post — the route is a prefill-only lever). Recipe: build the orkpack from an f16 gguf with ORK_ORKPACK_I4_FFN=1 ORK_MOE_CPU=1 (experts int4/NF4, attn/dense int8), no ORK_QUANT. A first attempt with ORK_QUANT=4 forced dense int4-on-NPU → bcreate IOVA wall → a contaminated 6.17 "not a win" (a scrutinize-negatives miss). Full writeup: Exp-2026-08-18-NF4-CPU-Prefill. Sibling experiments stashed this saga (patches in attachments/, snapshot branch wip/moe-saga-2026-08-18): big+little CPU dispatcher ❌ parity/HW-defeated · int4 expert-parallel coalesce ⚖️ 1.6× on matmuls, unwired · GPTQ int4 🧪 unvalidated (task #56) · int4 MoE A/B knobs (mpad/m1npu flat, removed).

Commit: ork-driver feat/ork-async-stream — src/npu.c (ork_dyn_begin_colsplit CONTIG default-on w->Bbc; is_valid_dma_addr +Bbc/Bgap; adaptive sentinel detect = 1.5× measured slice; reap_stuck+ork_dummy_probe recovery net; reset-first + short reap-timeout), examples/test_f16colsplit.c (ORK_CAMPAIGN per-iter bit-exact), diagnostics ORK_F16_TRACE. Kernel source cross-ref: rknpu_job.c (rknpu_job_abort/timeout_clean), rknpu_iommu.c (iommu_domain_refcount).

The fp16 doorbell colsplit had an intermittent (~every op at K=3584) drop that made it unshippable (gated off). Root-caused and fixed by prevention. Two things had been conflated:

• The drop is a wild DMA, not the benign int8 dispatch-miss. The 3 cores loop K-slices independently; the NPU prefetches Bb[ks+1]'s weight base while Bb[ks]'s fp16 fetch still drains → two distinct dma-buf fds on the CDMA → mis-addressed fetch. Always the boundary slice (Kslice=1, never Kslice=0) — the cross-boundary-prefetch signature. A CPU lockstep barrier can't gate a HW prefetch.
• A real fp16 drop is UNRECOVERABLE from userspace. Exhaustively tried: in-place resubmit → sticky re-drop + D-state hang; close(fd) → refcount underflow; timeout_clean reap + fast-detect + reset-first → kernel panic. Any BLOCKING fallback that can't dispatch enters the 60→180 s continue_wait → rknpu_job_abort, which puts the IOMMU domain 100 ms before halting DMA; the device-global iommu_domain_refcount (gates iommu_detach/switch on ==0) is over-decremented by the dropped job's double-cleanup → detach under a live matmul → wild write → slab-corruption panic in unrelated procs, delayed 1–2 min.
• THE FIX = CONTIG: concat all K-slices into ONE weight buffer (w->Bbc) → no dma-buf boundary → prefetch stays in one mapping → no wild, no drop. 3 cores read regions of the one buffer (benign concurrency). What hid it for weeks: is_valid_dma_addr had no case for Bbc, so a valid CONTIG weight base was false-flagged "wild/unallocated" → validate_regcmd refused → fell back to the wedging per-slice path. One-line fix.

Validation (256×3584×3584, the drop-prone shape; per-slice dropped on iter 0):

Full shape suite (Sk=1 and Sk=2) BIT-EXACT; make test all-PASS. CONTIG is now the default for Sn==1 (ORK_F16_NO_CONTIG opts out). The recovery protocol (adaptive detect + clean timeout_clean reap + int8 health-gate) is retained as the safety net for the benign miss class (int8/int4 doorbell drops, proven 12000 iters) and any future Sn>1 fp16 shape. Verdict: prevention beats recovery — for a wild-DMA drop the corruption is done before any detect can fire, so the boundary must be removed, not the drop caught. Full writeup: Exp-2026-08-05-fp16-Colsplit-CONTIG.

Commit: ork-driver feat/ork-async-stream — src/npu.c (ork_npu_fold_run_i8, ork_npu_fold_batch, ork_npu_sdp_stamp), src/regcmd_fold_refs.h (14 baked per-size templates), tools/re/{fold_tiler,bigm_run,ab_fold}.c + persistence-aware regcmd_decode · commits c07c4a2→96c7c42. Offline analysis: full_sdp.py / full_regmap.py / decompose.py over the 583 MB ~/pf.dump prefill capture.

The 2026-07-28 M-fold RE concluded the fold was chain-required behind a 12-bank CBUF-capacity wall, 1.66× as a single-core replay only, not integrated. Both halves of that verdict were wrong, and the fold is now reproduced end-to-end and integrated as a lib run-path.

What the fold actually is (decomposition, offline over 119k–134k captured tiles). Not a chain of accumulating K-slices — plain token-tiling: a batch of M_total tokens is split into fold sub-tiles of M ≤ 36 (set {36,32,28,24,20,16,14,12,10,8,6,4,2,1}) that sum to M_total; each sub-tile sets surfstr(0x4024)=16·M_total, writes its M rows at a row offset into ONE shared M_total×N output cube, streams its own weight (weight IOVA constant, full N, WEIGHT_REUSE=0 in every tile), and the sub-tiles run as a single multi-task submit. There is no weight reuse, no weight retention, and no CBUF-capacity wall. The earlier "wall" (a self-contained M=36 tile hard-wedging) was a self-inflicted bug: forcing surfstr=16·M (standalone) onto a tile that is really a sub-tile with surfstr=16·M_total. Fix the stride + pack the input at the batch width and a real M=36 tile runs standalone, bit-exact (bigm_run). Only 4 registers depend on M_total (0x4024=16·M_total, 0x107c=M_total, 0x1080=M_total−M, 0x40c0=128·M_total); everything else is per-M geometry. The output stage is one invariant config across the whole prefill, spanning two blocks (0x1001 DPU/SDP + 0x801 PDP/aux) and reconstructable from first principles (ork_npu_sdp_stamp, bit-exact). The persistence-aware regcmd_decode (never presumes a self-contained tile) quantified it: an M=36 sub-tile inherits 23 registers (20 output-stage + the op-enable doorbell) — none weight-related.

A/B — token-fold (ork_npu_fold_run_i8) vs ork-normal (ork_mm_run_i8), single-core, K=3584 N=1216, bit-exact (identical C):

The fold beats ork-normal by 1.6× at the prefill sweet-spot M (36–72), tapering to ~parity at M=128 where baseline's mg_max·64 tile catches up. This retires the "M-fold has no perf niche" note — that verdict was about the M≤8 chain (re-streams weight); the real fold is big-M tiles in one submit. (fold µs = pure submit avg; base µs = full call — gap holds either way.)

Shared-context fix (integration blocker): fold_batch reused c->task, which lives in the default IOMMU domain, while submitting in dom_active — an EINVAL once the standard path activated a domain. It now allocates its own task buffer in dom (like the multicore c->mtk[i]); fold + ork_mm_run_i8 now coexist in one context, bit-exact.

3-core — the real full-op test (measured): fold wins small M, LOSES large M. The single-ioctl core_mask=0x7 split wedges ork; the proven path is concurrent per-core submits (core_mask=1u<<c). Built ork_npu_fold_op_i8 = rkllm's real wide-projection scheme (N-split 3584 across 3 cores, each core single-core-folds its ~1216 slice, concurrently). Measured vs ork-normal ork_mm_run_i8, both 3-core, K=3584 N=3584, bit-exact every M: M=36 1.35× · 72 1.31× · 100 1.09× · 128 0.89× — crossover ~M=100. Cause: the fold re-streams its slice weight per 36-row tile (4×/core at M=128) and 3 concurrent cores saturate DRAM, vs ork-normal's 64-row tiles (2×/core). So at prefill-typical ubatch (M≥128) the fold is wash-to-loss on the attention projections and won't close the 7B prefill gap as-is. (This supersedes both the earlier mis-framed "wash" and an over-corrected "1.6× at 3-core" — the honest answer is M-dependent.) The unlock is weight reuse: stream each slice's weight once and fold all M rows against the resident copy → lifts the DRAM-BW ceiling → the fold would win at every M. That is the open lever (the WEIGHT_REUSE bit measured 3.4× amortization but currently miscomputes without weight retention). Loader-recipe sweep ran (ORK_WR_MODE, commit 4c06063): definitively negative. WEIGHT_REUSE mechanically skips the weight DMA and hits exactly the target (M=128 1639->1034 us, 1.6x) but computes WRONG (weight not retained); +DATA_REUSE faster still + wrong; +FC_SKIP_EN errors the submit. Triangulated with the header (no retention register) and rkllm never setting the bit: cross-tile weight reuse is NOT achievable on RK3588. The fold's per-tile weight re-stream is unavoidable, so it stays a small-M-only win. Vendor capture — definitive (commit 7efe58c): captured librknnrt's own int8 matmul (rknn_matmul API, B set once, run 4x on one ctx) under the LD_PRELOAD shim. Every run byte-identical: 0x1040=0x84 (WEIGHT_REUSE=0), 0x1060=0 (FC_SKIP_EN=0), weight-IOVA re-fetched each run. So the vendor's own matmul re-streams the weight every run and never uses reuse/skip — it is NOT an RE gap, it's a real limit: WEIGHT_REUSE/FC_SKIP are conv-only (intra-op spatial), the matmul/FC path has no cross-op weight retention on RK3588. rkllm = librknnrt = ork's fold all re-stream weight per matmul tile (a hardware floor). That is why the fold is small-M-only and ork's mg_max·64 (more rows/weight-stream) is the correct lever, not reuse. #39 fully and definitively closed. -- #39 closed: RE-complete, integrated, measured, honestly bounded (no end-to-end 7B lever; ork is already ~98% without it).

Why 1.6× — the mechanism (compute-under-DMA overlap, NOT fewer fetches). Neither N-tiling nor amortization explains it: nmax=8192, so at N=1216 both paths do one weight stream per tile over full N (no N-tile), and ork-normal actually amortizes better (64-row mg_max·64 tiles vs the fold's ≤36). The single-core timing decomposes cleanly and the model fits every point. ork-normal per tile ≈ weight-DMA (~396 µs = the 4.36 MB K·N int8 stream @ ~11 GB/s) + ~8.5 µs/row of SERIAL, exposed compute — it streams the weight, then computes, un-overlapped. The fold per tile ≈ ~420 µs FLAT (≈ the weight-DMA floor) for its ≤36 rows — the row-compute is HIDDEN under the weight stream. Model vs measured: ork-normal Σtiles(396+8.5·rows) → 702/1404/1642/1880 µs at M=36/72/100/128 (meas 722/1400/1654/1861); fold ~420·⌈M/36⌉ → 420/840/1260/1680 (meas 440/843/1236/1639). So the fold's NC1HWC2 M-in-WIDTH layout keeps the MAC array fed from the streaming weight (weight-stationary reuse across the width = the token rows), turning a DMA-then-compute op into a DMA-bound one. It does NOT reduce weight fetches — at M=128 the fold streams the weight more times (4 tiles of 36 vs ork-normal's 2 of 64); the win is purely hiding compute under the fetch. Hence largest at small M and gone by M=128, where ork-normal's bigger tiles amortize their serial DMA. This is a fetch/compute-overlap finding: the lever is pipelining compute under the weight fetch, not the fetch count.

Why rkllm's single-ioctl 3-core works but ork's wedges (UMD setup, not a HW limit). ork has never issued a single core_mask=0x7 submit — its whole multi-core path (mcworker) is PER-CORE submits (core_mask=1u<<i, one ioctl/core, own c->mtk[i]). rkllm's task_number=36/subcore 12/12/12 single-ioctl works under librkllmrt; ork's 0x7 + subcore_task split + our in-regcmd descriptor chaining hard-wedged. Notably rkllm's tiles mostly OMIT the chain descriptor (0x0010 present in only 6.5% of tiles → kernel re-arms per subcore) whereas fold_batch writes it — a plausible culprit. ork doesn't need the 0x7 ioctl: its proven per-core concurrent submits give the same 3-core parallelism, and that's the correct path for real 3-core.

Verdict: the M-fold is real, reproduced, and a 1.6× single-core prefill kernel at M≤72 — RE-complete and integrated (baked templates, no proprietary blob, no /tmp deps). But ork already reaches ~98% of rkllm on 7B prefill via mg_max·64 + the K-schedule + 3-core doorbell — without the fold — so the fold's marginal end-to-end value stays open (largest at small M, gone by M=128) — and 3-core does NOT compound it: measured a wash (the fold under-parallelizes at small M and ties at large M, while ork's 3-core splits any M evenly). Next, if pursued: a gated end-to-end llama-bench A/B — though this A/B predicts the fold won't move 3-core prefill, consistent with ork already at ~98% of rkllm without it.

Stack: ork-driver feat/ork-async-stream (+ ggml-ork fork). Board RK3588 (10.3.0.236), governors=performance, single-stream, resident orkpack. Full writeup: Exp-2026-07-28-7B-Prefill-Parity.

The Prefill Parity Grind left Qwen2.5-7B at ~40 t/s (54% of rkllm) with the diagnosis "the rest is integration, not a matmul wall." That held: the levers that landed after the grind now stack to 71.89 t/s, ~98% of rkllm's 73.07 (both warm, resident, churn=0) — and none of them is rkllm's M-fold. The path, in contribution order: (1) weight-DMA-amortization M-tile mg_max*64 — ork-native, the biggest lever (Weight-DMA Amortization); (2) ported rkllm K=3584 (Kp,M) schedule dial — the one genuine rkllm graft, the K-axis, not the fold (task #38); (3) 3-core NONBLOCK doorbell colsplit — run_multicore→ork_dyn_begin_mc, so the parity number is a doorbell number (Dynamic-Submit & Doorbell); (4) fully-resident churn-free orkpack across 8 IOMMU domains. The M-fold (task #39) stays out: 1.66× as a single-core kernel replay only, A-layout proprietary, not integrated. Caveats: quant paths differ (ork Q8_0 / rkllm w8a8), tokens not matched (ork 256 / rkllm 485 — favours rkllm), decode not parity (prefill is the point). Equal-token doorbell A/B is the one open confirmation.

Commit: ork-driver · src/npu.c (synth_i8_mfold, ork_npu_synth_i8_dump ORK_MFOLD, ork_npu_replay_i8 0x40c0 patch) · probes tools/re/{synth_dump,validate_mfold,nc1hwc2_test}.c. Full writeup: Exp-2026-07-28-M-Fold-Schedule-and-CBUF-Bank-Capacity. See regcmd ISA Reference.

Reverse-engineered rkllm's per-program-faster int8 schedule: it folds M into the CNA WIDTH dimension with an NC1HWC2 feature layout (weight unchanged = ork's tiling). Built synth_i8_mfold, which reproduces the captured regcmd bit-for-bit (M=2 and M=24). But run as a single task it computes a deterministic partial reduction — all-ones M=16/K=3584 → uniform 2400 = 150 of 224 input surfaces. Root cause via rocket_registers.h: the "schedule" regs are the CBUF bank/entry allocation — 0x1040 = CNA_CBUF_CON0 (DATA_BANK/WEIGHT_BANK), 0x1044 = CNA_CBUF_CON1 (DATA_ENTRIES). A single submit reduces exactly DATA_BANK × bank_capacity channels (DATA_BANK=3 → 2400; override to 5 → ~4000 > K → stall; small M → stall). ork's normal path uses DATA_BANK=1, WEIGHT_BANK=11 (weight-resident) and streams full K — hence correct for any K. The capacity steps never equal an arbitrary K, and raising DATA_BANK evicts the resident weight → the fixed 12-bank trade-off that forces rkllm's K-slice chain. Also decoded the hard-wedge cause: 0x40c0 = DPU_SURFACE_ADD is an absolute output IOVA left unpatched (a size value) → wild WDMA write → AXI hang; patching it makes the path soft-fail (errno-110) instead. Verdict: mfold is a chain-required kernel (its ~1.6×/program needs the multi-submit K-slice accumulate); RE-complete + synth-proven, banked as a documented candidate — payoff speculative since ork already beats rkllm on 7B prefill. Corrected register names feed the ork-driver register-naming refactor.

Commit: ork-driver feat/tiered-i4-residence · src/npu.c (mcworker K-slice loop chunk/sched at ~:3888, synth_i8 0x1040 at :854) · probe tools/mc_prof.c (ORK_KTILE). Full writeup: Exp-2026-07-27-K-Slice-Size-Sweep-and-the-Sched-Optimum. See Weight-DMA Amortization + PC-chaining M-tiles.

Swept the K-slice size KS (default 1024) on the 7B ffn_down (K=18944, N=3584, M=228 — the only matmul forced onto the K-slice accumulate path, K>kmax=4096). Result: 1024 is a sharp, provable optimum (764 GMAC/s); every other value is 30–60 % slower (512→486, 896→536, 1184→579, 4736→313; 2048 WEDGED core0). The cause is a single flag, sched = (Kp==1024||Kp==512), that gates two coupled effects: (1) the M-tile chunk = mg_max*64 (320 rows at 1024, so M=228 streams the weight slice once) vs sched=0's (RB/2)/Kp ≈ 54 rows (~5× weight re-stream — dominant, since the matmul is weight-DMA-bound); and (2) the 0x1040 K-reduction schedule (validated only at Kp∈{512,1024}). They're inseparable — the big M-tile is only bit-exact with the matching 0x1040, so leaving 1024 forces the small tile. 1024 beats 512 (both sched=1) because it's the larger validated point → 19 slices vs 37, halving the per-slice fixed overhead (submit + A-gather + host accumulate). Correction: an intermediate "17/19-slice bracket" claimed to hold the M-tile constant and isolate the schedule — it does not (sched=0 forces the 54-row tile); M-tile and schedule vary together, and the M-tile/weight-DMA is the dominant term. Verdict: KS is not a lever — the axis is exhausted. The two live levers for wide-K: (1) pipeline the host int32 partial-accumulate against the NPU (~94 ms/core vs ~235 ms NPU @M=228 → 1.69× ceiling, wedge-safe — the one that can still move it); (2) extend sched=1 to larger Kp (HW-RE, uncertain — mg_max shrinks with Kp, so no M=228 win, and it wedges).

Commit: ggml-ork fork feat/ork-async-stream · ggml/src/ggml-ork/ggml-ork.cpp (ggml_backend_ork_ffn_swiglu_chain; OrkFfnPool work-stealing pool, OrkNpuDriver spinner) · env ORK_FFN_SCHED/_DRIVER/_POOL_CORES (doorbell, production) + ORK_FFN_APIPE (async-worker, documented+removed). Full writeup: Exp-2026-07-26-CPU-NPU-FFN-Overlap.

The int8 FFN SwiGLU chain runs the exact CPU fp32 SiLU (the fused-fp16/int8 SiLU is quality-dead); it's serial and hideable. Overlapping silu(gate) with the independent up matmul + parallelizing it across cores gives +6.5% prefill (141.7 → 151 tok/s, PPL bit-identical 18.3891) on qwen3-1.7b-q8_0 (RK3588, P=256). Two mechanisms reach the same ceiling: apipe (up via ork_mm_run_i8_async = blocking kernel-IRQ worker that SLEEPS, freeing cores; silu on a work-stealing pool) = 150.97, and the doorbell (ORK_FFN_SCHED two-chain + a dedicated spinner) = 150.93.

The durable finding: it's core placement, not bandwidth. Throughput swung 112↔151 purely on which cores do silu. DRAM-bandwidth is refuted (the config with the MOST silu cores was FASTEST); doorbell-fundamental-limit is refuted (a dedicated spinner didn't change it); starvation is refuted for the pooled case. The doorbell's colsplit threads pin to 0..nc-1 and its spinner occupies a core, so silu must avoid them (explicit nc + poll-only spinner on a little core + silu on the freed big cores). apipe tolerates silu-on-big because its worker sleeps; the doorbell must keep silu off its spinning cores. The generalized ork_spine_run dynamic scheduler is correct but slow (orchestration tax); the fix is a prescribed schedule on the spine units, planned once — see PRESCRIBED_TUNING_PLAN.md (pack-time calibration bakes the per-(chain,shape) core-map into the orkpack). apipe is documented here and stripped (it duplicated the doorbell's role at equal perf); the dead ORK_FFN_SILU_OVERLAP (fp16 async) and ORK_FFN_GATE_M48 (fused-fp16-SiLU gate — two-signed band rejected by the single-signed fp16 LUT index) removed with it.

Commit: ork-driver feat/tiered-i4-residence · src/npu.c (ork_dyn_begin/begin_mc/progress/halt/end, ork_dyn_queue_*, ork_pc_*) · tools ork_dyn_test.c/ork_dyn_queue_bench.c/ork_pc_bench.c/dyn_spin_probe.c. Commits b94a4aa (API+queue), b9ce004 (multi-core NONBLOCK), 83ea8b2 (precompiled cache). Full writeup: Exp-2026-07-16-Dynamic-Submit-API-and-Doorbell.

A userspace NONBLOCK submission layer over the chain path (flags|=0x2 + per-output dc civac doorbell) that lets the host watch per-op progress, halt mid-flight to free the NPU early, and rendezvous without a thread — all on the stock kernel. It pinned the submission model: a job runs exactly task_number programs then completes; the count is fixed at submit. So you can HALT (run fewer — 0x0014=0, validated) but you can NEVER LOOP or EXTEND (run more): 0x0010-redirect isn't a runtime jump (the walk is task-indexed) — both ork_dyn_append (forward extend → job abort -22) and a circular spin (0x0010 self-loop → walked forward, corrupt) fail. Two validated wins that stack: multi-core NONBLOCK stream (begin_mc) = 2.27× over single-core, 1.38–2.30× over the blocking 3-core stream (doorbell rendezvous beats thread-join) — the dispatch half; precompiled-program cache (ork_pc_*, fixed chain + pinned buffers) = 6.03× lower per-token build (2652→440 µs, ≈ raw NPU floor) — the build half. The wrap is chunk+resubmit (persistent-job needs a driver/MMIO change — not viable via 0x0010). Wired into ggml-ork's #14 MoE decode (fork feat/ork-async-stream @ 0a80450ef, ORK_DYN_MC, gated single-domain) and validated END-TO-END on LFM2.5-8B-A1B: ~1.6× faster decode (6.5–6.7 vs 4.0–4.2 tok/s, reverse-order confirmed), coherent. ork_bench/ork_ppl are now cmake targets.

Commit: ork-driver feat/ssm-npu · src/npu.c (ork_bmm_fp16_fused sched fix) · tools/f16_shape_probe.c (sweep + ork_f16_fuzz_* register bisect) · reproducer tools/ssd_fusedchain_probe.c. Full writeup: Exp-2026-07-12-Matmul-Shape-Boundaries. See NPU-Quirks + regcmd ISA Reference.

The on-NPU SSD/Mamba-2 scan's fused batched-GEMM (ork_bmm_fp16_fused) emitted zero output at K=64 where the per-op ork_bmm_fp16 is correct. Root-caused to a single register: bare synth() with sched=1 programs the 0x201:0x1040 K-reduction schedule from a formula (base=177−15·(K/256−1)) that overshoots the template default 0xB1 below the prefill band (0xBC@K64, 0xBE@K32) → an invalid reduction atom → zeros. A silicon register-bisect (override only 0x1040 0xBC→0xB1) flips it zero→bit-exact (3.8e-8). The boundary is purely K-driven, N-/M-independent: fp16 breaks at K<96 (K∈{32,64}), int8 at K<256 (formula scales by K/512). Production paths (ork_mm_run/mcworker) were already safe (they pick sched per shape); the fused chain + raw probes hardcoded sched=1. Fix: match the run-path predicate (sched = pow2 K∈[128,2048)); validated bit-exact across the scan shapes (K∈{32,64,96,128}, N∈{64,128}). Reconciles the earlier FLOOR_DECOMP "raw-synth_i8 N=1024 half-cols" note — same K-schedule defect.

Commit: ork-driver feat/ssm-npu · src/npu.c (ork_npu_probe_silu_std_i16 task save/restore fix; mm_timeout_ms/ORK_MM_TIMEOUT; ork_npu_mode_invalidate/ork_npu_mode_reset hooks) · tools/mode_probe.c (the 2-op transition probe) · reproducer examples/test_ssd_chunk_npu.c. Full writeup: Exp-2026-07-12-NPU-Mode-Transition-Task-Descriptor. See NPU-Quirks.

test_ssd_chunk_npu mode 2 (whole SSD/Mamba-2 decay-elementwise path on NPU) reproducibly wedged (errno=110) on the fp16 CumBA bmm (K=64,N=16) mid-sequence, while mode 1 (same exp→bmm) survived and the CumBA shape was clean in isolation. A minimal 2-op transition probe (tools/mode_probe.c, sweeping {matmul dtype} × {SDP op} as separate submits, with a fast ORK_MM_TIMEOUT so wedges self-terminate) pinned it:

Root cause: the LUT-activation ops (act_lut_i16 → ork_npu_probe_silu_std_i16) overwrite the shared c->task descriptor (regcfg_amount 1097→69, enable_mask 0x18) and never restored it; the single-core matmul path relies on the persistent init descriptor (108/0xd) and doesn't rebuild it, so it submits a 108-word matmul regcmd under a 69-reg SDP task → no task dispatched (task counter 0x0) → timeout. Multi-core matmuls rebuild per-core c->mtk (immune); ewmul/add already save/restore. This is why mode 1 (all N≥64 → multi-core) survived and mode 2 (adds CumBA N=16 → single-core) wedged. Confirmed not the NVDLA "SDP-can't-feed-MAC" ordering law (refines the earlier Activation→matmul ordering quirk): ewmul→matmul is same-direction yet safe, and ACT_RESET/ORK_NO_PINGPONG/mode-invalidate all fail to clear it (it's a poisoned software descriptor, not HW state). Fix: save/restore c->task in the LUT op exactly like ewmul does — every transition then SAFE with no interposed reset, and mode 2 goes to PASS (full-G1 rel-L2 5.8e-3, full-G2 5.9e-3, bit-coherent). Safe-sequencing rule: every op must leave the shared c->task as it found it; a fused single-submit chain is not required for per-op correctness.

Commit: ork-driver main 32db14c (+ f4b0154) · src/npu.c, src/neon_activations.{c,h}, examples/test_bmm.c (in make test), tools/re/test_rsqrt_lut.c. Gates ORK_NORM_NPU / ORK_ATTN. Full writeup: Exp-2026-07-10-On-NPU-Attention-Batched-GEMM-and-Norm.

New primitives for the ops the static-weight matmul path can't serve: ork_bmm_i8/i4/fp16 (batched dynamic GEMM — both operands dynamic, the attention QKᵀ/AV + Gated-Delta-Net chunked-matmul primitive; i8/i4 exact, fp16 4e-4); ork_npu_rmsnorm_f16/l2norm_f16 (the full-channel Σx² reduction runs on the NPU as a matmul contraction, any n via K-split, + on-NPU rsqrt; 1.3e-3); ork_mm_build_f16_rsqrt_lut (bakes 1/√(ss/n+ε) into the fp16 fused-output LUT — the fp16 fused-into-the-reduce twin of the standalone ork_npu_rsqrt from Exp-2026-07-04; fused K=n≤2048 6e-4 / decoupled K=512 any-n 5e-4); ork_l2norm_f32 (NEON CPU). All gated off + self-validating in make test.

Perf verdict: neutral end-to-end. ORK_ATTN on Ornith-9B: GDN chunked matmuls route but NPU stays 1.6% / prefill unchanged — the CPU wall is the sequential GDN recurrence (gated_delta_net/solve_tri/cumsum/exp), not matmuls. Pure-attention (Qwen3-1.7B): the real QKᵀ/AV are non-contiguous → declined, not offloaded. A standalone norm/attention op is submit-floor-bound and loses to the fused CPU/NEON pass; the win needs fusing into an adjacent matmul's output stage. RE landed: the DPU output-stage registers are named against the mainline accel/rocket header (0x4040 DPU_BS_CFG, 0x4070 DPU_EW_CFG+LUT, DATA_CUBE_* geometry), and NVDLA's CDP LRN window is HW-capped at n∈{3,5,7,9} — why a full-channel norm reduction must ride a matmul contraction, not the fixed-function LRN.

Commit: ork-driver main v0.6.66 · src/npu.c (live-buffer registry live_add/live_del in bcreate+bimport+bdestroy; ork_sig_teardown handler; sigaction in ork_npu_init; env ORK_NO_SIGCLEAN). Fork master (sub v0.6.66).

Surfaced wiring Ornith-1.0-9B (dense qwen35, W4A8 inflated int8 ≈ 9–19 GiB across 5–7 IOMMU domains). A process killed mid-run (timeout SIGTERM, Ctrl-C) that skips MEM_DESTROY strands its multi-domain IOVA mappings — the kernel's fd-close cleanup doesn't reclaim the domain-allocator state, so the NEXT run fails PRIME_FD_TO_HANDLE: Cannot allocate memory / domain N full / bcreate failed while RAM & CMA are free (it's IOVA, not memory). Single-domain (≤4 GiB) never hits it. Fix: track every live handle and, on SIGTERM/SIGINT, MEM_DESTROY them all + ACT_RESET before re-raising — the same reclaim a clean exit does. Validated: SIGTERM a fully-loaded 9B mid-decode (exit 143) → immediate re-run clean, no leak, no reboot (previously required sudo reboot); single-domain i4 still bit-exact. Also observed (open bug): multi-domain layout is non-deterministic — the same 9B W4A8 orkpack loads to 10 GiB/5-domain (46 tok/s, NPU active) or 19 GiB/7-domain (4.3 tok/s, swap-bound, NPU 0.4%) via Bf (wide-K inflate) bloat. Full-system telemetry (good layout, bench_monitored.sh): NPU 12.8%/32%, RAM-BW 23/61%, CPU 23/54% — nothing saturates; multi-domain is swap/orchestration-bound.

Commit: ork-driver main 848c6f7 (+ CI e98e30e) · src/npu.c (run_i4_bchain/i4_bchain_worker ORK_I4_BCHAIN; run_i4_cbatch ORK_I4_CBATCH gated scaffold; knobs ORK_I4_BCH_H/_W, ORK_I4_CB_REGS/_AMT). Fork master 2935a75d7 (sub v0.6.62). Full writeup: Exp-2026-07-08-INT4-BChain-Batch-Chain.

Unified STRATEGY A (in-task H-row batch, H=16384/K) with STRATEGY C (chained tasks, weight loaded once): chain the H-row batches, M/H tasks per weight stream. The naïve full-N batch hits the 0x1040 K-reduction wall (only row 0 at K=2048, N>64); the fix is tiling N by the weight-bank width Wb=131072/K (=64 @ K=2048, since one 64-block int4 weight at K=2048 = exactly one 64 KB CBUF bank, and there is no D_BANK) so each chunk is a wall-free single-bank batch, chained across (M-group × N-chunk). Multi-core splits the chunks across the 3 cores (core_mask=1<<i).

All bit-exact (PPL identical to 4 dp; i4 example maxerr=0). Matmul: 2.62× @N2048, 2.32× @N6144. CBATCH (the batching-chains dual — carry the batch in the lean compact task) is HW-rejected on RK3588: a multi-row compact errno-110s for every register config (kept gated for other NPUs). ⇒ the duality is asymmetric — chain-batches ✓, batch-chains ✗; BCHAIN's full-task-per-batch is required. 3-core scaling is sublinear (~1.6×, contention-bound), and with the run halved the prefill is now quant-bound (quant 36% / run 57%) — the next lever is CPU-side activation FWHT+int4-quant (NEON / reuse / overlap), not the NPU.

Commit: ork-driver main 473b90e · src/npu.c (ork_i4_msched, synth_i4 mc>1 ORK_I4_MREGS/ORK_I4_1040/ORK_I4_1010, i4_mcworker native-multi-M branch, ork_npu_probe_i4_mm 2× fix) · tools/i4_multim_probe.c · fuzzer tools/i4_multim_fuzz.c. Full writeup: Exp-2026-07-07-INT4-Multi-M-Schedule-RE. Resurrected native int4 batch mode (Exp-2026-06-19) behind ORK_I4_MSCHED (OFF; per-row path unchanged). Fuzzing proved reg 0x1040 (K-reduction schedule) is decisive — the int8 formula corrupts int4; without it (mregs=0x1f) batch mode is bit-exact at K=64 (disproves the 06-20 "structurally single-row" verdict). Remaining wall: row count is capture-K-tuned (K=64→4 rows, K≥512→1) so production K=2048 computes only row 0 — needs a systematic 0x1040 × K × mc × CBUF sweep.

Commit: ork-driver branch ewmul-re — d7482fe (register-operand EW-mul validated) · 8d30bbd (REGCMD_MUL standalone element-wise op executes) · 55babcf (calibrated bit-exact) · src/npu.c (ork_npu_probe_i8_mul, ork_npu_probe_i8_ewmul[_tmpl/_lin], set_i8_ewmul, splice_ew_lane, apply_ork_geom, ew_timeout_ms) · src/regcmd_ewmul.h (REGCMD_MUL, REGCMD_EWMUL[_LIN], REGCMD_EW_LANE) · tools/ewmul_probe.c · tools/regcmd_capture.c (RKDUMP_WORDS). All rk3588-gated; RE/probe-only (not yet wired into ork_mm_run).

Chased the second half of the FFN-inner fusion — the SwiGLU elementwise silu(gate)⊙up — onto the NPU. The path was long but ended with a genuine result: the on-NPU element-wise multiply works and is validated bit-exact on silicon.

The breakthrough: the RK3588 NPU is an NVDLA derivative. After the initial register-patch attempts wedged the NPU (core0 raw status 0xc0000015, no IOMMU fault), a web-research pivot found Tomeu Vizoso's open-source "Rocket" driver (mainline Linux 6.18 / Mesa 25.3), which reverse-engineered this NPU using NVIDIA's NVDLA documentation — "the programming model matches very closely to NVDLA." That turned black-box guessing into datasheet engineering. From rocket's registers.xml + rkt_regcmd.c + the NVDLA SDP docs, the whole output stage decoded: 0x4070 = EW_CFG (element-wise mode; EW_OP_TYPE bit2 = mul, EW_OP_SRC bit6 = operand-from-RDMA), 0x4074/0x4078 = EW_CVT (operand convert), 0x4084/0x4088/0x4080 = OUT_CVT scale/shift/offset, 0x40c0 = SURFACE_ADD, and the 0x50xx block = DPU_RDMA (0x5018 SRC_BASE, 0x5038 EW_BASE, 0x5034 ERDMA_CFG, 0x5008 OP_ENABLE, dims/strides). The gain decode matched the earlier multi-capture correlation exactly.

Two approaches — the second is the clean one:

• Fused into the matmul output stage (dead end for ork's program). RKNN fuses the mul into the up-projection's flying-mode conv output stage. Replicating that onto ork's synth_i8 wedges: the SDP element-wise, in that mode, expects RKNN's tiled conv-output geometry, but ork emits dense [M][N] — a cube-vs-dense format mismatch the RDMA operand can't bridge. (A register-constant operand does work fused — proving the SDP mul stage itself is fine on ork's dense geometry — but that only covers a scalar/8-value operand, not per-element silu.)
• Standalone SDP element-wise layer (works). The NVDLA guide: element-wise is typically a standalone SDP layer reading both operands from memory, not fused. Captured a pure ONNX Mul(a,b) → a clean 69-register op, enable=0x18, no conv lane (reads a via SRDMA 0x5018, b via ERDMA 0x5038, writes a·b to 0x4020). Templated it (REGCMD_MUL) and ork_npu_probe_i8_mul submits it with the three addresses patched. It executes on the board cleanly — no wedge, 27 µs.

Validated bit-exact. The full formula, confirmed on silicon:

out_i8 = clamp_i8( round( (a + za)·(b + zb)·gain + zo ) )
   za = BS_ALU_OPERAND (0x4044)     zb = EW_CVT_OFFSET (0x4074)     zo = OUT_CVT_OFFSET (0x4080)
   gain = OUT_CVT_SCALE (0x4084, SIGNED 16-bit) / 2^OUT_CVT_SHIFT (0x4088)

Checked against the baked capture: (64+12)(64+10)·0x7aa9/2²¹ + 21 = 105 = captured output. With the zero-points nulled and gain = 1 (mult=0x4000, shift=14): out = clamp(a·b) exact — 8·1=8, 16·1=16, 7·9=63, 10·8=80, 10·12=120, all bit-consistent (uniform-in → uniform-out). Two gotchas recorded: OUT_CVT_SCALE is signed 16-bit (0x8000 = gain −1, use ≤0x7fff); and rapid back-to-back sudo submits give spurious 0s (single clean runs are exact). The multiply semantics also matched an independent golden-tensor analysis: out = clamp(round(R·G·gain)), gain = s_up·s_silu/s_out (100 % within ±1 ulp).

For SwiGLU: ork's two int8-out matmuls write up and silu(gate) to memory, then this standalone SDP op multiplies them on-NPU (a=up, b=silu, gain = s_up·s_silu/s_out) — no conv-geometry coupling.

Remaining (bounded integration, not RE): (1) lay out non-uniform up/silu in the op's NVDLA cube order (uniform validated the math; a single ramp probe fixes the permutation, and ork controls all three buffers so a consistent permutation suffices); (2) wire ork's SwiGLU + validate out=clamp(round(up·silu·gain)) vs CPU + ggml-ork. Tooling: build_mul.py (rknn-toolkit2 on 10.3.0.239), run_rknn_multi (multi-input runner), the RKDUMP_WORDS full-buffer interposer, ewmul_probe mul.

Also fixed this session: the version-bump CI wedge (a stale-checkout commit had reverted the header below the latest tag → every push collided on tag already exists; now the bump anchors on max(header, latest tag)). ork-driver main 59ccc56.

Commit: ork-driver main 27a39ee · src/npu.c (set_i8_out8, set_i8_silu, ork_ppu_fuse_enabled, ork_npu_probe_i8_out8, ork_npu_probe_i8_silu[_cfg]) · src/regcmd_silu.h · examples/test_mm_i8_out8.c (in make test) · grid scripts build_silu_grid.py · captures ppu_captures/

Drove the PPU output-stage fusion from RE all the way to a landed, tested primitive. Gated on the detected SoC (ork_ppu_fuse_enabled returns 1 only on rk3588 via device-tree) — off elsewhere the CPU/NEON requant+activation path stays the default fallback (portability + known-good). Two results:

(1) int8-output requantize — LANDED + validated bit-exact. The matmul emits int8 = clamp(round_half_to_even(acc·mult >> shift)) in its output stage instead of int32 (6-reg delta to REGCMD_I8, RE'd from an rknn_matmul INT8→INT32 vs INT8→INT8 capture diff). make test runs test_mm_i8_out8: mism=0 across M=1/32, K=512..4096, N..4096, identity + scaled requant. The rounding is banker's (half-to-even), verified on silicon. This is the on-chip replacement for the CPU/NEON f32→int8 pack between chained matmuls — the key to keeping a whole FFN NPU-resident (the chaining synergy).

(2) Fused SiLU — mechanism fully decoded + validated (RE scaffolding, gated). Built a Linear+SiLU model grid across a weight-scale octave sweep, captured each under the interposer, and cracked the complete pipeline against silicon (a composite model reproduces the captured curve to ~2 int8):

ii = acc·R          (R = mult/2^shift, a SINGLE scale knob: sets both index AND output gain)
idx = ii >> 6        (6 low bits = PWL interpolation weight)
V16 = LUT[idx]*(1-frac) + LUT[idx+1]*frac    (FIXED silu·S curve, 2 banks, scale-independent)
out_i8 = clamp( R·V16 + bias )               (bias = 0x4080, asymmetric zero-point)

Key findings: the SiLU LUT is scale-independent (one fixed curve, md5-identical across the grid); the index is Q6 fixed-point PWL (the >>6 was the last structural piece); it's a two-knob system — R = 0x4084/0x4088 = f(out_scale) sets the output gain, cfg4068 sets the index's in_scale part (index_per_acc ∝ R·cfg4068), and the LUT is indexed by the int8-quantized pre-activation (a 3-scale model: in/preact/out). The parametric-RKNN-dump + post-requant-agnostic staircase/const-LUT probes cracked the mechanism (blind register search was 2⁹⁶ and wedged the NPU).

The program mismatch (and how it was bypassed): ork's matmul is a 108-reg program (REGCMD_I8) while RKNN's fused Linear+SiLU op is a 126-reg program. So RKNN's captured LUT + registers don't transfer (replaying RKNN's exact tensors+registers on ork gives mean err 109/127), RKNN's output can't be a calibration reference, and a CPU-computed reference can't be made bit-exact from the black-box arithmetic. The captures gave the mechanism; a working fused SiLU had to be built in ork's own frame — which it now is.

★ ork-NATIVE FUSED SiLU WORKS (tools/silu_native.c, commit 3693673): correct silu on the NPU, ~1 int8. Since ork controls both the LUT (the probe's lut override) and the output-stage registers, a 3-pass construction yields correct silu with zero RKNN dependence: (1) measure ork's index(acc) via a ramp LUT (with R = mult/2^shift chosen so a reachable acc range spans silu's transition band); (2) build ork's own silu LUT at those indices — LUT[idx(acc)] = silu(acc·in_scale)/out_scale / R; (3) validate. index(acc) depends only on (R, cfg4068) — measured once, LUT regenerated per (in_scale, out_scale). Validated mean|err| ~1 int8 (max 3–4), general across in 2.5e−4..1.5e−3 / out 0.04..0.20 — near the int8 rounding floor. The fused-SiLU generation problem is solved.

Packaged as a reusable API (ork_mm_silu_build_lut, commit cf3efb1): given (in_scale, out_scale) + a register config it measures ork's index(acc) (one calibration submit) and builds ork's own silu LUT — the caller runs matmuls via ork_npu_probe_i8_silu_cfg(.., lut, ..). examples/test_silu_native.c (in make test) builds the LUT then runs a fused matmul and asserts the output matches correct mathematical silu: mean|err| 0.4–1.2 int8, max ≤4 across several scales.

Matmul-level benchmark (tools/silu_bench.c, commit 4805dc5; RK3588, warm, LUT amortized):

SiLU is free in the output stage — the fused submit equals a plain matmul submit within noise — and it eliminates the CPU-silu pass (5 µs at M=1 up to ~158 µs at M=32/16k-out) plus the dequant/requant roundtrip. The win scales with M (largest in prefill), and it's the enabler for keeping a whole FFN NPU-resident (no CPU handoff between chain links).

The int8-output stage + gate are the shipped production win; ork-native fused SiLU produces correct silu (~1 int8), is a tested library call (ork_mm_silu_build_lut), and is free at the matmul level. Remaining: register auto-pick → ork_mm_pack_i8_silu (cache LUT+regs with a resident weight) + a run path, then EW-mul (SwiGLU dual-input), chain integration, ggml-ork wiring, and an end-to-end llama-bench. Full RE narrative + register tables + probe methods: PPU_FUSED_ACTIVATION_WIP.md.

Files: PPU_FUSED_ACTIVATION_WIP.md · scratch build_pair.py / run_rknn.c · interposer tools/regcmd_capture.c · captures /tmp/{none,relu,silu}_cap.log (board)

The standalone PPU-LUT replay was closed (Phase 1B: sigmoid replay wrote 0/6144 words — the PPU only comes live as a matmul/conv output stage, not standalone). This reopens the viable path and decodes it. Built rknn-toolkit2 2.3.2 (in the rkllm-converter x86_64 container on .239) and build_pair.py — one seeded conv weight, exported 3× as none / +ReLU / +SiLU so the only cross-model difference is the activation (killing the earlier random-weight diff confound). All three run on the NPU (rknn_run ret=0), and the output buffer proves correct compute: bare-conv int8 out 0f 12 17 0e… → SiLU out ae ad ac a4… (the SiLU-transformed values). Two distinct fusion mechanisms fall out of the regcmd diff:

• ReLU / clamp = FREE — 5 output-stage register writes, no LUT, no extra program (same 126+122+26-reg structure as bare conv, enable=0x1d). Key reg: 0x4080 → 0xff80 (= −128 = int8 zero-point = clamp-min = the ReLU), plus requant 0x4060/0x4084/0x4088 and 0x1184.
• SiLU / GELU / sigmoid = LUT path — enable=0x18 + a 1097-reg program. 0x4100 = LUT index/bank port; 0x4104 = LUT data port, streamed with the quantized activation curve (captured a monotonic SiLU ramp ff9d, ff9c, ff9b… 7fff); 0x4108/0x410c/0x50xx = base/step/PWL config; same requant output-stage regs. LUT values are per-output-scale (must be regenerated for each gate matmul's scale). regcmd word-pair encoding: <val16><regoff16> <op16><mask16> (mask flips 0000→ffff for a signed/negative value).

Only SiLU (pointwise) is fusable this way — RMSNorm/softmax are reductions and stay on CPU/NEON (same as librkllmrt). This targets the M=1 decode CPU-handoff bubble (gate→SiLU→down): the RE is done; what remains is implementation in src/npu.c — emit enable=0x18 + the LUT stream + requant on the FFN gate matmul, generate the 256-entry SiLU LUT per output scale, and validate bit-exact vs CPU SiLU. Bounded win (one activation handoff per FFN; RMSNorm/softmax residual), not yet wired end-to-end.

Commits: ork-driver bed4ccd · llama.cpp-rockchip master 4f244e5c2 · tool tools/domain_probe.c · docs AGENTS.md

The rk_iommu window is 32-bit (~4 GiB) per iommu_domain_id, so a >4 GiB model resides across several domains and each submit zero-copy-swaps the active domain (dom_activate parks/restores per-domain scratch; domain_probe confirms 8.32 GB across 2 domains, distinct IOVA). The bug was residence order: filling one domain to its ~4.16 GiB EFAULT soft-resets + corrupts the IOMMU, so the next domain dies on its first tile. Fix: ork_weight_domain() byte-balances and advances BEFORE the EFAULT (3.0 GiB cap, below the measured ~4.16 GiB point) — no EFAULT, no soft-reset, clean hand-off. Plus: RAM-ceiling residency budget (no eviction/churn), auto n_domains, all-core pack/tiling, and a .orkpack version token (= library MAJOR) with zero-copy import load. (The ork_stream_pool was the wrong mechanism — it only matters beyond ~64 GiB; left off. The earlier "domains share 4 GiB" was a 62-commit-behind Frankenstein checkout, not reality.)

7B resides 10.38 GiB across 4 domains, churn≈0, 0 soft-resets. Correct, not fast-but-wrong: ORK PPL 7.84 (in the ~7–8 band; broken prefill would be NaN/>50) + make test quant-RMS + domain_correct bit-exact. Open: churn=3 (target 0); close the 7B-prefill gap to rkllm (partly CPU-bound). Full writeup → Multi-Domain Residence.

Commit: ork-driver main (6775334) · tool tools/sparse_check.c · docs AGENTS.md "Weight-DMA amortization"

Characterized the residual single-core "per-row" gap from first principles: it's weight-DMA-bound, not row/compute-bound. mc_prof parametric sweeps show µs/submit is flat vs rows (mc=4→493, mc=31→516) but ∝ N (228/510/914 at N=512/2048/4096 ≈ 130µs + 0.19µs/N-col) — i.e. each M-tile submit re-streams the whole K×N weight from DRAM (~11 GB/s) and that dominates. So the lever is rows per weight-stream. The M-tile was capped at R-1 = pow2_floor(2*cbuf/K)-1 (~31), a disproven "CBUF-resident rows" RE finding — activations stream, and 0x1010 is a neutral hint. The real bit-exact ceiling is the 0x1040 K-schedule, mg_max*64 (full CPU-ref: mc=mg_max*64 correct / mc+1 miscomputes at every K — 704@K512, 320@K1024, 128@K2048, 64@K3584/4096).

Default-on, no knob; make test fully green. Corrects the prior "tile size is throughput-neutral" reading (those tests never broke past mc=31). Open: end-to-end prefill (partly CPU-bound), the wide-K K>4096 K-slice path (own caps), fp16 (latent large-tile bug). Full writeup → Weight-DMA Amortization.

Commit: ork-driver main (a95a655 ORK_DECODE_MC knob in budget())

The decode companion to the entry below. Per-core sampling of the clean, resident llama-server decode (churn 0) on Qwen2.5-7B-Q8_0 caught the resident decode running single-core — budget() hard-returned 1 core at M==1, so the big FFN projections (N/K=18944, which dominate the decode wall) ran on Core0 alone while Core1/2 sat idle. (The earlier "3-core 26%" reading was the churny llama-bench path masking it.) Gating that short-circuit behind ORK_DECODE_MC=1 lets M=1 split N-tiles across all 3 cores:

Bit-exact vs CPU ref (test_matmul/decode/layer all ALL OK); default-off (small-N decode matmuls lose the multi-core barrier to the single-core dispatch floor). run_rr_chain (no-barrier whole-matmul-per-core round-robin) scoped and rejected: for the FFN-dominated decode the barrier N-split already puts all 3 cores on each big matmul (≈T/3); whole-matmul-per-core would use only the 2–3 independent matmuls/group (QKV=3, gate/up=2) at full T each — strictly worse. Verdict: ✅ real 1.62× banked; the residual 1.49→4.09 is the same synchronous-execution wall (NPU 33% / cpu7 62%, vs rkllm's NPU 91% / CPU 20%) — cross-node async in ggml-sched, upstream-scale. ⏭️ NPU-bound decode parity needs M>1 batched/speculative verify, not better M=1. Full writeup → Decode Multicore and the Sync Wall.

Commits: ork-driver main (db46800 tools/overlap_probe.c diagnostic) · ggml-ork fork master fc0d618c8 (recovered + committed working multi-domain resident build)

Supersedes the next entry's "NEXT: regcmd-RE the per-core kernel" conclusion. Per-core NPU+CPU utilization sampling (/sys/kernel/debug/rknpu/load + /proc/stat, 150 ms) on the recovered, committed, trustworthy build (pp512 45.3 / pp256 40.1, 10 GiB resident) vs a fresh rkllm reference (rk_bench, librkllmrt 1.2.3: prefill 68.9–72.1, decode 4.06):

Both use all 3 cores at max clock — the matmul/multi-core/governor are not the lever. The gap is per-core saturation (ork ~40% vs rkllm ~65%, ≈ the 1.52× throughput gap): ork's NPU idles ~60% between matmuls. Proven by mc_prof back-to-back (no graph) = 77% vs live = 40% — the delta is per-matmul CPU work the NPU can't overlap. Root cause in code: ggml-ork is a synchronous BLAS-style backend (quant→run_i8 blocking→dequant, nodes one at a time); NPU idles during quant/dequant, the serial run_multicore setup+copy on cpu7, and the between-matmul graph ops. New overlap_probe diagnostic bounds the within-backend overlap ceiling at ~1.15× (the bigger overlap needs the next node's input = the same async wall). Easy wins already in code (NEON+OMP quant/dequant, chained submits, 3-core); OMP big-core pinning = no win (45.05 vs 45.72). Verdict: prefill banked at trustworthy 40.1 (0.56× rkllm) with a fully-mapped, async-gated ceiling; the real lever is cross-node async execution in ggml-sched (architecturally hard, upstream-scale), so higher-leverage fronts are decode (0.22×) and the distributed/speculative path that puts the NPU on separate M>1 work. Full writeup → Prefill NPU Utilization and the Sync Wall.

⚠️ Superseded (2026-06-30): the "matmul is the lever → regcmd-RE the per-core kernel next" conclusion came from a per-matmul comparison; direct per-core utilization sampling shows the matmul saturates fine (77% back-to-back) and the gap is the synchronous execution model — see the entry above. The 31→40 lever numbers were also measured on uncommitted dirty builds; the recovered, committed baseline is 40.1.

Commits: ork-driver main (af5bf78 per-N-slice chain prefill + aceffb5 test_sn3; cb99c1c multi-domain Af fix) · ggml-ork fork branch (lever D attn_output routing, B′ NEON act-quant, C overlap gated)

Lever-by-lever grind on Qwen2.5-7B-Q8_0 prefill, each measured + bit-exact-gated vs librkllmrt (pp256 73.89 / decode 4.07). Net pp256 ~31 → ~40 (54% of rkllm). Only two levers were big: (1) per-N-slice chain prefill — chain M-tiles within one N-slice buffer so wide-N gate/up stop going per-tile → ioctls/matmul 55→5, pp256 30.94→36.11 (+16.7%); (D) route attn_output to NPU — a strstr("output") guard was over-matching attn_output.weight and forcing 28 [3584,3584] matmuls to CPU → 35.79→40.2 (+12.6%), one line. Three more measured null/neutral/tiny (B′ NEON chain-quant null; A fuse dequant→requant a structural no-op — 0 matmul→matmul edges, every matmul fed by a norm; C overlap ~3%, gated). Also fixed two push-blocker wedges (multi-domain SIGSEGV; chain-prefill Sn>1 errno-110) and repaired the build/make test suite.

The diagnostic (rknn_vs_ork, ork vs the closed RKNN matmul API): ork-3core is 1.27–1.31× of Rockchip's closed kernel on q/o + gate/up (fair layout — RKNN's "native" layout was 1.28× slower). So the matmul is competitive; the 40→74 gap is host-side integration/engine, not a matmul wall. Nuance: ork-1core is 3.5×/core slower than RKNN-1core (we use 3 cores to reach 1.3×) → the per-core kernel has real (hard) headroom. Verdict: ✅ +27% prefill banked; activation/dequant/quant levers spent, matmul competitive. NEXT: the per-core matmul kernel via regcmd RE (close the 3.5×/core gap) — a both-regimes win (prefill + decode/M=1 + future M>1 verify). Full writeup → Prefill Parity Grind.

Commit: ork-driver 33905ff (branch re/rkllm-mtile-decode, local) · tools/autotune.c · ORK_KTILE knob

Built a per-shape NPU matmul tiling autotuner (sweep cores / M-N-subtile / K-tile, bit-exact-gate each, time warm run_i8, cache the winner) and ran it on the Qwen2.5-7B prefill shapes. It found a real matmul-level lever — CBUF K-tiling lifts GOPS +36–60% (KT=1024 raises the row-ceiling R 16→64: q/o_proj 700→1118, gate_up 731→994) — but end-to-end pp256 was flat (11.96→11.46, ~16% of rkllm 73.6). Decisive: a +36–60% matmul speedup moved prefill 0% → prefill is host-bound (ggml-ork CPU quant/dequant/scheduling), not NPU-matmul-bound.

Verdict: ⚖️ Autotuner works + K-tiling is a genuine matmul win, but it doesn't move prefill — the 12→73.6 gap is host/ggml-ork-side, not the NPU matmul (which is already adequate; existing knobs saturated, core-count is the dominant lever and already default). Closes the NPU-side prefill arc (submit-floor → quant → tile-size → K-grouping → multi-task-M-offset[doesn't exist] → autotuner). ORK_KTILE gated-off; the matmul-GOPS win is banked for the M>1 batched-verify (spec-decode) regime where the matmul binds. Autotuner + K-tiling = durable reusable assets. NEXT: pivot to spec-decode. Full writeup → Tiling Autotuner + K-Tiling.

Commit: ork-driver 8bc4004 (tools/domain_probe.c + ORK_IOMMU_DOMAIN knob; ork_npu_set_pack_domain API)

The RK3588 NPU's 32-bit IOVA limit is per iommu_domain_id, not per device — the rknpu driver exposes up to 8 domains, so ~32 GiB of weights can be NPU-resident at once (ceiling = physical RAM, not IOVA). Reachable from a field already in rknpu_ioctl.h (rknpu_mem_create.iommu_domain_id + rknpu_submit.iommu_domain_id), no kernel patch, no librknnrt, bit-exact. Overturns the "only ~4 GiB DMA-mappable → must stream" conclusion. Found via: arbiter (librkllmrt beats CPU on a streamed 7.65 GiB w8a8 7B — prefill 73.6 vs 15.7 = 4.7×, decode 4.05 vs 3.42 → big-model-on-NPU is implementation, not hardware) → dump (librkllmrt does 8 MEM_CREATE total, zero per-token churn, residing 6.6 GiB across 2 domains; our stack did ~100–170 MEM_CREATE/token — the gap) → de-risk (confirmed from our uABI: 2 domains = 8.46 GB, 8 = 32.28 GB, domain-1 matmuls bit-exact).

Verdict: ✅ Multi-domain residence escapes the 4 GiB cap for any model ≤ physical RAM — no streaming, zero churn, open uABI. Then beatable via int4 storage (~2× decode), SRAM tier (vendor uses none), async dispatch, submit fusion. NEXT: end-to-end 7B 2-domain residence (target librkllmrt 73.6/4.05) + 27B scale test (RAM-bound now). Full writeup → Multi-Domain Residence.

Commit: ggml-ork d2e8526e8 (board, ORK_STREAM_POOL=1, off by default; submodule → ork-driver a443b64)

Wired the RAM-int8 stream pool into ggml-ork's streaming LRU and measured end-to-end. Bit-exact, zero regression at the resident budget — but decode does not convert: forced-stream 1.7B decode (-n 8, IOVA 1536 MB, RAM 22 GB) is identical pool-on vs pool-off (0.03 t/s; resident ceiling 3.85). ORK_PROFILE: NPU compute is ~0.6% of the decode wall; ~99% is per-token IOVA map/unmap churn (~100–170 weights through the window every token, each MEM_CREATE+MEM_DESTROY). The pool removes inflate+copy, but that's not the bottleneck — the 4 GiB cap as churn is, and the pool doesn't reduce it.

Verdict (banked): ⚖️ NPU-streamed decode is churn-bound → big-model decode runs on CPU (native Q4_K, fused, no churn); the NPU's lane is resident-hot + M>1 prefill/verify. The pool's open upside (M>1 prefill streaming — import removes the copy, prefetch hides the inflate at M ≳ 256) still needs a real >4 GiB dense transformer to demonstrate (the board has none — only Q4 MoE 27B + hybrid-SSM 9B). Detail folded into → Prefetch-Overlap + Stream Pool.

Commit: ork-driver a443b64 (v0.5.0) · probe tools/stream_prefetch_probe.c · API ork_stage_* / ork_stream_pool_* / ork_mm_load_i4a8_import

Measured the int4-streaming swap (FILL int4→int8 inflate+tile into a non-mapped staging dma-buf → MAP MEM_CREATE import → SUBMIT, with a prefetch thread running FILL(N+1) behind SUBMIT(N)). MAP is cheap (48–169 µs); MEM_DESTROY UNMAP is expensive (~0.5–2 ms) — the irreducible per-swap killer. A transient unmap-every-swap ring therefore LOSES (overlap only ~1.0–1.17×, never reaches t_submit; the fill hides only at crossover M ≈ 256). But this validates the RAM-resident-cache design: since unmap is the dear op, cache the inflated int8 in RAM, keep hot entries mapped, and unmap only on eviction — a hit pays just the cheap map and never re-inflates. Realizes the routing-skew hit-rate lever in RAM budget (8–15 GiB) vs the 1.2 GiB usable IOVA that capped the old cache.

Built ork_stream_pool_* (RAM-resident inflated-int8 cache; add_i8/add_i4a8; cheap map/run/unmap; entries persist across unmap; LRU policy left to the caller) + ork_mm_load_i4a8_import. All bit-exact vs CPU ref, make test green.

Verdict: ✅ RAM-resident inflated-int8 cache with cheap map-on-hit + evict-time unmap is the right streaming structure — and the measurement is what ruled out the simpler unmap-every-swap ring. ⚠️ Wins M>1 prefill/verify (fill hides) + warm hits (cheap map, no re-inflate); not M=1 decode (un-hidden + bandwidth-bound), not the 4 GiB cap. NEXT: Step 2 — retarget ggml-ork's hot-expert/wcache LRU to RAM-int8 + map-on-hit via the pool, keeping hot entries mapped.

Full writeup → Prefetch-Overlap + Stream Pool

Commit: ork-driver 57feca4 (v0.4.0) · API ork_dma_import/_sync/_free + ork_mm_load_i8_import · uABI src/rknpu_ioctl.h (DRM PRIME fd↔handle, dma-heap alloc, dma_buf_sync) · test tools/zerocopy_import_test.c

The closed runtime's rknn_create_mem_from_fd zero-copy import turned out to be not a new ioctl — an LD_PRELOAD strace showed it's an unused input mode of MEM_CREATE: PRIME_FD_TO_HANDLE (generic DRM, dma-buf fd → GEM handle) then MEM_CREATE{handle≠0, flags=0, size=0} ⇒ the kernel IOMMU-maps the existing pages, no alloc/copy. Built it: a dma-heap buffer of pre-tiled weights imported as a matmul tensor is read directly by the NPU, bit-exact vs alloc+copy, make test green. ★ Coherency bug found+fixed: rknpu's MEM_SYNC does NOT flush a foreign imported cacheable dma-buf (only its own buffers) → all-wrong output on the cacheable system heap; fix = the dma-buf's own DMA_BUF_IOCTL_SYNC (START|WRITE / END|WRITE) around the CPU fill → correct and fast. A fixed-uniform-size-slot pool cycled a 4295 MB working set / 1024 swaps through one slot → 0 fails, 0 mismatches — the earlier streaming wall was variable-size MEM_CREATE/DESTROY churn, not the swap.

Verdict: ✅ Zero-copy import + uniform fixed-slot streaming, on the existing uABI — load 2.2–4.4×, swap 2.1×, crash-free past 4 GiB churn. ⚠️ Does NOT escape the 4 GiB IOVA cap (imported dma_addr still 32-bit) or change M=1 decode bandwidth — the win is the eliminated copy in the M>1 streaming swap/load path. NEXT: wire ork_mm_load_i8_import into the ggml-ork streaming path (27B/35B verify / MoE) for the end-to-end number.

Full writeup → Zero-Copy Weight Import + Fixed-Slot Streaming

Commits: ggml-ork edf5e855b (M1b hot-expert partition) · ggml-ork a72692188 (cold-path vec_dot reuse) · ork-driver dd6e7e6 (v0.3.5) · probe tools/moe_m1_probe.c · knobs ORK_MOE_HOT=<K> / ORK_MOE_HOT_GIB / ORK_PROFILE=1

Wired the decode partition (M1b) into ggml-ork — backbone pinned + a static hot-expert cache (top-K/tensor, LRU under a byte budget) resident on the NPU + cold experts on CPU — which also fixes the broken live MoE path (the wedge was over-residency hitting the 4 GiB IOVA cap, MEM_CREATE EFAULT, not a bad regcmd). Coherent vs CPU ref, peak 1.8–2.0 GiB (<4 GiB), no wedge. Cold path now reuses ggml's vec_dot (asimddp/SDOT) instead of a scalar+f32-dequant loop. ★ Then adversarially verified a wrong claim that the partition was bottlenecked by a "~21 ms/expert NPU submit floor" (→ "MoE-on-NPU decode is dead"). It isn't: isolated resident per-expert M=1 submit = run_stream 54–67 µs, which BEATS CPU 73–214 µs (real Qwen3.6-35B-A3B dims 2048×512 / 512×2048). The 21 ms = residency thrash (MEM_CREATE/DESTROY churn at ~25–30 % hit serializing on the single-stream NPU) + live-pack (3–5 ms) + dead-code profiling (the moe_pack accumulators were never written → read as "repack ~0%").

Verdict (updated 2026-06-27 — M2 end-to-end): the resident GEMM beats CPU at the micro level, but the end-to-end M2 decode LOSES on RK3588: 6.56 t/s vs CPU 19.09 (~2.9× slower). Walls are RK3588-specific: usable IOVA only ~1.2 GiB (GGUF backbone eats ~2.8 of 4 → max ~17 % hit), the cold-expert per-GEMV is memory-bandwidth-bound (loses to ggml's fused batched MUL_MAT_ID), and ork_mm_load_i8 is 5.9 ms/load (Bf rebuild). The lever is M>1 (batched verify/prefill), not residency — there load+submit amortize and cold becomes compute-bound. ⚖️ Kept as an experimental, default-OFF setting (ORK_MOE_NPU), wired into oRKLLM with an RK3588 "not recommended" disclaimer; revisit on wider-IOVA / DDR5 hardware or the M>1 path.

Full writeup → MoE Decode Partition + the "21 ms submit floor" that wasn't

Tool/data: llama-moe-skew (~/llama.cpp/examples/moe-skew) · /tmp/skew_{lfm2,qwen36}.csv (box 10.3.0.241, CPU-only, no NPU)

The 4 GiB IOVA wall makes resident-all-experts impossible, so the M3 hot-expert cache only works if routing is skewed. Measured it via the contiguous ffn_moe_argsort tensor per layer (⚠️ GOTCHA: the ffn_moe_topk view is non-contiguous → reading it gives a FALSE "uniform" result). Profiled on the UD-Q4 GGUF (routing skew is precision-invariant + UD/k-quant keeps the router gate high-precision → q4 routing ≈ f16; an earlier f16-convert path was wasted effort). Both models SKEWED (Gini 0=uniform): LFM2-8B-A1B (32 exp, top-4) 0.44; Qwen3.6-35B-A3B (256 exp, top-8, the deploy target) 0.65. Skew RISES with layer depth (deeper/majority layers specialize → good for caching). Per-expert int8 only 3.15 MB on the 35B.

Verdict: ✅ Hot-expert cache VIABLE. Operating point ~3–3.5 GB cache (top ~26–28/256, ~10–11%/layer) → ~50% of decode expert lookups NPU-resident (alongside the ~0.5–1 GB pinned backbone, within the 4 GiB window); cold ~50% → CPU/GPU; >50% hit (≥6 GB) exceeds the window → stream the tail. Policy = per-LAYER top-K (global pin underperforms). Confirms the Branch-A decode partition.

Full writeup → MoE Routing Skew — sizing the hot-expert cache

Commits: ggml-ork c89cf8654 (synthetic per-expert key + persist/load) · ork-driver 6b0d568 (ork_mm_load_i8 full-K Bf) (RK3588 10.3.0.236 validate / box 10.3.0.241 convert, governors pinned)

Implemented the sparse-MoE expert persist/load path so a mixed .orkpack captures the routed _exps (MUL_MAT_ID) tensors (before: backbone + router + dense shared-expert persisted, ZERO routed experts). ggml-ork c89cf8654 gives each expert slice a synthetic key <name>#<e>, persists ALL n_expert slices via ork_orkpack_tier at convert, and on load calls ork_mm_load_i8/ork_mm_load_i4a8 (read pre-tiled int8) instead of dequant+repack. ork-driver 6b0d568 makes ork_mm_load_i8 rebuild the full-K Bf (gated K%512==0 && K<=4096) so disk-loaded int8 experts are first-class chain inputs. Result (LFM2.5-8B-A1B): .orkpack now has 2112 _exps entries (264 × 8) vs 0 before; reload-from-disk NPU-expert output bit-identical to a CPU int8 reference (889 experts loaded from disk on a deterministic prompt). Gate = reload-vs-CPU-ref because the live NPU MoE path (mode 0) is itself broken on this model (pre-existing, orthogonal).

★ KEY FINDING: running real experts on the NPU surfaced the 4 GiB IOVA cap for real — an 8B MoE's int8 experts exceed 4 GiB, so compute status -1 once ~1909 distinct experts accumulate resident (tokens correct up to the cut). NOT a persist bug — it's the wall that mandates streaming (M2) + a hot-expert cache (M3); resident-all-experts is impossible beyond a tiny model.

Verdict: ✅ M1a DONE — routed experts persist and reload bit-identically; ★ the same run confirmed the 4 GiB IOVA wall bites real 8B MoE experts.

Full writeup → MoE Expert Persist/Load + the 4 GiB IOVA Wall

Tools: tools/moe_expert_probe.c (1840abb) · src/npu.c (ORK_DIRECT_I4 570cb68; chan_scales_f32 NULL-guard 980cea9) (RK3588, governors pinned)

Milestone M0 of the MoE-on-NPU saturation build, isolating just the per-expert FFN GEMM (the prior MoE Expert Offload ~7×-under-CPU was an integration number walled by the CPU-side Q4_K→int8 repack, not the NPU). With the new primitives — direct int4→int8 tiled inflate (ORK_DIRECT_I4, skips the f32 round-trip) + cacheable 0x403 fill — the NPU beats CPU at every Qwen3-A3B expert shape (gate/up K2048×N768, down K768×N2048, + 4096/1408) and every per-expert M∈{8..128}: 3–10× at M=8, no crossover, bit-exact vs the CPU int8 reference. The direct inflate cut per-expert weight materialization 16–30× → cost collapses to the GEMM itself. Dispatch: run_i8 (single submit) is already multi-core (auto-tuner spreads N-tiles over 3 cores); run_chain_i8 is single-core (slower/expert) → use run_stream_i8 for chainable gate/up (K%512==0, K≤4096), run_i8 for down. Caveat: CPU baseline was single-thread NEON; a 4-thread CPU GEMM narrows the M=8 margin to ~1.5–2.5× — NPU still wins and frees all 4 CPU cores (the saturation goal).

Verdict: ✅ VIABLE — flips the matmul-level go/no-go the integrated ~7× couldn't. Follow-on (M1a): the mixed-.orkpack convert does NOT yet persist routed _exps/MUL_MAT_ID expert tensors (the backbone incl. router + dense shared-expert does).

Full writeup → MoE Expert FFN GEMM on the NPU — M0 viability

Tools: tools/prefetch_headroom.c (f959e49) · tools/disk_stream_bench.c (f3ef81a) · tools/dmabuf_fill_probe.c (c70f71d) · src/npu.c (ork_dma_alloc flags b53c59b; ORK_DIRECT_I4 570cb68) (RK3588, governors pinned)

Streamed-prefill prep (int4 inflate + tile into the resident DMA buffer) is 6–26× the matmul at M≤128 and TILE/transfer-bound — and since t_tile is M-independent while t_npu ∝ M, the prefetch-viable crossover is large-M (prefill), not decode. The fill was pinned at a flat ~1 GB/s — but that was an artifact of benching ork_dma_alloc's write-combine zero-copy flag (0x401), not slow RAM: the production weight tiles already use cacheable 0x403. Benched correctly, cacheable + clean-before bsync = ~10 GB/s (8.3–8.8×), NPU read correct (cached memcpy ~12 GB/s; NEON stnp doesn't help the WC path — the uncached write is the wall). A second lever, a direct int4→int8 tiled inflate (ORK_DIRECT_I4) skipping the f32 round-trip, adds 2.2–2.7× UNIFORM / 3.6–4.9× NF4, bit-identical. disk/mmap-warm == RAM (page cache hides it, 0.5% delta); cold disk only matters when the model exceeds RAM.

Verdict: ✅ The "1 GB/s wall" was a benchmarking artifact; production cacheable tiles run ~10 GB/s, and the direct inflate makes the nibble inflate near-free. Levers compound (cacheable 8× × direct inflate 2–5×). NEXT (identified): clean-only bsync(TO_DEVICE) steady-state re-fill, and an on-chip SRAM/NBUF tier (RKNPU_MEM_TRY_ALLOC_SRAM).

Full writeup → dma_buf Fill Bandwidth

Tools: src/npu.c (ork_mm_pack_i4a8/_im, ork_w_dump_i4a8/ork_mm_load_i4a8, ORK_NF4/ORK_SR) · tools/gguf_tier_map.c (RK3588, governors pinned). Commits dc782b2, 93364b7, d656955, 378045d, 2c38a27+f78bc04.

Three accuracy levers on the int4-stored/int8-computed path: an NF4 codebook (bitsandbytes' 16 non-uniform levels, NEON vqtbl1q_s8 inflate via an int8 LUT, bscale=absmax/127; reconstruction err on Gaussian weights 0.1102 < uniform-int4 0.1581, RMS 0.873%); stochastic rounding into the bracketing NF4 levels (ORK_SR; a dot-product/PPL effect, raises per-weight RMS while removing accumulation bias); and an imatrix-weighted per-channel scale (ork_mm_pack_i4a8_im) — imatrix is per-INPUT-channel (K), scale per-OUTPUT-channel (N), so it can't re-weight a level pick; instead a clip-ratio grid search minimizes Σ_k imatrix[k]·(w−dequant)² → 6.4× weighted-error reduction, imatrix=NULL byte-identical to absmax. Plus a compact 'O4N1' persist (~½ the int8 dump, bit-identical reload, d656955) and gguf_tier_map (the allocation oracle: Qwen3-1.7B-Q4_K_M → 140 q4_K→int4 + 28 q6_K→int8, matches the convert split).

Verdict: ✅ Mixed int8 + int4-NF4 from the fp16 source with imatrix scales = 11.33 PPL @ ~58% int8 storage, beating double-quant by −1.15 → validates "fp16 is the source; the mixed GGUF is an allocation oracle only, never re-quant its bytes." Residual gap to Q4_K_M = the int8-ACTIVATION cost (bounded by the all-int8 anchor), not the weights.

Full writeup → NF4 Codebook, imatrix, and Mixed Precision

Tools: src/npu.c (ork_mm_free/ork_w_dump/ork_mm_load_i8/arena) + ggml-ork.cpp (budgeted-LRU wcache, .orkpack, ORK_EVICT_SRC) + tools/dma_probe.c·stream_probe.c·persist_probe.c (RK3588, governors pinned)

Established the NPU IOMMU addresses only ~4 GiB at once (hardware: rk_iommu 32-bit IOVA — dma_probe fails at exactly 4.23 GB in both CREATE-only and CREATE+mmap modes; v2's "40-bit" is physical-PTE reach, not the device IOVA). So a 13.5 GB int4 27B / 9 GB int8 9B cannot be resident. Added streaming (ork_mm_free reclaims IOVA on LRU eviction — stream_probe cycled 26.8 GB through the window, ≤256 MB resident; budgeted-LRU wcache ran a 9B with 0 wall-hits / 352 evicts) and .orkpack persist (self-populating disk cache of pre-tiled weights; persist_probe byte-identical roundtrip, load 8.3× faster than pack). Custom loader (ORK_EVICT_SRC) cut convert peak RAM 30 GB → 5.7 GB.

Verdict: ✅ A model larger than the NPU's address window runs on the NPU (streaming), and the conversion is paid once then skipped (.orkpack) — the foundation for a streamed 27B verify in spec-decode. The 4 GiB cap is permanent; stream the verifier, never resident.

Full writeup → Layer-Streaming and .orkpack Persist

Tools: tools/regcmd_capture.c + tools/parse_trace.py + src/npu.c (RK3588, governors pinned)

Successfully intercepted the complete register command stream of a real 0.6B RKLLM model-forward pass natively on the board, and implemented/validated a zero-overhead single-submit hardware wait-chained monolithic pipeline.

Key Findings:

1. The Secure Preload and 32-Bit Range Mismatches Resolved: Solved ld.so security clearing under sudo node by preloading the interposer from a trusted root-owned directory (/usr/lib/regcmd_capture.so). Fixed a 32-bit sign-extension mismatch in by_dma_range allowing full hexadecimal capture of 11,586 regcmd blocks.
2. The Closed-Source CPU-NPU Ping-Pong Exposed: Decoded traces proved that the proprietary librkllmrt.so runtime ALSO executes non-matmul layers (RMSNorm, Softmax, and RoPE) on the CPU (using NEON SIMD), round-tripping to the NPU only for matmuls!
3. Zero-Overhead wait-chaining (CORE_CTRL = 0x000f): Completely bypassed the broken kernel-level task_number > 1 path by concatenating the [RMSNorm -> MatMul -> RoPE] chain into a single, unified regcmd buffer, writing trigger command CORE_CTRL = 0x000f between each sub-task to pause the hardware parser and wait natively for each step to finish.

Verdict: ✅ Success. The monolithic chained pipeline compiles and runs cleanly on the physical RK3588 NPU, establishing the only contiguous, on-die LLM execution path in the world that bypasses the CPU-NPU barrier.

Full writeup → End-to-End RKLLM Trace Capture and Zero-Overhead wait-chaining

Tools: test_chain_i4_s1.c + examples/test_matmul.c (RK3588, governors pinned)

Investigating errno=110 (Timeout) and errno=22 (Invalid argument) during chained/multi-task submissions in ork_mm_run_chain_i4 and ork_mm_run_chain_i8. Progress, methodologies, and next steps documented.

Full writeup → Hardware Chaining Debug (INT4/INT8)

Tools: examples/i4.c + examples/test_matmul.c + examples/test_speed.c (RK3588, governors pinned)

Resolved regressions involving bounds calculation errors during manual M-dimension slicing in the INT4 emulated scheduler, and physical context state leakage when switching between disjoint and chained submissions in INT8.

Key Findings:

1. Explicit NPU Reset Required: Context leakage during NPU mode switching causes corrupt results. We instituted a self-healing RKNPU_ACT_RESET ioctl intercept when switching to multi-submissions on dirty context handles.
2. Strict Metric Gatekeeping: We augmented the validation suite to strictly catch these artifacts. test_matmul.c actively tests disjoint-to-chain leakage, i4.c exercises exact manual loop bound allocations, and the new test_speed.c enforces absolute speed requirements ($&lt;9000 \mu s$ prefill latency and $\ge 1.4\times$ multi-core scaling) to detect performance regressions immediately.

Verdict: ✅ Success. The fixes definitively resolve the regressions while erecting new testing guardrails that mathematically and performance-wise bound future degradation.

Full writeup → INT4 Scheduler Emulation and INT8 State Leakage

Tools: tools/chain_probe.c + src/npu.c (RK3588, governors pinned)

Programmed the NPU's internal Program Control (PC) sequencer directly via Block 0x0101 trailer registers (0x0010 and 0x0014), bypassing the CPU-to-NPU submission latency floor (~134 µs per submit) and avoiding kernel multi-task limitations.

Key Findings:

1. Direct Sequencer Jumps: Programmed register 0x0101:0x0010 (Next PC Pointer) with the next query's DMA address, and register 0x0101:0x0014 with trigger flag 0x0037 to instruct the hardware sequencer to jump natively from query to query in physical memory.
2. Halt Signal: Marked the terminal task by clearing the Next Pointer and writing 0 to register 0x0014 to stop execution and fire a single completion interrupt back to the host CPU.
3. Core Binding Constraint: Chained tasks must be explicitly pinned to a single physical core (RKNPU_CORE0_MASK) to avoid cache and execution pipeline desynchronization.

Verification Benchmark: Executing chain_probe 8 200 natively on the RK3588 target board yielded an 11.21× speedup, slashing the average latency of executing an expert matrix multiplication from 464.6 µs down to 41.4 µs.

Full writeup → Mixture of Experts (MoE) Sequencer Programming and PC-Chaining Validation

Tools: tools/i4_multim_fuzz.c + examples/test_layouts.c (RK3588, governors pinned)

Conducted an exhaustive userspace register sweep (836 runs, 29 hits) mapping unmapped and inactive register ranges across the CNA, DPU, and PPU blocks.

Key Findings:

1. The Quantization Engine Formula: Mapped the exact PPU fixed-point scale-and-shift quantization math done in hardware before committing results to DRAM: $$\text{Output}{ij} = \left( \text{Accumulator}{ij} \cdot \text{Multiplier} \gg \text{RightShift} \right) + \text{Bias}$$
2. Quantization Control Registers: Identified bias 0x4080, integer scale multiplier 0x4084, and right shift scale 0x4088 (with duplicate/alias 0x301c in DPU) as the primary tuning knobs.
3. Stride & Shape Registers: Identified CNA 0x1014 (activation row stride) and CNA 0x1028 (logical macro-row bounds) as the registers dictating the multi-batch dimension traversal.

Verdict: ✅ Success. The mapped registers enable a direct architectural path to perform 100% hardware-accelerated quantization and dequantization, bypassing CPU loops completely.

Full writeup → Exhaustive Register Fuzzing and Quantization Mapping

Tools: tools/stride_test.c + test_matmul + quant (RK3588, governors pinned)

Successfully reverse-engineered the output row-stride configuration in NPU registers 0x403c and 0x4038, enabling robust Output Zero-Copy (ORK_ZC_OUT=1) support on wide matrices with multiple column slices ($S_n &gt; 1$).

Key Findings:

1. The Row-0 Constraint: Overriding the Row 0 gap settings in register 0x405c causes pipeline mismatches in the CNA/DPU engine and wedges the physical NPU (errno=110).
2. Offset-Stride Output Layout: The physical NPU writes outputs using a deterministic formula where Row 0 has a gap of exactly the active slice width $N_c$ elements, and all subsequent rows have a pitch of $N_s$ (total stride) elements: $$\text{Row_m_start} = \begin{cases} 0 & m = 0 \ N_c + (m - 1) \times N_s & m \geq 1 \end{cases}$$
3. Zero-Allocation In-Place De-Tiling: To reconstruct a standard contiguous layout for CPU reference buffers from the strided outputs, we perform a back-to-front CPU gather copy (memmove looping row $M-1$ down to $1$ and column slices right to left). This avoids overwriting uncopied data and requires exactly zero extra memory allocation.

Validation: Verified first via a standalone hardware diagnostic probe in tools/stride_test.c (STRIDE TEST PASSED BIT-EXACTLY WITH OFFSET STRIDE!). Followed up with native compilation and test runs of the complete driver regression suite on the RK3588 target board under active ORK_ZC_OUT=1, passing all FP16 and Int8 checks with zero mismatches (mism=0, maxerr=0.0000 EXACT).

Full writeup → Output Zero-Copy and Strided Tiling

Tools: tools/int4_bench.c + make test (RK3588, governors pinned)

Implemented manual ARM NEON SIMD optimizations for four key CPU-side data layout and processing hot spots in src/npu.c:

1. Grouped int4 Dequantization (i4_mcworker_g): Widened and converted 16-bit NPU outputs to single-precision float vectors and accumulated with vector scale products (vmlaq_f32).
2. Nibble-Packing (tile_i4_Aslice): Used vld1q_s8, unzipped even/odd elements via vuzpq_s8, shifted odd elements, and combined with bitwise OR (vorr_s8) to pack 16 bytes into 8 bytes of 4-bit nibbles in 6 vector instructions.
3. Widening Accumulator (i4_mcworker): Leveraged vaddw_s16 and vaddw_high_s16 to accumulate int16_t outputs directly into int32_t buffers.
4. Partial Matrix Accumulator (mcworker): Vectorized the multi-core pre-fill accumulation using 128-bit vector additions (vaddq_f32, vaddq_s32).

Validation: We successfully verified this by running the full test suite (make test), passing all tests with zero mismatches (mism=0). Standardized benchmark results show that GCC's auto-vectorizer already generates highly optimized NEON streams for simple loops under -O2, and execution is dominated (>85%) by the kernel ioctl submit floor. Explicit vectorization ensures high performance across all compilers/optimization flags.

Full writeup → CPU NEON SIMD Optimizations

Tools: examples/i4.c + examples/bench.c (RK3588, native compiler, governors pinned)

The Solution: Completely eliminated the instant hardware wedges on the RK3588 during INT4 multi-core regression testing, and resolved the multi-row prefill accuracy degradation (M > 1) on native INT4 pipelines.

Key Findings:

1. The regcfg_amount Config Length Alignment: Configured regcfg_amount = 116 for INT4 task structures (matching REGCMD_I4_N = 232 registers). Previously, setting it to 108 (the INT8 default) truncated critical address-mapping and stride configurations, causing immediate physical bus lockups and hardware wedges.
2. Decoupled Asynchronous Core Execution: Removed unstable user-space thread barriers (pthread_barrier_wait) and unified multi-core ioctl submissions. Replaced them with independent asynchronous single-core submits (sub.core_mask = 1u << i) on isolated, core-specific task buffers (c->mtk[i].obj), completely bypassing the unstable RKNPU driver multi-core scheduler and race conditions.
3. Forced Sequential Prefill Scheduling: Verified that the physical INT4 datapath is structurally single-row (mc = 1) and ignores multi-row scheduling. Forced sequential single-row scheduling (chunk_M = 1) inside INT4 workers, completely restoring mathematical exactness (0.0000 maxerr vs. CPU reference) on prefill-sized shapes without any perceptible latency penalty.

Validation & Benchmarks: Added automated regression test cases inside examples/i4.c for prefill-sized shapes on multiple active cores (test(c,8,512,256) and gtest(c,8,1024,256,64)). All tests now compile and pass flawlessly. Evaluated throughput with examples/bench.c on Qwen3-1.7B:

Full writeup → INT4 Stability & Multi-Core Refactoring.

Tools: tools/i4_multim_fuzz.c + examples/test_layouts.c (RK3588, governors pinned)

The Breakthrough: Definitively resolved the outstanding reverse-engineering task for Fuzzing for the Undocumented int4 Multi-M Scheduler on the RK3588 silicon. We have proven that native multi-M execution (computing multiple rows in a single hardware submit) is fully supported by the INT4 hardware datapath.

Key Findings:

1. The PPU Stride (0x405c) Constraint: Unlike INT8 which sets the output PPU stride register 0x405c to (mc-1)<<16, the INT4 datapath requires 0x405c to be left at its default value 0 (disabled). Overriding it with a non-zero value completely corrupts the output and halts rows $&gt;0$ from being written.
2. The Stride-2 Output Layout: When multiple rows are computed natively, the output DMA step logic operates by int32_t width while the computed elements are int16_t. This results in computed logical row m being written to physical output buffer row 2 * m (a stride of exactly 2 rows of int16_t).
3. The Register Formula: Setting mc = 2 * H allows the NPU to compute exactly H rows in a single pass.

Validation: We successfully verified this by expanding examples/test_layouts.c to run a large-scale $M=32$ (mc=64) native INT4 submit. After implementing a stride-2 de-tiling gather copy (reading from even rows 2 * m * N in the output buffer), the NPU output matched the CPU reference bit-exactly across all rows with 0 mismatches. This completely eliminates the PC-chaining row-by-row fallback, unlocking optimal prefill speeds for INT4.

Tools: ork_mm_pack_i4_to_i8 API (RK3588)

Implemented the Tier 4 Memory Optimization fallback: storing weights as 4-bit on disk/RAM but executing them on the NPU's int8 physical path.

The Problem: The native int4 hardware path physically executes 2× slower than int8 due to silicon-level constraints (ceil(K/128) K-pass loop + int16 intermediate accumulation), and its accuracy is severely impacted by activation outliers without heavy CPU-side Hadamard rotations. The Solution: We introduced ork_mm_pack_i4_to_i8, a dedicated API that takes int4 weights (unpacked in int8_t containers) and packs them physically into the NPU's highly optimized int8 resident memory layout.

Verdict: ✅ By deferring to the fast physical int8 datapath, we achieve the best of both worlds. The caller (e.g. ggml-ork) maintains the massive 50% footprint reduction on disk while the NPU runs at its absolute maximum speed, completely avoiding the intrinsic int4 execution latency penalty and preserving high accuracy without requiring runtime rotations.

Tools: examples/perplexity_i4.c (RK3588, 32 tokens, stories15M.bin)

Rank 2 of the reinvestigation list: Does the Hadamard rotation proven in pure-CPU diagnostics (Tier 4a) actually improve perplexity when routed through the NPU's physical W4A4 int4 math?

Built a standalone perplexity_i4.c to evaluate the Karpathy 15M model natively. It pads the dim=288 dimension to dim=512 to satisfy the power-of-2 requirement for Fast Walsh-Hadamard Transform (FWHT), quantizes to int4, and runs autoregressive generation, comparing the NPU's probability outputs against the pure-CPU FP32 ground truth token.

Note: 1.00 is a perfect match with the FP32 unquantized probabilities.

Implementation Architecture (Option 2): Instead of embedding the FWHT math inside the low-level ork_mm_pack_i4 and ork_mm_run_i4 driver calls (which would hide massive memory padding bloat and silently block the NPU submit threads on heavy CPU matrix math), we exposed void ork_fwht_norm(float* v, int n) as a public API utility.

Verdict: ✅ The Hadamard rotation successfully stabilizes W4A4 probability retention natively on the NPU, particularly saving outlier steps (e.g. at step 23, the baseline W4A4 probability for the correct token was 0.18; Hadamard restored it to 0.99). Exposing it as an independent API utility preserves ork-driver's philosophy as a raw, asynchronous submitter while enabling the higher-level framework (ggml-ork) to fuse the operation into its own compute graph (e.g., merging it with RMSNorm) or apply it offline during GGUF conversion.

Tool: ork_mm_run_chain_i4 API (RK3588)

Implemented the Rank 1 optimization to bypass the NPU's structural M=1 limit for int4. The int4 datapath physically ignores multi-row scheduling registers (the "Graft" was proven inert), which previously forced us to send M separate ioctl submits from the CPU. This resulted in a catastrophic latency floor (~134µs per row).

To fix this, we exposed ork_mm_run_chain_i4, which decomposes any M>1 task into M separate single-row hardware tasks. Instead of submitting them independently, it chains them using the NPU sequencer's built-in 0x0101 trailer registers (rc[216..219]). All M rows are submitted in a single hardware ioctl (up to 512 rows).

(Note: Independent ioctls incur a ~134µs kernel latency floor per row. At M=512, this is a ~68ms dead-time penalty. PC-chaining entirely eliminates this penalty).

Verdict: ✅ The NPU correctly executes the chained single-row tasks. While PC-chaining forces the hardware to re-fetch the weights from DDR for every single block (which previously ruled it out for int8/fp16), the math flips for native int4. Int4's half-size weights reduce the memory bandwidth penalty, making it entirely worth the sacrifice to completely eliminate the massive kernel ioctl overhead. The latency floor is successfully amortized.

Tools: ork_npu_probe_i4_mm + tools/i4_multim_probe.c (RK3588)

Pursuing the full int4 path, the first sub-task is giving int4 the multi-row M-scheduler int8 has (to fix the 1-submit-per-row prefill wall). Attempt #1 — feed M rows and set int8's M-count regs (0x4034=mc-1, 0x4038=(N/4-1)) — corrupted even row 0: those output-side M-regs don't apply to int4. Leaving the M-regs verbatim and laying A per-row-contiguous (ORK_I4_ALAY=1) makes row 0 bit-exact while rows 1–3 are uncomputed garbage.

Finding: the captured int4 task[0] is single-row (M=1); the "M=4" in regcmd_i4.h is the closed runtime's 12-task program (M tiled across tasks). So — unlike int8, which had a multi-M capture to anchor its scheduler — multi-M int4 has no capture reference. It must be derived from scratch: graft the int4 precision regs (0x100c=0x360, 0x1080, 0x3010=0x601, 0x4010, weight-bytes (K*N)/2, K-pass ceil(K/128)) onto int8's validated multi-M skeleton (REGCMD_I8 + its 0x1010/0x1040 scheduler), reconciling the different A/B/C byte layouts (int4 B (N/64,K/32,64,32) 2-per-byte vs int8 (N/32,K/32,32,32)). Materially more than the "~day" first estimate — and the ceiling is still ~2× int8 (intrinsic per-submit cost), so the payoff remains footprint, not speed. RE tools landed (ork-driver main); the graft is the open work.

Tool: tools/int4_bench.c (RK3588, warm, governors pinned)

Tier 4b: does per-channel int4 (one scale/output channel → ork_mm_run_i4, a full-K single submit like int8) kill the grouped path's K/G-submit explosion and reach int8 speed? Timed per-channel int4 vs grouped (G=128) vs int8 on the same shapes.

Per-channel beats grouped 7–8× ✅ — but int4 is still 2–13× slower than int8, for two independent reasons: (1) pc-i4 scales linearly with M (~0.05 ms/row) — it does one submit per row (the captured int4 regcmd is M=1; int8's multi-row M-scheduler regs are unset); (2) even at M=1/one-submit, int4 is 2× int8 (the int4 ceil(K/128) K-pass + int16 accumulate). So porting the M-scheduler would only get int4 to ~2× int8 — never parity, because the NPU is submit-bound (int4's half-size weights buy nothing) and int4's per-submit cost is intrinsically higher.

Verdict: there is no int4 path that beats int8 speed on this NPU. int4's only value is memory footprint, which forks into: int4-stored / int8-computed (smaller download, free, = int8 speed+accuracy) or the full int4 matmul (fit a bigger model in NPU RAM, ~days of regcmd RE + Hadamard, ~2× int8). Full writeup → Per-channel int4 submit.

Tools: hadamard_int4.c (synthetic) + hadamard_real.c (real GGUF Q4_K weights) · CPU math, no NPU

Started Tier 4. Before wiring a Hadamard pre-rotation into the model pipeline, two CPU probes measure whether it actually reduces W4A4 matmul error — synthetic Gaussian, and a real Q4_K weight dequantized from a GGUF (random data is worst-case for quant, so the real-weight run removes that artifact).

Three robust facts (synthetic ↔ real agree): (1) Hadamard cuts W4A4 error 40–60%; (2) the benefit scales with outlier severity (×20: 46→19; ×4: 23→13.5) — the QuaRot signature; (3) it's the activations — Hadamard helps W4A4 far more than W4A8 (int8 acts already absorb outliers), confirming activation-int4 is the dominant error source, and W4A8 would be far better but the NPU can't do it (same-precision MAC).

Verdict: ✅ mechanism confirmed — but matmul RMS is a diagnostic, not the verdict. int4 error is unbiased and averages out across the matmul + residual layers (why QuaRot int4 LLMs work despite double-digit per-matmul RMS), so go/no-go needs end-to-end perplexity from a Hadamard-wired path. Sequencing call: 4b (speed) gates 4a (accuracy) — int4 is still ~23×/1.8× slower than int8, so it's only worth the Hadamard integration if submit-count reduction lands first. Full writeup → Hadamard int4 accuracy.

Branch: feature/zerocopy-buft (llama.cpp-rockchip, not merged) · evidence ORK_BUFPROBE

Pressed past the "blocked" hypothesis by actually building the DMA buffer type and measuring. Result: it works partially, but delivers no practical end-to-end benefit.

1. Baseline ACCEL placement (get_buffer_type = CPU): ORK_BUFPROBE → src1=CPU | dst=CPU — both matmul operands in CPU buffers, so zero-copy can't fire.
2. DMA-backed is_host buffer type (get_buffer_type = ORK-DMA): ORK_BUFPROBE → src1=CPU | dst=ORK-DMA. So ggml-gallocr does place the matmul output in ork's buffer (even for ACCEL), and is_host=true means no cross-backend copies; generation stays coherent. But the input (src1, a CPU-op output) stays in a CPU buffer → input zero-copy still can't fire.
3. Output zero-copy end-to-end (single-core + ORK_ZC_OUT=1, since the multicore mcworker path has no zero-copy and would need a strided write): fires + coherent, but pp512 unchanged: 143.11 → 142.85 tok/s — the writeout it saves is a negligible fraction of end-to-end prefill.

Verdict: zero-copy is conclusively closed for this integration — not "deferred, might help" but implemented and measured at ~0 gain. The input half (the bigger one) is unreachable without a GPU device type (which would cross-backend-copy inputs anyway + risk the decode path), and the output half saves too little to matter. The validated matmul-level capability (ork_dma_alloc, input+output zero-copy correct) ships in main; the integration buffer type lives on feature/zerocopy-buft as the documented no-gain result (measured, parked). The open stack stays at prefill 178 (97%) / decode 12.8 (109%) — parity; no further lever here measured worth pulling, though a future GPU-input device type could revisit it.

Commit: 3fad74a · validated vs the CPU reference (examples/quant)

Attempt #3 (after #2 was reverted) cracked it with a localized fix — no global submit/scheduler changes. Root cause was not a second output-address register (my earlier hypothesis was wrong): the output DMA buffer's dirty CPU cache lines evict over the NPU's DMA writes, and stale lines are read back. Fix = the standard coherency dance on the zero-copy output buffer: bsync(cbuf, TO_DEVICE) before the matmul (clean dirty lines) + bsync(cbuf, FROM_DEVICE) after each submit (invalidate). That exactly explains the single-tile-~90%-wrong / multi-tile-mostly-right symptom (dirty-line eviction hits the first/only tile hardest).

Scope: validated at the matmul level (single-N-slice, single-core), still opt-in (ORK_ZC_OUT, off by default) and matmul-only — output zero-copy fires only when the caller allocates C via ork_dma_alloc. The ggml integration doesn't yet, so llama-bench is unchanged (cbuf NULL → no-op). Wiring it in = Stage 2 (a ggml-ork DMA buffer type), now in progress. (Lesson: the earlier reverted attempt failed by making global changes; the fix was localized coherency — diagnose the data path, don't perturb the scheduler.)

Commit: 22887de (local, not pushed — reverted) · evaluated against the CPU reference + mc_prof

A follow-up attempt to fix output zero-copy changed three things globally instead of localizing to the output-address bug: submit1's post-submit sync bsync(Cc) → bsync(c->task); synth_i8's 0x1010 row register 16*(mc+1) → 16*mc; and the int8 full-K chunk 4*R → R. Measured (Qwen3-1.7B-w8a8, RK3588):

The bsync(c->task) change is almost certainly wrong (the result is in Cc/the output buffer, not the task descriptor) and would break the K-split path + RE probes that read Cc after submit1. The 0x1010/chunk edits touch every int8 matmul, regressing the already-validated input-zero-copy and non-zero-copy paths. The working tree was also left non-compiling (TK typo in mcworker) with scratch RE tools (brute_ac.c, test_second_reg.c, print_dma*.c) from an unfinished hunt for the output-address register. Reverted to the validated state (fe47dcd: input zero-copy default-on + correct, output zero-copy off) and re-validated: quant ALL OK; full-K 1-core M=512 = 4694 µs normal / 4020 µs input-zero-copy (matches the pre-attempt baseline). (Note: the DDR governor had drifted back to ondemand@528 MHz, briefly inflating the restored number to ~9800 µs — a parked-DDR artifact, re-pinned to performance@2112 MHz per step 3 of the benchmark standard; not a regression.) Output zero-copy remains open — the right fix is still localized regcmd debugging (dump where the result lands vs 0x4020), not global submit/scheduler changes.

→ Optimization Roadmap, 1c-i · commit 39d5369 · ork_dma_alloc / ORK_TEST_DMA

ork copies the activation into a DMA buffer + writes the result back per matmul; RKNN keeps A/B/C in NPU-coherent memory (zero-copy). Built ork_dma_alloc/ork_dma_free (NPU-coherent CPU-mapped buffers + a registry); the matmul auto-detects when A/C are resident and points the regcmd at them in place — no host memcpy. No API change to ork_mm_run_i8.

Input zero-copy: validated correct vs the fp32 reference (quant ORK_TEST_DMA, RMS unchanged) — gather eliminated, default on. Output zero-copy: buggy — redirecting only the regcmd's 0x4020 result address gives single-tile ~90% wrong output (a 2nd output-address register needs RE); opt-in ORK_ZC_OUT, off by default, deferred. End-to-end still needs a ggml-ork DMA buffer type so the activation lands in a DMA buffer (it's produced by a CPU op into a host buffer; ggml-sched placement is the open question) — Stage 2, deferred (we're at 97% prefill, so this is push-past-parity).

→ matmul page, prior-art · commit bc8739b · rknn_vs_ork <n> a

Dug the closed RKNN matmul header (it documents its own native layouts) + RK3588 GEMM blogs to find the residual prefill lever. Two findings: ork already uses RKNN's native int8 B layout (N/32,K/32,32,32) (no native-B win exists), and the only unmatched layout — native A — is 14% slower at the prefill shape, measured risk-free on RKNN itself:

The fp16 blog's "+20%" doesn't transfer to int8 large-M. So both layout levers (1c-i) are closed without touching ork's risky regcmd. Only residual = zero-copy DMA buffers (host copy/writeout).

→ Optimization Roadmap, Tier 3

Feasibility verdict. 3a persistent per-layer program: BLOCKED — multi-matmul-per-RKNPU_SUBMIT times out (kernel rejects task_number>1, re-confirmed ntask 2&4 → rc=-1), and a layer's non-matmul ops (attn/softmax/RoPE/norm/SwiGLU) force CPU round-trips so intermediates can't stay on-NPU. 3b async overlap: net-negative — all-NPU decode 213ms/token (NPU 72ms + CPU 141ms), perfect overlap → ~7 tok/s < shipped hybrid 12.8. Goal (beat rkllm decode) already met by Tiers 1–2; no Tier 3 work pursued.

→ CPU-affinity page · commit b9a5c4e

Tried forcing process affinity to the big cluster at init (so llama.cpp's pool inherits it — the taskset effect, automatic). It doesn't help -t 4 (already optimal) and craters -t 8 decode (9.3 → 2.3, 8 threads oversubscribed on 4 cores). Reverted; the driver pins only its own NPU-driver threads. Big-core win stays a thread-count choice (-t <big-core-count>).

→ Optimization Roadmap, Tier 2 · commit 0a25292

Split the M=1 decode submit to see if the per-matmul cost was trimmable. It isn't: 98% is the irreducible NPU ioctl/IRQ round-trip floor; regcmd-prep is 2%, writeout 1%.

Resident regcmd / cache-sync / poll-vs-IRQ all chase ≤ a few µs (or a risky busy-spin). And ork's per-matmul decode is already at RKNN parity and beats librkllmrt end-to-end (12.8 vs 11.71). Tier 2 closed.

→ full writeup · commit c1bb9cc · knob ORK_FULLK_PREFILL (default on)

Prefill (M>1) now does one submit per M-tile over the full K (reusing the decode Bf layout + M-scheduler) instead of K-splitting — removing the second K-slice's host accumulate + result cache-sync. Validated correct vs the CPU reference (power-of-2 K; non-pow2 like 6144 falls back).

Combined with big-core threading, the open stack is now prefill 178 (97% of closed) / decode 12.8 (109%, beats) on the same model.

→ full writeup

Tier 1b answer: the integration's prefill was never NPU-matmul-bound. With -t 8, llama.cpp's CPU threadpool spins on all cores (starving the ork NPU-driver threads → matmul 7596 vs 4365 µs standalone) and the 4 little A55 cores drag the threadpool barrier. Running -t 4 (big A76 cluster only) fixes both — and decode then beats librkllmrt.

Decode 12.84 > 11.71 (110% of closed); prefill 84%. The remaining prefill gap is ork's single-core large-M kernel (3.3× vs RKNN per core).

→ full writeup · commit 34b539e · tools/mc_prof.c · knob ORK_NO_AFFINITY=1

The multi-core pool workers were never CPU-affinity-pinned → the scheduler parked them on the A55 little cores, ~2× slower than the calling thread on every phase, collapsing prefill multi-core scaling. Pinning each NPU-driver thread to a big core fixes the worker imbalance. End-to-end int8 is unchanged (the integration's budget-3/Bf pack is a separate bottleneck — new Tier 1 follow-up), but int4 prefill improved 1.57×.

→ full writeup · tools/rknn_vs_ork.c

The same int8 (M,K,N) matmul run through ork_mm_run_i8 (open) and rknn_matmul_run (closed librknnrt) — both on the same NPU. The true kernel-level comparison the full-model tables can't give. Splits the gap: decode ≈ 1.24× (per-matmul kernel efficiency), prefill ≈ 3× (large-M tiling). (RKNN matmul API is single-core; ork-1core is the clean comparison, ork-3core shown for context.)

µs/matmul, 200 warm iters. The ~1.24× decode matches the end-to-end decode gap to librkllmrt; the 3× prefill gap pins ork's large-M tiling as the prefill lever.

→ full writeup · commit 4956b054 · knob ORK_MINM=32

The offload gate (M>=min || N>=min) always passed (every weight has large N), so M=1 decode was dragged onto the NPU — a ~2× loss. Gating on M only routes decode→CPU, prefill→NPU. Keeps the prefill win, kills the decode penalty.

→ design blueprint · active architecture validation

Analyzed and resolved the "0% NPU load" anomaly during Mixture of Experts (MoE) prefill. The static $M \ge 32$ gating threshold (amortizing the 365µs NPU driver submission latency) is a dense-model bias that fails catastrophically on MoE models. Under MoE routing, token splitting reduces individual expert batch sizes to tiny fractions (e.g., $M_{\text{expert}} = 4$), which naive gates mistakenly offload to the CPU/GPU.

• The Core Bottleneck Inversion:
  • CPU/GPU Offload: CPU/GPU must stream all non-resident expert weights (e.g., 25 MB per expert layer, or 800 MB per forward pass at 32 layers) from LPDDR system RAM. At 25 GB/s bus bandwidth, this introduces a massive ~32 ms memory transfer bottleneck per token.
  • NPU Resident Run: The NPU keeps weights resident in physical IOVA buffers. It only pays the 365µs driver context-switch latency. At 3 submits per active expert, NPU latency is only 1.1 ms (a massive ~30× speedup).
• The Gating Policy Fix: Transition from a naive "Batch-Size Gating Policy" to a "Weight-Resident Gating Policy" in the inference engine's dispatcher (llama.cpp / ggml-ork). If weights are resident (like 4-bit experts), the $M \ge 32$ guard is bypassed (threshold set to $M \ge 1$), keeping all expert math strictly on the NPU.
• KV-Cache Compression: Walsh-Hadamard coordinate rotation (TurboQuant/PolarQuant) implemented via Mali-G610 OpenCL shaders (or CPU NEON fallback) allows the KV cache to be compressed to 4-bit (GGML_TYPE_Q4_0), shrinking the 8k context window RAM from 2.1 GB to 525 MB.
• The 35B RAM Math: Keeps the total footprint of Qwen 35B-A3B at ~21 GB (14 GB 4-bit experts, 7 GB 8-bit attention, 525 MB KV-cache, 1.5 GB OS), comfortably fitting inside the 32GB RK3588 memory limit.

→ full writeup · commit bd900ae · ork_npu_run_timing(), ORK_PROFILE=1

Tested whether the open-vs-closed gap was in run() CPU machinery. Hypothesis refuted — machinery is 3%, the rest is the genuine NPU submit. Redirected effort to which matmuls offload (→ M-gate).

rkllm comparison: librkllmrt is closed (no per-matmul instrumentation); its effective decode of 11.71 tok/s (~86 ms ÷ ~197 matmuls) implies it pays a similar ~1-submit-per-matmul floor — i.e. the floor is the runtime-independent wall, not an ork-driver run() artifact.

→ full writeup

Forced a fair comparison after earlier over-optimistic claims. Standardized (llama-bench/rknpu_bench, warm, same model) the gap is 1.36× prefill / 1.23× decode — not the "3.9× / 2.6×" from the old llama-cli numbers (those were polluted by short prompts + the lazy weight-pack).

→ full writeup · commit b72a84a4 · knob ORK_NOREUSE=1

Skip re-quantizing the activation when consecutive matmuls share src1 (Q/K/V, gate/up). The standard relocates the benefit to prefill (~8%); decode is neutral (quant is ~2% of the NPU-floor decode). Free + correct → kept on.

→ full writeup · commit 1a0a39e8 · knob ORK_FUSE=1

Fuse same-input matmuls into one wide submit. Revised under the standard: on the 1.7B it's neutral (within noise) — not the decode 9.4→6.4 loss seen earlier on 0.5B/llama-cli. Buys nothing measurable here + adds complexity, so off by default.

→ full writeup · commit 95f2481 · tools/batch_probe

Tried task_number>1 (many matmuls per ioctl). Kernel-level submit probe (not tok/s): the batched ioctl times out, NPU dispatches nothing. The closed runtime's 12-task submit is one matmul's program, not batched matmuls — the submit floor is inherent.

rkllm comparison: librkllmrt hits the same ~11.7 tok/s decode wall on this model — consistent with it also paying ~1 submit/matmul (it does not batch matmuls per submit either).

→ full writeup · commit 20186ed2

SIMD-friendly loops kept; OpenMP threading reverted (net-hurt: ~400 tiny parallel regions/token). The reference profile shows quant/dequant is too small to matter — no quant micro-opt can move decode.

rkllm comparison: the closed runtime's decode (11.71) is bounded by the same NPU run, not by quant — so this CPU-side lever is irrelevant to both.

→ full writeup

No software wattmeter (USB-C PD reports the 45 W contract, not draw). Thermal proxy: NPU ~7–9 °C cooler (lower instantaneous power), but energy/token likely favours the CPU on small models. Power is the missing axis — needs an external meter + a 1.5B+ model.

→ full writeup · commit 60a83221 · ORK_PROFILE=1

Decode is submit-latency-floor-bound: a small projection costs the same ~ per-matmul as a giant one because the RKNPU_SUBMIT round-trip dominates. ~197 matmuls/token → the floor. The central finding framing every decode experiment.

rkllm comparison: librkllmrt decode 11.71 tok/s ≈ 86 ms/token ÷ ~197 matmuls ≈ ~440 µs effective/matmul (with some CPU/NPU overlap) — the same order as our 359–386 µs. The closed runtime hits the identical floor; it just shaves the per-submit cost and overlaps better, which is the remaining ~1.2× decode gap.

→ full writeup · commit 81533e27 · knob ORK_QUANT=4

Per-group int4 weights + activations (same-precision MAC ⇒ no w4a16). Runs, but output is incoherent (~9.5% error, needs Hadamard) and ~23× slower on prefill — K/G submits multiply the floor. Strictly worse than int8 here. (int4 numbers are the post-affinity build; pinning lifted int4 prefill 3.66 → 5.73.)

rkllm 4-bit note: rkllm's own 4-bit is w4a16 (weight-only 4-bit, ≠ our W4A4) — a different scheme — and its 1.7B w4a16 .rkllm failed to load on this board (target_platform does not match), so no rkllm 4-bit point is available here; the w8a8 reference stands in.

→ full writeup · commit 0eabf5f

Per-channel int8 weights + per-row int8 activations, int32 accumulate, fp32 dequant. Coherent, ~1% error, half the bytes of fp16. The practical default; natural target from GGUF Q8_0 — and the same scheme librkllmrt uses (W8A8).

→ full writeup · commit 4484b5b · knob ORK_NOCACHE=1

Pack+quantize each weight once, keyed by host pointer, NPU-resident. Standardized A/B is binary: without it (re-pack every matmul) the run dies with CREATE: Bad address — at 1.7B it's a correctness prerequisite, not just a speedup. (On 135M, when no-cache still limped: decode 17×, 0.53→9.34.)

rkllm comparison: librkllmrt packs weights once at model-load (resident for the session) — the same strategy this experiment introduced; it never re-packs per matmul.

• Exp 2026-06-23 Hardware Chaining Zero Overhead: Zero-overhead pure hardware chaining bypassing kernel driver queueing via CORE_CTRL=0x000f and regcmd concatenation. ⚠️ Superseded — not reproducible in integration; see Resolution below.
• Exp 2026-06-23 Chaining Resolution: ✅ run_chain_i4/i8 fixed by mirroring the validated i4_mcworker PC-chaining recipe (per-task entries, task_number=S, timeout=60000, warmup). Different-weight S=4 verifies; full make test green. Corrects the concatenation "breakthrough" (it timed out in run_chain_i8). Includes the kernel rknpu_job_subcore_commit_pc mechanism + RK3588 PC constants.
• Exp 2026-06-23 MoE Expert Offload: ⚖️ Qwen3.6-35B-A3B MoE experts on NPU — not yet viable but closing. Full path built (MUL_MAT_ID handler, group-by-expert, resident repack pool). Prefill pp128 walked from ~80× under CPU (group-by-expert) → ~23× (expert chaining) → ~7.9× (NEON-fused repack) → ~7.2× (fused Q4_K→int8 dequant, b520ee7/b5bc5ce8b, current) — see the 2026-06-24 follow-up on that page. New wall: the NEON f32→int8 quant+tile itself (77%; dequant now folded in at 0 ms). Experts stay CPU-default (ORK_MOE_NPU off). Gotcha: ORK_MOE_NPU=0 still enables NPU MoE (getenv-presence check) — unset it for a CPU baseline. ★ Partly superseded (2026-06-26): that net-loss story was repack/thrash/single-core, not the GEMM — the resident per-expert submit beats CPU even at M=1. See Exp 2026-06-26 MoE Decode Partition + Submit Floor.
• Exp 2026-06-23 Decode Profiling: ✅ GGUF-Q8_0 decode ~4 t/s is not a regression — ork-driver is only ~12% of decode wall time (~88% is CPU/ggml + weight bandwidth); the ~11–12.8 figure was the w8a8 path. ORK_PROFILE per-matmul timing added. Chain-batching confirmed not firing for decode (non-consecutive matmuls; Qwen3 QK-norm) — but the upside for dense decode is <4%; the lever is MoE/prefill.
• Phase 1B PPU LUT Silicon Verification: ❌ FAILED / INACTIVE STANDALONE. The captured PPU Sigmoid LUT/PWL regcmd does NOT execute as a standalone replay on RK3588 — the submit returns rc=0 but writes nothing (output buffer comes back byte-identical to the input sweep, idx0=0xc400=fp16 −4.0). PPU activation appears to require full compiled-model-graph init (clock/power-gating). The prior "✅ bit-exact" claim was an artifact of a CPU-sigmoid fallback that falsified the gate (removed, ork-driver d400d93). PPU activations stay on CPU. ★ SUPERSEDED 2026-07-04 (below): PPU activations DO run standalone bit-exact — that attempt drove the wrong program shape.
• Exp 2026-07-04 PPU Element-wise / Activation Ops: ✅ On-NPU SiLU / GELU / EW-mul / residual Add / exp / rsqrt — bit-exact on silicon, in make test (branch ewmul-re). The RK3588 NPU is NVDLA-derived; element-wise + activation are STANDALONE SDP ops (enable_mask=0x18, 69 regcfg) reading operands from memory — the activation is a LUT op whose contents define the function (SiLU/GELU/exp/rsqrt = one op + curve swap). int8 bit-exact (linear index idx≈2·in+512); int16 RKNN-class (~0.24% FS). assemble-from-dump (tools/re/assemble_op.c): transplanting RKNN's LUT + index regs as a matched pair reproduces an op bit-exact-to-RKNN with no index decode. RMSNorm reduction = matmul (RKNN builds ReduceMean as einsum→matmul→conv; no dedicated reduce op) → RMSNorm/Softmax/RoPE compose from these blocks; RKNN's fused exNorm(78)/exSoftmax13(42) are multi-task chains for framework composition. Reusable RE toolkit added under tools/re/. Two walls: (1) fp16 activations need NVDLA's LE-table exponential mode + multi-task submit (LE SRAM doesn't persist across separate submits like LO did); (2) int16 residual Add — the SRDMA/X1 operand halves negatives (int16-specific sign/shift), needs an X1/X2 sub-module scale decode. Both need multi-task submit (shared unlock, also for the fused composites).
• Exp 2026-07-13 Linear-Attention Scans on NPU: Mamba-2/SSD ✅ wins ~2× at Q8 2.7B (byte-coherent; the 130M "no crossover" verdict was a small-scale + Q5_K-dilution artifact — CPU recurrence craters at scale, NPU amortizes) — size-gated default-on (nh≥64≈≥1.5B). Gated DeltaNet (Qwen3.5/Qwen3-Next/Ornith) ❌ built + coherent end-to-end but not a win: ORK_GDN_PROF isolates the scan → loses ~1.9×@0.8B / ~4.8×@9B (44% NPU-submit, 38% CPU UT-solve the CPU recurrence skips; no crossover — GDN state fits cache). Keep-warm (int8↔fp16 no-thrash) is the general 5× unlock, now default on. Reusable machinery banked (weightless-op offload + densify marshaller + fused-multicore fp16 stream) → next: attention-on-NPU. GDN parked on ork-driver feat/gdn.
• Exp 2026-10-11 Weight Ownership Under Mmap: ✅ ggml-ork disowned every weight whenever mmap was on. The device iface aliased .get_host_buffer_type to .get_buffer_type, so llama_model_loader::create_tensor's "avoid a host buffer under mmap" rule fired on ORK_Weights itself: Qwen3-1.7B-BF16 moved 311 tensors to the CPU backend, 115 after the fix. The iface also contradicted the device's own get_props (host_buffer = false) — the loader trusts the iface, not the props. Compute was never stranded (ork runs buffer-less via offload_op, 196/196 weights resident either way); ownership was. The fix is two halves that must land together: dropping the alias alone routes weights through alloc_buffer's posix_memalign + set_tensor memcpy, which over a stub source copies gigabytes of holes into anon RAM — so buffer_from_host_ptr also had to stop returning a CPU-buft buffer and return ORK_Weights over llama.cpp's mmap window (free_buffer = NULL, borrowed mapping). Result: ORK_Weights model buffer = 2688.46 MiB, 0 bytes copied, prefill flat (165.85/164.57 → 166.21/165.51), ork_ppl bit-identical (1.6587, NLL 0.506032). New [ORK BUFTYPE] line states ownership positively and names the mechanism (mapped vs copied) — "owned via memcpy" and "owned via mmap" otherwise print the same word. ⚠️ Side finding: tests/test-ork-cache.cpp is RED on pristine diverged-history (returns 0.0, not the 64.0-twice stale-cache signature), and never checks graph_compute's return status, so a backend failure is indistinguishable from a wrong answer — the backend's only automatic correctness gate is dead and CI never runs it (LLAMA_BUILD_TESTS=OFF). Use test-backend-ops support -b ORK instead (safe, no submits; baseline 0/1638 MUL_MAT claimed — a supports_op-change detector, not a correctness gate).
• Exp 2026-10-12 W4A4 Tier Trade: ✅ The 4-bit ladder works — and mixed-precision packs cost 18× for a reason that is one XSPEC row per direction, not an intrinsic tier cost. Four packs, one model (Qwen3-1.7B), one baseline, disjoint GPTQ calibration: i4a8 DT_I4 PPL 34.30 / 169.46 tok/s; rotation-only native 50.73 / 46.84; arm T (+SmoothQuant +GPTQ +attn_v,ffn_down promotion) 26.36 / 9.08; arm T with ORK_I4_PROMOTE=i8 25.78 / 9.10. So native 4-bit built with the landed ladder BEATS the i4a8 tier on quality by 1.30× (promotions verified by census: 56 I8/I8_ROT = attn_v+ffn_down × 28 layers, ≈5.16 bpw). ⚠ The 18× was initially mis-attributed to the DT_I8_ROT tier — that was an inference from an average and it was wrong. ORK_I4_PROMOTE=i8 (plain DT_I8, off the rotated path entirely) measured 9.10, unchanged; and ORK_DEBUG_RESET=1 named the caller: mixed pack = 221 ACT_RESETs (110 × profile=8 from=3 to=4 + 110 × profile=2 from=4 to=3), homogeneous packs = 1 each (cold entry only). 110 = the promoted-weight invocation count; every promoted weight drags the datapath I8_CHAIN ↔ I4_CHAIN and back. At the 105 ms/reset this repo already measured (mode.c:148, 27B mixed-tier run), 220 × 105 ms = 23.1 s against a measured 25.0 s run — closes within 8%. ORK_MIXED_NOTHRASH=1 didn't help because it governs XP_MC_MM, not profiles 8/2 — so the knob could never affect it. ⚠ CORRECTION (same day): the reset is LOAD-BEARING and must not be relaxed — fc06bc6, the commit that added ORK_XSPEC_NORESET, already measured it on the real 27B mixed-tier workload at PPL 10.4421 → 3,442,988.92, and 42cc655 measured the complementary keep-warm widening at 5,029,965. A clean suppressed-reset run on the 1.7B (K=2048, single domain, DT_I8_ROT) is the expected non-result — fc06bc6 says its own positive control stays coherent at that scale and added the control precisely to stop such a run being misread. The 6.17× stands only as a ceiling. The lever is transition COUNT, not the reset: promote contiguous whole layers and the forward crosses the boundary twice instead of ~110 times, every reset intact (42cc655: "the uniform G=512 pack takes the same model from 160 resets to 2"). Consequences: every promotion scheme is mixed by construction, and MoE packs (int4 experts + int8 attn/dense) are mixed by design and may have been paying this silently. Corrections forced: the source's "~54×" native-prefill claim was stale by ~15× (real 3.62×) and its stated cause ("single-row/no-M-amortization") wrong — the path chains 281.7 programs/ioctl; native buys nothing at decode (1.02 tok/s on all arms — dense M==1 never reaches the NPU); ORK_GPTQ_CLIP_MIN=0.531 must NOT be set (live default 0.10, self-tuning — 0.531 would raise the floor); GPTQ already skips the int8 tier (finalize: 140 weights = 196−56). Build cost measured: GPTQ 23.7 of 28.5 min on the 1.7B → ~18 h extrapolated for a 27B (O(K³)), the largest open risk to defaulting it. ⚠ Latent surprise: stub vs full GGUF changes which engine runs decode for every packed weight (:11393 early-accept fires ahead of the N-cap, sbits and M>=min_m gates) — benchmark the same pack two ways and get different answers.
• Exp 2026-10-12 INT4 BCHAIN Dispatch Envelope: ✅ The int4 dispatch envelope is one law at two DRAM-clock points — not two dies, and not a capacity limit. ~190 + ~66 graded cells across both boards. .236 (LPDDR4X @ 2112) has the uniformly LARGER envelope and the SLOWER DRAM; .235 (LPDDR5 @ 2400) the narrower, and lowering DMC grows the envelope monotonically on BOTH parts. ⚠ A per-program weight-tile capacity law was proposed and refuted the same day — equal byte counts land on both sides (K=512 Wb=192 passes, K=768 Wb=128 misses, both 49152 B), sweeping the CBUF bank split moves nothing, and raising H rescues a byte-identical tile: it is a RACE with a lower bound on H, and it looked like capacity only because it was first measured at M=H=2, the worst corner. So orki_i4_hcap is "the .236 measurement", never hardware — do NOT add a second table keyed on the part, that encodes a race as a law; the ladder is the remedy, and its first rung must be today's shipped width (aeb50f7: rung 1 used to jump below the default and cost −25% permanently). Cost is governed not by program count but by r, the FINAL M-group's row count (T ≈ 1140 + 8·r; at fixed H=32, M=128→129 is +0.8% work and −21% time). ★ Four methodology traps that each produced a wrong answer: arm order changes the PASS RATE not just timing (6/6 @436 µs ascending vs 1/6 @16.1 s descending); checksum-only grading ships self-heals; a 3-rep arm cannot distinguish 1/6 from 0/n; and grading against the probe's own iteration 0 lets a consistently-wrong config agree with itself — use an absolute CPU reference.
• Exp 2026-10-12 Shared Doorbell Drain Defects: ✅ Three defects in the SHARED NONBLOCK doorbell spine (src/npu/core/) — not int4 bugs, merely exposed by an int4 workload that misses constantly. (1) the miss-window EWMA key omitted the geometry, so a geometry the board takes and one it will not shared a slot and the window was wrong for both (fix: orki_geomkey(H, Wb)); (2) the 50 ms miss floor sat below one ~50 ms poll pass, so miss detection could not resolve finer than ~15× the op's correct runtime — now measured as 3× the first complete pass rather than a constant calibrated on one shape; (3) ★ the recover re-seeded all of nout including the 75% of each slot that is padding the NPU never touches, so after any recover the adjudicator was structurally unable to accept on this path ever again (census: 92–94% survivors, ~75 pp of it padding). ★ The generalisable lesson: a correctness-preserving recovery path hides its own defects, because the only symptom is time. All three are failure-path code that a board which never misses never executes — and the ladder returning a bit-exact answer is exactly what kept them hidden. Anything that self-heals needs a timing gate as well as a correctness one.
• Exp 2026-10-12 The fp16 M Envelope Is A Predicate: ✅ Started as a performance question and turned out to be a correctness bug on the MAIN fp16 path, at exactly attention's shape — 54 of 126 swept cells returned wrong answers with no error and no decline. The rule is entirely the 0x1040 bank split: CORRECT iff dreq <= DATA_BANK AND DATA_BANK + min(wreq, WEIGHT_BANK) <= 12. At every well-formed K the encoding keeps DATA_BANK + WEIGHT_BANK == 12, so the model reduces to the banks × 32 KB cap already in the tree; K=128 is the sole exception (it needs half-bank steps the encoding cannot express) and that is head_dim, so ork_f16_mm_run was silently wrong at K=128 N=2048 M=64 (max|diff| 5.12 against |C|max 2.17). It looked correct only because the per-core column split happened to shrink Ncore below the boundary at the N values anyone had tried. ★ An envelope expressed as a CEILING cannot describe a rule with a hole in it, and a caller that only knows a ceiling cannot decline what it cannot compute — expressed as a predicate (orki_f16_mc_ok), three callers respond differently by what each can do: orki_run pads, colsplit narrows to the round's widest per-core N, the stream pair refuses (−2). Result: 0 of 126 cells failing, and the chain's usable region GREW (every M ≥ 80 at K=128 N=1024/2048, including 144/160/192/288/320/448, which were wrong before). Cost stated plainly: K=128 N=2048 M=64 goes 66.1 → 35.3 GMAC/s — a 2×-tall program buys the correct answer. Open: ork_f16_mm_run_stream is ~100× slower than the chain at every K, correct or not (538 ms vs 4.5 ms) — a separate, unfixed defect in the fp16 doorbell path.
• Exp 2026-10-13 INT4 BCHAIN Completion Gate: ✅ Where the int4 call's time actually is, measured rather than apportioned — and ~50–55% of the available bound taken. At M=160 H=32 K=512 N=2048: seed + flush + spin together are worth ~4%; the drain-verify is ~26% (the per-poll scan is cheap because it early-exits; the expense is the post-completion full-surface verify, serialised behind a kernel sync covering 4× the written cells, then a second pass re-reading the same cells to de-tile). Shipped three changes that alter how the gate executes and not what it checks — fan the verify out over the existing pool, invalidate exactly the lines a program wrote (★ fail-safe by construction: a still-stale line reads back the seed, so it can delay completion but never declare it early), and fuse verify + de-tile into one pass — for 1.15–1.24×, ORK_BCH_SLOWGATE=1 restoring the old path for bisect. ★★ A kernel fence IS available (CONFIG_ROCKCHIP_RKNPU_FENCE=y, RKNPU_JOB_FENCE_OUT = 1<<4), worked first try with no reflash — and it is 13% SLOWER. Do not build on it: it signals at job_done, and int4's DPU writeback lags job_done, so the whole drain-verify still runs afterwards, cold, while poll() wake-up costs more than the 430 µs spin it replaced. A fence is a correct completion signal, not a data-visible one. (This is also what "fencefix" in uname -r means.) Correctness graded at a calibrated sensitivity — the negative control (no gate at all) is caught in only ~2–4% of iterations, so "it passed" is meaningless without that number — and the page states honestly that the harness cannot demonstrate the drain-verify is load-bearing on these shapes. ★ The residual ~13% is the cost of not trusting job_done, and the only thing that would delete it is an out-of-band completion witness, which §6 records as not constructible (task order does not imply write-visibility order: 1,998 violations in 34,480 gates).