Status: DRAFT specification, not yet started. Target: the RK3588 CLASS of edge SBCs — 32-bit NPU IOVA, A76-class NEON, Vulkan iGPU, shared LPDDR (validated on RK3588 board .236; RK3576 inherits the code path with untested caps). Training host: Apple Silicon M5 Max / 128 GB UMA.

Provenance note. Some [M] figures below are measured on the training host (Apple Silicon M5 Max / 128 GB) by the sibling qwen35-lora-moe project rather than on the RK3588 board — every such figure says so inline. They bear on Phase 1 and Phase 2, never on board behaviour.

How to read this page. Every number is tagged. [M] = measured on our board, with the experiment linked. [D] = derived arithmetic from [M] figures. [P] = proposed, not yet validated. [?] = external claim we have not verified and which gates part of the plan. The design below is coherent under [M] and [D]; the [?] items are the real schedule risk and are collected in §7.

The thesis: prefill and decode are different machines' problems, so stop running them on the same engine.

Three independent measured reasons, any one of which is fatal:

Measured end-to-end this is not close: NPU M=1 decode ≈ 7.2 t/s against CPU 12.8–13.67 t/s on the same model [M]. Decode has been the CPU's job in this project since that measurement, and this design makes that split architectural rather than a fallback.

One correction to the premise in the brief. M=1 decode is described as "strictly memory-bandwidth bound". On this board it is latency-bound first: a single A76 achieves ~8 GB/s of the ~21.9 GB/s it can stream, which is why a one-cache-line PRFM prefetch is worth 9–11% [M] (test_nf4_decode). Bandwidth only becomes the binding constraint once several cores are streaming concurrently. This matters for the plan: quantization alone will not deliver the projected token velocity — thread count, prefetch distance and expert-gather locality carry comparable weight. Budget engineering time for all four.

MoE expert offload to the NPU is closed in all regimes measured — the binding constraints are the fused-CPU-MoE gap, fusion lost on partition, and the 4 GB IOVA budget; even a free NPU loses [M]. The measured-optimal profile for a routed MoE on this board is already asymmetric, and this roadmap generalises it rather than inventing it:

• prefill (M>1): experts on CPU (batched 4×4 NF4 NEON GEMM), attention/dense int8 on NPU [M]
• decode (M=1): all-CPU fused gemv [M]

Reference point on qwen3.6-35B-A3B: prefill 36.9 t/s (1.59× native) but decode 6.3–6.6 (0.85×), PPL +5.5% [M].

Read that decode figure correctly — it is easy to misuse. Both rows are the same routed-MoE model: it compares our runtime against ggml's, both sparse, both on CPU. It is not evidence about sparse vs dense. Two known contributors: NF4 re-quantization instead of Q4_K (which also produces the +5.5% PPL), and it is an accepted trade for the 1.59× prefill — ORK_MOE_AUTO=0 exists to decline it. Treat it as a ~15% implementation gap in our expert path, not as a verdict on the architecture.

The sparse premise itself is arithmetic, not conjecture: decode reads active parameters while dense reads all of them, so at equal decode bandwidth sparsity buys total capacity — precisely the currency a bandwidth-bound device lacks. This blueprint's 2.1 B / 636 M active costs the same per token as a dense 636 M model while carrying 3.3× the parameters.

The IOVA wall and the bandwidth wall have different solutions, and disaggregation lets us apply each where it works:

• IOVA wall — only the prefill slice needs to be DMA-mappable. The NPU IOMMU (rk_iommu v2) is 32-bit; we measured EFAULT at 4.03 GiB and the driver reports a 3900 MiB ceiling [M]. Keeping the NPU-resident set at 2.1–3.0 GB leaves real headroom rather than living on the edge.
• Bandwidth wall — decode never enters IOVA at all. It runs INT4 in 64-bit host memory where the format is unconstrained by the NPU datapath.

The cost this brief does not price: weight duplication. The same tensors exist in two precisions — W8A8 for prefill, INT4 for decode. Under config (B) (§1.6) only the dense backbone is duplicated into IOVA; the routed experts are held once as INT4 for decode plus a W8A8 copy in host RAM for CPU-side prefill. For the §2 model: 0.51 GB (IOVA) + 1.59 GB (host W8A8 experts) + 1.05 GB (host INT4) ≈ 3.15 GB of weights [D]. Affordable on an 8 GB class member and comfortable on 16–32 GB, but it must appear in the RAM budget from day one rather than be discovered in Phase 3. If RAM is tight on an 8 GB board, the W8A8 expert copy is the one to drop — accept NF4 experts in prefill too and take the quality/throughput hit, which Phase 1.4 should quantify.

