Prefill NPU Utilization — measuring the real ceiling, and why it's the synchronous execution model, not the matmul

Supersedes the "NEXT: regcmd-RE the per-core kernel" conclusion of Prefill Parity Grind. That entry inferred (from a per-matmul rknn_vs_ork comparison) that the matmul kernel was the remaining lever. Direct per-core NPU+CPU utilization sampling shows it is not: the matmul kernel is fine; the gap to rkllm is that the NPU sits ~60% idle between matmuls on CPU work the synchronous ggml backend can't overlap. RK3588 10.3.0.236, governors max, complete .orkpack resident across 4 IOMMU domains, -t4, warm.

Why this entry exists: every "fast" number was a config artifact

A cross-session audit found the working 40-t/s build was uncommitted in both layers (ggml-ork glue + the ork-driver submodule) and lived only in agent-mutated build dirs; the fork's committed master had an older, broken ggml-ork that loaded 0 weights onto the NPU. Recovered, committed, and verified reproducible: fork master fc0d618c8 → pp512 45.3 / pp256 40.1, 10 GiB resident across 4 domains, loaded 168. All numbers below are on that pinned, committed build. (Methodological note for the community: pin the exact config — build commit, "loaded N / packed M", resident GiB, governors — before trusting a t/s. The whole audit was undoing numbers that weren't.)

The trustworthy gap (re-verified, same board, same governors)

rkllm measured fresh with ~/rkbench/rk_bench (standalone C++, rkllm.h + librkllmrt v1.2.3, big cores, governors max) on Qwen2.5-7B-Instruct-rk3588-w8a8.rkllm:

	prefill tok/s	decode tok/s
ork (open stack)	40.1 (pp256) / 45.3 (pp512)	~0.9 (churn-bound)
rkllm (librkllmrt 1.2.3)	72.1 (163 tok) / 68.9 (1390 tok)	4.06
ratio	0.56× / 0.66×	0.22×

Confirms the documented 73.89 / 4.07 reference within variance. Prefill is the focus here (decode is a separate ggml-ork/CPU/M=1-churn problem).

The measurement that overturned the matmul theory: per-core utilization

Sampled /sys/kernel/debug/rknpu/load (per-NPU-core) + per-CPU-core /proc/stat at 150 ms during prefill, ork vs rkllm:

	prefill tok/s	NPU Core0/1/2	big-core CPU (cpu4/5/6/7)
ork	45.3	38 / 40 / 41 %	14 / 20 / 20 / 56
rkllm	68.8	66 / 64 / 64 %	36 / 27 / 27 / 21

• Both use all 3 NPU cores (an earlier sampler had a field-parse bug reading Core1/2 = 0 → a false "single-core" claim; corrected here). Multi-core is not the lever; it's already done (ORK_NPU_MC=3 is a no-op).
• NPU clock is maxed (1000 MHz) on both — the dmc_ondemand NPU governor is a non-issue.
• The gap is per-core saturation: ork ~40 % vs rkllm ~65 % — and that 1.6× ≈ the 1.52× throughput gap. ork's NPU cores idle ~60 % of the time between matmuls.
• ork has one hot CPU core (cpu7 = 56 %), rkllm is balanced — because run_multicore's caller thread (pinned to cpu7) drives NPU core 0 and does the serial setup + cres→C copy and orchestration.

Root cause, traced to the code and proven

mc_prof back-to-back matmuls (no ggml graph between) vs the live prefill:

	NPU Core0	hot CPU core
back-to-back matmul (no graph)	77 %	cpu7 = 20 %
live prefill	40 %	cpu7 = 56 %

With nothing between matmuls the NPU runs hot and the CPU is idle; in the live prefill the NPU drops and the CPU rises. The delta is exactly the per-matmul CPU work the NPU can't overlap. From the code: ggml-ork is a synchronous, BLAS-style backend — each MUL_MAT node does quant → run_i8 (blocking) → dequant, and ggml's scheduler runs graph nodes strictly one at a time. So the NPU idles during (1) per-matmul activation-quant + dequant, (2) the serial run_multicore setup + cres→C copy on cpu7, and (3) the CPU graph ops between matmuls (RMSNorm/attention/softmax/RoPE). Nothing overlaps NPU compute with CPU prep of the next node. rkllm hits 65 % because its own integrated executor pipelines that prep against NPU compute.

How much is recoverable: the overlap-ceiling probe

tools/overlap_probe.c (new diagnostic, committed db46800, like mc_prof) pipelines the per-matmul quant+dequant behind the NPU run (serial vs double-buffered):

shape	quant	run(NPU)	dequant	(q+dq)/run	serial→pipelined
FFN 512×3584×18944	11 ms	94 ms	11 ms	24 %	1.20×
attn 512×3584×3584	11 ms	18 ms	2 ms	72 %	1.64×

The double-buffer reaches ~95 % of the theoretical ceiling — the mechanism works. But two caveats collapse the realistic number: (1) the probe's quant/dequant are scalar; ggml-ork's are NEON + OpenMP-parallel, so the model's real costs (t_actq 2822 ms + t_deq 2143 ms vs t_run 32521 ms = 15 %) give a ~1.15× ceiling, not 1.2–1.6×; (2) overlapping quant-of-next with run-of-current needs the next node's input, which the synchronous backend doesn't have — that's the ggml-sched async wall again.

What's already done, and what's left

Within the backend, the easy wins are already in the committed code: activation-quant and dequant are NEON + #pragma omp parallel for; submits are chained (lever-1 cut them 55→5/matmul); all 3 cores run. Pinning OpenMP to the big cores (OMP_NUM_THREADS=4 + places) measured 45.05 vs 45.72 — no win.

The last within-backend candidate — pipelining run_multicore (overlap the host synth/copy/scatter behind the NPU submit) — was measured directly with a synth/ioctl timing split (mc_prof, committed 25f89ce): the NPU ioctl is 94–95% of per-submit time, the synth (the part a double-buffer would hide) is 0.4%, and the overlappable host total (copy+synth+scatter) gives a theoretical pipeline ceiling of only ~1.06× (FFN) / ~1.10× (attn) — before double-buffer/threading overhead. So it was not pursued: ~1.06× on the densest, most wedge-prone path is below the threshold worth risking the working build (measure-first paid off — no risky rewrite built for a capped ~5%).

Verdict & next

The prefill gap to rkllm is an execution-model problem (synchronous vs pipelined), not a matmul/kernel problem. The matmul saturates the NPU fine in isolation (77 %); per-matmul auto-tuning (R/KS/cores/tiling) can't help a bottleneck that lives between nodes; the within-backend overlap ceiling is ~1.15× and the bigger overlap hits the async wall. The real remaining lever is cross-node async execution in ggml-sched (overlap NPU compute of matmul N with CPU prep of N+1 + the intervening graph ops) — architecturally hard, upstream-scale (the same wall the MoE concurrency work hit; meta-backend territory). ⏭️ Prefill is banked at a trustworthy, committed 40.1 (0.56× of rkllm) with a fully-mapped, async-gated ceiling. Higher-leverage fronts now: decode (0.22×) and the distributed/speculative path (EAGLE-3/DFlash batched verify), which put the NPU on separate M>1 work and sidestep this sync wall instead of fighting it.