Heterogeneous engines are additive only on different resources — measured directly: a memory-bound third engine adds nothing (Mali streaming cut CPU+NPU aggregate 29.3 → 16.9 GB/s), while a compute-bound third engine is fully additive (Mali memory-free kernel held 99.9% while CPU+NPU held 99.3%) [M]. The DRAM ceiling is ~30 GB/s at DMC 2112 MHz, not the 20–25 GB/s in the brief [M].

Consequence: sequential prefill→decode is safe. Pipelined prefill(chunk N+1) ‖ decode(chunk N) puts two memory-bound engines on one bus and will be zero-sum. If pipelining is wanted later it must be measured, not assumed, and /sys/class/devfreq/dmc/load is the ground truth.

test_nf4_decode was built for this question: 64 distinct expert weights (32 MB ≫ 3 MB L3) walked once, so every K-step is a cold DRAM read like real MoE decode. It measured the M=1 loop latency-bound — one A76 reaching ~8 GB/s of the ~21.9 it can stream — with a one-cache-line PRFM worth 9–11% [M]. That is a tuning surface with headroom, not a wall, and it is the best evidence we hold that sparse decode on these cores is viable. Use this probe (not test_i4_gemm, whose single 0.5 MB weight stays L3-resident) for any expert-path work.

The target board is actively cooled, which has two consequences worth stating because they are usually assumed away on SBCs:

• Every figure in this document is a sustained-clock measurement, not a burst one — the ~30 GB/s bus ceiling, the 21.8 GB/s achieved decode, and the three-engine additivity results all hold under continuous load [M].
• Sustained multi-engine operation is a supported design point, which is what makes the Phase 4 Vulkan path (fully additive for compute-bound, memory-light kernels [M]) worth evaluating at all.

It does not unlock a persistent GPU spin-wait: Vulkan gives no forward-progress guarantee between workgroups, and the KMD watchdog kills long-running jobs. Those are correctness and driver constraints, not thermal ones. A bounded polling kernel behind timeline semaphores is the version worth measuring.

Neither engine runs at its datasheet number, so both projections below are anchored to a rate this board has actually achieved rather than to a theoretical ceiling.

Decode. The reference point: Qwen3-1.7B-w8a8 decodes at 12.8 t/s [M]. That is 1.7 GB/token × 12.8 = 21.8 GB/s achieved on 4×A76 — which also settles §1.1's nuance: per core the loop is latency-bound (~8 GB/s of 21.9 [M]), but in aggregate at -t4 it does reach the bus. Applying that achieved rate to this model:

Prefill. Measured 178 t/s on 1.7B dense = 3.4 GFLOP/token → ~578 GFLOP/s [D], which is ~10% of the 6 TOPS headline. That is not a compute limit: prefill is CPU-bound on this board — activation quant ~45%, resolve ~43%, NPU only ~11% [M]. Applying the same achieved rate to 636 M active params (1.27 GFLOP/token):

The actionable consequence: the NPU is not the prefill bottleneck and adding NPU TOPS would not help. The lever is the host side — the precompiled regcmd cache and quant/NPU overlap — which is another reason §6 matters more than it looks.

The target is the RK3588-class edge SBC — 32-bit NPU IOVA, A76-class NEON cores, Vulkan iGPU, shared LPDDR4x/5 — not one board. That changes which constraint binds, and the answer is not the one the brief assumes.

Two configurations, and the choice determines everything downstream:

(A) is the trap sparsity is supposed to avoid. If every expert must be DMA-mappable, the 4 GB window caps total parameters at ~3.5 B — so sparsity buys decode speed but not capacity, which defeats the point of a sparse MoE on a memory-constrained device.

(B) is the measured-optimal profile already and it inverts the constraint: the NPU only holds the dense backbone. For the §2 blueprint that is mixing 0.239 B + shared experts 0.198 B + embeddings 0.076 B = ~0.51 GB at W8A8 — seven times under the window [D]. IOVA stops being the binding constraint entirely.

Adopt (B). The consequences are worth stating plainly, because they reverse §4's emphasis:

• The real ceilings become RAM (total params) and bus bandwidth (active params), not IOVA.
• There is spare IOVA to spend: a wider dense backbone, more attention layers, or resident activation scratch for larger prefill chunks are all affordable.
• The design ports downward: a class member with a smaller window still holds a 0.5 GB backbone.

active_params   ≤  (achieved_bus_GB/s ÷ target_tok/s) ÷ bytes_per_param      # decode velocity
npu_resident    ≤  iova_budget − activation_scratch − egress_ring            # prefill residency
total_params    ≤  (ram_budget − kv − state − activations)
                   ÷ (bytes_W8A8_slice + bytes_INT4_slice)                   # dual residency

Worked for this class at 21.8 GB/s achieved [M] and INT4 decode:

The §2 blueprint at 636 M active is conservative — there is room to roughly double active parameters and still clear 40 t/s. Given (B) frees the IOVA budget too, the natural rebalance is a larger dense backbone (better quality per active parameter, and it runs on the NPU where prefill is cheap) rather than more routed experts. Settle it in Phase 1.3/1.4 against measured expert-gather locality, not by intuition.

We have run this, so it does not need re-litigating from first principles. Mali-G610 exposes VK_KHR_shader_integer_dot_product with integerDotProduct4x8BitPackedSignedAccelerated = true, and on a real int8 GEMM, verified bit-exact, it measured ~32× behind the NPU [M].

Worse for a three-engine plan, the additivity rule is sharp [M]:

The iGPU shares the same ~30 GB/s bus, so for anything memory-bound — which is all of decode and most of prefill — it cannot help and will take bandwidth from the engines that can. The architecture must not depend on it.

The one regime that survives is compute-bound and memory-light. Plausible candidates, none yet measured [P]: the GDN chunked-form UT-solve (38% of scan time on CPU [M], small matrices, arithmetically dense), sampling/logit post-processing, and expert index permutation. Treat Vulkan as an opportunistic Phase 4, gated on a DMC-load measurement showing the candidate kernel is genuinely memory-light.

(Test-harness note: any Vulkan probe must reject llvmpipe explicitly — a software rasterizer will happily produce plausible-looking numbers that mean nothing [M].)

ork-driver already detects the SoC at runtime and carries one caps file per chip (src/soc/<chip>.c: core count, CBUF budget, output-width cap, K-slice) — RK3588 is hardware-validated, RK3576 shares the code path with inherited, untested params [M]. The model must consume these as parameters, never constants: mcap, nmax and core count all move across the class. A blueprint that hardcodes mcap = 192 will silently miscompute on the next chip.

Sized so that active parameters × 0.5 B/param ≈ 300 MB/token, which is the brief's decode target.

d_model            1536
n_layers           28          21 GDN  (75%)  +  7 attention (25%), attention at layers 3,7,11,15,19,23,27
vocab              49152       tied embeddings

Totals [D]

The dimensions above are not free choices — they are picked to satisfy every limit this project has measured. Any future edit to the blueprint must re-check this table.

Two notes that fall out of the table:

• mcap = 192 at K=1536 sets the prefill tile geometry. A 512-token prefill chunk is 3 M-tiles per program; a 2048-token chunk is 11. Weight DMA amortizes over rows within a tile, so larger prefill chunks are strictly better for the NPU — bounded by IOVA scratch, not by correctness.
• The weight-reuse gate we just shipped (K*segw ≤ WEIGHT_BANK*32KB and mc == mcap, eligibility N ≲ 192 [M]) engages on exactly one tensor in this model: the MoE router (N = 16). Correct and free, but negligible — do not design around it.

GDN/LRU gives O(1) state update, which removes the KV cache from 21 of 28 layers [D]:

Caveat, measured and load-bearing: the on-NPU GDN scan loses to ggml's solve-free CPU recurrence — ~1.9× at 0.8B and ~4.8× at 9B, and the gap widens with scale [M]. Root cause: the chunked form adds a CPU UT-solve (38% of scan time) that the plain recurrence skips, plus submit overhead, and there is no crossover because the GDN state fits in cache. So: GDN belongs on the CPU decode engine and must not be scheduled onto the NPU. That is consistent with this architecture, but it forecloses a "just move the scan to the NPU" optimisation someone will inevitably propose.

If sequence-mixing offload to the NPU is ever genuinely wanted, Mamba-2 / SSD is the variant with evidence: it loses 0.7× at 130M but wins ~2× byte-coherent at 2.7B (9.72 vs 4.37 P128) [M]. That crossover is real and is the reason this document does not simply say "linear attention is CPU work".

§1.1 corrects the brief to say decode is latency-bound first (~8 GB/s of a possible 21.9 per core [M]). That is exactly the regime where speculative decoding pays: it converts latency into compute the board already has idle. It is not mentioned anywhere else in this document and should be, because the 75% GDN choice prices it.

What GDN costs a speculator [M]. A recurrent state cannot be rolled back by truncating cache slots the way a KV cache can — once history is folded into the state, a rejected draft token cannot be subtracted out. The standard remedy is snapshot-and-recommit, ~2 target forwards per accepted step, which puts break-even at acceptance length τ > 2 rather than τ > 1. On a model that is 75% GDN this is the dominant term, not a detail.

What it is worth if paid [M]. Measured on Qwen3.5-4B, which shares this document's ~75% GDN pattern, on Apple Silicon:

The third row is the warning: a drafter 9× too small and trained on ~140× too little data returns nothing, while sitting in the serving path costing latency. Speculative decoding is not a switch to flip; it is a second model with its own data requirement.

Two cheaper mitigations than snapshot-and-recommit, both measured elsewhere: cache the SSM inputs rather than the state so rollback becomes a pointer move (vLLM carries an RFC for the same idea), and note that a fully-accepted block needs no rollback at all — the "2 forwards always" model is pessimistic and the gap between it and measurement is itself worth reporting.

Recommendation: treat this as a Phase 4-class option, gated on a measurement, and do NOT size the decode velocity projection in §1.5 as if speculation is available. If it is wanted, the drafter's data budget belongs in Phase 1 alongside the base model's, not discovered later.

KV quantization measured as a 4× attention tax, not a bandwidth win [M] (Experiment Log 2026-09-03). With only 7 attention layers the fp16 KV cache is 344 MB at 32k, which is affordable. Keep it fp16.

The brief's ~2.5 GB static state per 1B params decomposes as:

For the §2 model (2.1 B): ~5.3 GB static [D], leaving ~120 GB of a 128 GB UMA for activations, batch and sequence length. The binding constraint on this host is Metal compute throughput, not memory — the memory plan is comfortable and should not drive architecture decisions.

§3.1 concludes that Metal compute throughput binds this host rather than memory. Both are true and neither is what stops a long run. Metal limits a device to 499,000 buffers per PROCESS, and the count climbs with iterations — it is a COUNT, unrelated to how much RAM is free.

Measured on the same M5 Max / 128 GB training host, by the sibling qwen35-lora-moe project:

The failure text is RuntimeError: [metal::malloc] Resource limit (499000) exceeded, and a sibling form is [Event::Event] Failed to create Metal shared event — the latter raised with 101 GB free. mx.synchronize() plus mx.clear_cache() between iterations slows the approach to the ceiling but does not prevent it; only process exit releases the resources.

Consequence for Phase 1.5. A ~2.1 B pre-training run is many thousands of steps. In one process it will hit this, and a crashed run that only writes weights at completion loses everything it did. The mitigation is cheap and must be in the training loop from the start rather than retrofitted:

• chunk the run so each process performs at most a few hundred optimizer steps, and
• persist after every chunk, off-host where possible, so a dead process costs one chunk.

Cost is one model load per chunk (~7 s measured). The sibling project's auto_resume.sh and train_experts_chunked.sh implement exactly this pattern and are the reference implementation.

MLX native FP8 training is unverified [?]. MLX's mature quantization surface is 4/8-bit for inference; FP8 training with Metal-side scaling is a materially different claim. Do not schedule around it until proven. Phase 1 exit criterion: a 2-layer toy model trains stably in the intended precision for 1k steps, loss curve matching a bf16 control. If it fails, the fallback is bf16 master weights + FP8 storage, which costs ~2 GB/1B more and still fits comfortably.

Similarly [?]: Muon's Newton–Schulz orthogonalization and GaLore's projections both need Metal kernels that may not exist in MLX at production quality. Prefer Muon if either works — dropping Adam's FP32 second moment is the larger and more reliable saving, and it has fewer moving parts than GaLore's periodic subspace re-estimation.

                 FP8 / bf16 checkpoint (MLX)
                            │
          ┌─────────────────┴─────────────────┐
          ▼                                   ▼
  W8A8 symmetric per-channel            INT4 / NF4 per-channel
  → NPU prefill slice                   → CPU decode slice
  (int8 A · int8 B → int32 C)           (grouped, NEON gemv/GEMM)

Hard constraints from measurement:

• Symmetric only. RK3588 matmul is symmetric-quantization only [M] — no zero-point on the NPU path.
• W4A8 does not exist [M]. Int4 on the NPU means W4A4 and incoherent activations. The INT4 slice is therefore CPU-exclusive by hardware, not by preference.
• Both slices must agree numerically. Prefill produces the state that decode continues from; a quantization mismatch shows up as a quality cliff at the prefill→decode seam, which is invisible to per-op tests. Gate: perplexity measured across the seam, not per-engine.
• Attention precision is a file property and the single biggest decode lever measured: q8_0 attention is the balanced optimum (36.9 prefill / 6.59 decode) vs f16 (38.4 / 5.26) and Q4_K (26.8 / 7.42, below the sbits<5.0 floor at which the NPU declines the tensor entirely) [M].

╔═══════════════════════════════════════════════════════════════════════════╗
║  32-BIT IOVA WINDOW  —  NPU DMA-mappable, hard ceiling 4.03 GiB [M]       ║
║  driver-reported budget 3900 MiB [M]                                      ║
╟───────────────────────────────────────────────────────────────────────────╢
║  W8A8 prefill weights — DENSE BACKBONE ONLY (config B)   ~0.51 GB         ║
║    mixing 0.239 B · shared experts 0.198 B · embed 0.076 B                ║
║    routed experts are NOT here — they live in host RAM (§1.6)             ║
║  regcmd program cache (ork_pc_*, precompiled)            ~16–64 MB        ║
║  activation / A-stage scratch (M-tiled)                  ~128 MB          ║
║  ┌─────────────────────────────────────────────────────┐                  ║
║  │ EGRESS RING  32 MB (double-buffered 2 × 16 MB)      │                  ║
║  │ computed states / recurrent tokens → host memory    │                  ║
║  │ NOT the KV cache, NOT the context                   │                  ║
║  └─────────────────────────────────────────────────────┘                  ║
║                                                                           ║
║  TOTAL ≈ 0.67 GB   →  ~3.2 GB headroom under the 3900 MiB budget          ║
║  IOVA IS NOT THE BINDING CONSTRAINT under config (B). Spare window can    ║
║  fund a wider dense backbone or larger prefill chunks — see §1.6.         ║
╚═══════════════════════════════════════════════════════════════════════════╝
              │  DMA out (doorbell-signalled, see §4.1)
              ▼
╔═══════════════════════════════════════════════════════════════════════════╗
║  64-BIT HOST RAM  —  CPU-addressable, no IOVA constraint                  ║
╟───────────────────────────────────────────────────────────────────────────╢
║  INT4/NF4 decode weights (full model, incl. all experts) ~1.05 GB         ║
║  W8A8 routed experts for CPU-side PREFILL (batched NEON)  ~1.59 GB        ║
║  MoE router weights (fp16, latency-critical)             ~2 MB            ║
║  GDN recurrent state, 21 layers                          16.5 MB (const)  ║
║  KV cache, 7 attention layers, fp16                      344 MB @ 32k     ║
║  expert gather staging / NEON worker scratch             ~64 MB           ║
╚═══════════════════════════════════════════════════════════════════════════╝

The brief specifies 64–128 MB. The arithmetic does not support it. Per 512-token prefill chunk, everything that must cross the boundary is:

32 MB double-buffered (2 × 16 MB) covers a 512-token chunk with margin, and IOVA is the scarcest resource in the system — spending 96 MB of it on ring headroom we can prove is unused is the wrong trade. Size the ring from the chunk, and if chunk size grows past ~2048 tokens, recompute rather than round up.

On-chip SRAM is not an alternative here. This board has 956 KB of NPU SRAM enabled (custom kernel, CONFIG_ROCKCHIP_RKNPU_SRAM=y) [M], which is ~30× too small for this ring. It is also the wrong port: SRAM is a confirmed separate port for CPU access but measured null for NPU weight reads (67.7 vs 67.5%, sign flips across 3 sizes) [M]. Do not plan around it.

The brief specifies streaming egress "via DMA fences". The RKNPU uABI exposes fence_fd, but the mechanism this project has actually built and measured is different and better:

• NONBLOCK submit (flag 0x2) + DRAM doorbell busy-poll on stock driver, no reflash: dispatch 130 µs → 5 µs [M].
• Precompiled regcmd cache (ork_pc_*): 6.03× lower per-token program-build cost [M]. The two stack — one is the dispatch half, the other the build half.
• Zero-copy input (ork_dma_alloc) is validated and default-on, −17% on the full-K prefill matmul [M]. Zero-copy output is correct but ~0 gain and unsafe under concurrent multi-core (per-core bsyncs do not serialize against NPU writes) — leave ORK_ZC_OUT off for the egress ring [M].

Design the ring as a single-producer/single-consumer buffer with a last-word sentinel poll, which is the pattern already proven in the colsplit path, rather than a fence-based handshake.

Everything above fits one domain under 4 GB, and that is a design goal, not an accident. Multi-domain operation on this silicon has a history: two IOVA leaks (abrupt-kill teardown; a kernel bug filing every dma-buf import under domain 0 so destroy unmapped the wrong page table), a broken domain-switch path (default_domain lie + force-reset), and non-deterministic layout costing 46 vs 4 t/s. All are fixed by kernel patches carried on our board [M], but a design that never needs a second domain avoids the entire class. If the model ever exceeds the window, stream layers through it (LRU weight cache) rather than adding a domain.

Decision deferred to 1.4, deliberately: whether routed experts belong in the NPU prefill slice at all. Our measured MoE profile puts experts on CPU even for prefill [M], but that was measured on 35B-A3B with large experts; at 768-wide experts the balance may invert. The memory map in §4 assumes experts are NPU-resident for prefill; if 1.4 says otherwise, the IOVA slice drops from 2.1 GB to ~0.5 GB and the design gets easier, not harder.

Do not target ONNX/RKNN. See §6.

Board-safety constraints that bind Phase 3 work [M]: never kill -9 an in-flight submit (wedges the IOMMU → power cycle, and a hard cut mid-submit can corrupt SPI and require physical reflash); run anything touching the NPU under sudo tools/util/npu_guard.sh -- <cmd> on a shared board; use timeout so board-side jobs self-terminate cleanly.

Only after Phase 3 is measured end-to-end. Entry gate: a candidate kernel must be shown memory-light (/sys/class/devfreq/dmc/load near-flat while it runs), because a memory-bound Mali kernel measured negative [M]. Candidates, in order of plausibility [P]: the GDN chunked-form UT-solve (38% of CPU scan time [M], small dense matrices), logit/sampling post-processing, expert index permutation. Non-goal: int8 GEMM — the NPU is ~32× faster [M] and shares the same bus.

The brief specifies a runtime "using librknnrt". This project exists specifically to replace it, and that replacement is measured at parity: open stack 178 prefill / 12.8 decode vs librkllmrt 184 / 11.71 on Qwen3-1.7B-w8a8 — 97% of prefill and 109% of decode [M].

Building on librknnrt would forfeit, concretely:

• the precompiled regcmd cache (6.03× lower per-token build [M]) and the NONBLOCK doorbell (130 µs → 5 µs dispatch [M]) — neither is reachable through the closed API;
• hardware task chaining (task_number > 1), which is how the per-task cost drops toward the ~1.3 µs floor [M];
• direct control of 0x1040 bank splits and the M-tile cap mg_max*64, worth ~2.1× single-core / ~1.6× three-core at K=2048 [M];
• the ability to fix driver-level defects ourselves — the multi-domain IOVA bugs in §4.2 were diagnosed and patched by this project, not by the vendor.

It would also reintroduce a proprietary runtime dependency this repo is explicitly free of.

Recommendation: Phase 3.1 targets ork-driver directly. Phase 2.1 exports the .orkpack tile layout, skipping ONNX/RKNN entirely. The RKNN toolchain remains useful for calibration capture when adding a new op or SoC, which is a development-time need, not a runtime one.

• It does not claim NPU decode can be rescued. It cannot, for three independent measured reasons (§1.1).
• It does not claim heterogeneous execution is free. It is additive only across different resources [M].
• It does not claim GDN-on-NPU works. Measured, it loses and the gap widens with scale [M].
• It does not claim the iGPU contributes. Measured, Mali is ~32× behind on int8 GEMM and a memory-bound Mali kernel reduces CPU+NPU aggregate throughput [M].
• It does not claim speculative decoding is available. §2.1a prices it; the decode projections in §1.5 do not assume it, and a drafter is a second model with its own data budget [M].
• It does not claim class portability. Only RK3588 is validated; RK3576 caps are inherited and untested [M].
• It does not claim 6 TOPS is reachable in practice. That is 3 cores × ~2 TOPS on dense W8A8 with all three saturated; our multi-core colsplit measures ~2.2× single-core, not 3× [M].

Optimization-Roadmap · Experiment-Log · Benchmark-Standards-and-Methodology · Model-Recipes · Vendor-Kernel-Behaviour · Kernel-Modifications