Activity

1. bong-water-water-bong

   added

   enhancementNew feature or requestNew feature or request

   on Nov 2, 2026
2. bong-water-water-bong commented on Nov 2, 2026

   

   bong-water-water-bong

   on Nov 2, 2026

   CollaboratorAuthor

   More actions

   Sub-issues: P1 #800, P2 #801, P3 #802, P4 #803, P5 #804, P6 #805, P7 #806.
3. bong-water-water-bong commented on Nov 4, 2026

   

   bong-water-water-bong

   on Nov 4, 2026

   CollaboratorAuthor

   More actions

   Measured on Strix Halo, pi-onebp's lane (ZR1-1.5B, the Q8 container 1bit-MONSTER/ZR1-1.5B-GGUF/npu-container-q8): npu_plan.py writes decode: {"M": 16, "enabled": true} into plan.json, but nothing builds the *_m16.elf family, so the forward's decode-family probe finds 0 of 4 roles and single-token decode silently runs on the M=128 prefill designs with am=1 — 128× the useful rows per launch.

   What the log says (no code change needed to see it)

   [plan] …/plan.json: 28 layers, gate/up fused, family "rm", 1 decode dir(s) from the plan
   [forward] derived 4 designs (fused_gu=1, H=1536 NH=12 NKV=2 HD=128 IM=8960)
   [forward] decode family: 0 role(s) at M=16, 4 on M=128      <-- the whole problem
   

   Cost, same container / prompt / box, back to back in one ticket

   device path	32 greedy tokens	rate
   M=128 designs used for decode (no _m16 on disk)	189.14 s	0.17 tok/s
   four _m16 designs present (decode family found)	40.13 s	0.80 tok/s

   4.7×, and it puts the device route at parity with the same container's host-GEMM matvec (41.07 s, 0.78 tok/s, same ticket). Both device arms produce byte-identical text to each other ("… The capital of Spain is Madrid. The capital of7 is New York. The"), so the M=16 design is not a different-numerics path here; the difference from host appears only as the device's per-column int8 pack, i.e. the documented quantization, not a design defect.

   The four designs, and how they were built

   full_i8_{QKV,O,GU,D}_K…_N…_m16.elf for ZR1's geometry (1536×2048, 1536×1536, 1536×17920, 8960×1536), M=16, microkernel 16×64×128, generator n1_core_i8_v27.py -r 1, emit_full_elf … --get-full-elf, with the lock the arena already pins (mlir-aie d53582d3e0f + Peano 22 from ~/amd-oss/iron-upstream-venv) — full ELF only, no xclbin, no aietools compiler. Files land in the --kernels dir the plan's decode_dirs resolve to ({"root": "kernels", "path": "."}), which is exactly what the forward scans.

   Why I am commenting here rather than just fixing it

   The gap is in the plan layer's own contract: plan.json names a decode family that the kernel set has no way to satisfy, and the failure is silent — nothing refuses, nothing warns, and the route just runs 128× the rows. Two options, your call:

   1. npu_plan.py (or the packer) emits/validates the decode family: if decode.enabled, the pack must either carry *_m16.elf or say so in the plan, and the forward should refuse/warn by name instead of falling back to M=128;
   2. or the plan stops declaring decode: M=16 until the ELF set ships it.

   Two other facts worth having in this issue while I am here: (a) the same route also re-dequantizes and re-packs every layer's weights inside the per-token loop (step() → load_layer() → npu_gemm() → packB()), ~1.3 G weight floats per token for this model, which caps any device number on the model-generic path at ~1.1 tok/s and is why host-GEMM and device land at the same rate — that is the bigger lever, but it is a larger change and not mine to take; (b) elf_emit.sh's comment that --full-elf-name alone writes the ELF is stale for mlir-aie d53582d3e0f: without --get-full-elf aiecc runs the whole pipeline, writes no ELF and exits 0, so a build can look green and produce nothing.
4. bong-water-water-bong commented on Nov 4, 2026

   

   bong-water-water-bong

   on Nov 4, 2026

   CollaboratorAuthor

   More actions

   Correction to my numbers above: they were taken on the pre-0007 driver, and the driver changes the conclusion

   The arm table I posted earlier today ran on amdxdna srcversion EBC78CDA6FD7F3F6E824064 — the module with the ubuf
   BO-invalidation BUG (the one that put list_add corruption in dmesg on this box at 01:49). This box now boots the
   0007 TDR-recovery module (srcversion 56AB94758C18FEEC838D18C, the same one the c23 release row asserts), so I
   re-ran the same A/B: same container, same prompt, same four _m16 ELFs, same script, one ticket, back to back.

   ticket	host-GEMM	device, M=16 family	device/host
   14:1x, pre-0007 driver	41.07 s / 0.78 tok/s	40.13 s / 0.80 tok/s	0.98× (parity)
   17:0x, 0007 driver	47.13 s / 0.68 tok/s	30.21 s / 1.06 tok/s	1.56×

   Text is unchanged between the two tickets, arm for arm, token for token: host still ends …is Madrid. If7 capitalizes the capital of a country, device still ends …is Madrid. The capital of7 is New York. The. So the driver moved
   throughput, not numerics — the device number is the same quantization, computed faster.

   Attribution, stated carefully because the absolute rates are contended-box rates:

   • the host arm is a pure-CPU matvec and got slower (41.07 → 47.13 s, +15 %) — that is box contention (load 30, fleet
     restart), not the driver;
   • the device arm got faster (40.13 → 30.21 s, −25 %) while sharing that contention, so the delta is the array path;
   • ~10 s over 32 tokens × 112 launches/token ≈ 2.8 ms per launch of per-launch cost that the 0007 module removed —
     consistent with the ubuf BO revalidation the bug was doing on every submit.

   What this changes in my earlier framing. I said the device route was host-bound and that the device "adds nothing"
   over host GEMM. On the patched driver that is no longer true: device is 1.56× the host matvec in the same ticket, so
   the route is now bounded by the per-token host-side load_layer() dequant + packB (unchanged, q4nx_forward.cpp:1275)
   plus a device GEMM that is cheaper than the host one. The ceiling argument stands (that per-token re-pack is still why
   any model-generic device number sits near 1 tok/s), but the sign of the device's contribution flipped.

   Also: my --- amdxdna errors: 0 line in the scripts is worthless on this boot — dmesg is restricted here and returns
   nothing. Re-checked through journalctl -k: 0 list_add / list_debug / aie2_cmd_submit / TDR lines for the whole
   boot (4 amdxdna lines total, all init), and nothing in the ticket window 17:05–17:09. Fixing the scripts to read
   journalctl -k before the next long run.
5. bong-water-water-bong commented on Nov 5, 2026

   

   bong-water-water-bong

   on Nov 5, 2026

   CollaboratorAuthor

   More actions

   Correction #2, and this one retires my own ceiling argument: the per-token host re-pack is 5 % of a decode token

   I said twice in this thread that the q4nx device route is bounded by the host-side load_layer() dequant +
   ctx.packB() that runs inside step(), "~0.9 s/token of host work". That was inferred from the code plus the
   host-vs-device wall clocks, not measured. Measured now — the serve process's own utime+stime (thread-group total)
   read around a fixed request, both arms in one npu-hold ticket, same container/ELFs/prompt, 8 greedy tokens, 0007
   driver:

   arm	wall/token	host CPU/token	of which wall	residual (waiting on the array)
   host GEMM (NPU_INFER_HOST_GEMM=1)	1.541 s	11.20 CPU-s (727 % of wall, ~7 threads)	1.541 s	0
   device GEMM (M=16 family)	1.156 s	0.058 s (5.0 %)	0.058 s	1.099 s (95.0 %)

   Both arms produce the same text, Paris. The capital of Germany is Berlin, so this is a timing split, not a
   different path. Script ~/lb/zr1-npu-cpusplit.sh, log ~/models/npu/zr1-elfs/cpusplit.log, 0 amdxdna lines in
   journalctl -k.

   What follows, and what it costs me:

   1. The re-pack is not the ceiling. On the device arm the host is idle 95 % of a token; caching packed weights per
      layer would buy ≤ 5 % (≈1.06 → 1.11 tok/s). My "host-bound" framing was wrong, and the forward-side change I was
      about to propose to that effect would have been wasted work. Withdrawn.
   2. The token is dominated by the 112 per-layer submissions (28 layers × 4 designs at 1.156 s/token ≈ 9.8 ms per
      launch). Cross-checking against the M=128 arm measured earlier (5.9 s/token, same 112 launches ≈ 53 ms per launch),
      per-launch time is roughly ~3.7 ms fixed + ~0.38 ms per row, i.e. ~38 % of an M=16 launch is fixed cost and the
      int8 GEMM itself is a rounding error inside it (12 M MACs per launch). Those two arms were measured under
      different box load, so treat the split as an order-of-magnitude statement, not a fit.
   3. So the lever on this route is submission granularity — one command list covering the whole token (or a
      persistent chain across layers, or fewer/larger roles per submission) — not weight residency. If the fixed
      per-launch cost is what it looks like, a token could leave the ~1.1 s it sits at today by a large factor.
   4. Separately, and probably worse for us than the NPU question: the host GEMM arm burns 11.2 CPU-seconds per
      token for ~2.6 GFLOP of matvec on this model — an effective ~0.23 GFLOP/s with ~7 threads busy. That is not the
      NPU's problem, but it is why the host route is CPU-hungry on the fleet and it may be worth someone's look on its own.

   Next step on my side is a clean per-launch cost measurement (vary tokens and rows in one ticket) before anyone acts on
   point 3. Flagging here rather than opening an issue because the plan layer is this thread's subject and the fix, if it
   is a fix, is in the same place.
6. bong-water-water-bong commented on Nov 5, 2026

   

   bong-water-water-bong

   on Nov 5, 2026

   CollaboratorAuthor

   More actions

   The per-launch measurement I promised: 5.9 ms fixed + 0.30 ms/row, so half of a decode token is submission overhead

   Four points, one ticket, same contention, 0007 driver (56AB9475…), ZR1-1.5B container, 4 ELFs per family, bare serve.
   Both decode families at two token counts, --kernels ~/models/npu/zr1-elfs/, host CPU from the serve process's own
   utime+stime:

   family	tokens	wall/token	host CPU share	per ELF invocation
   M=16 (decode family: 4 role(s) at M=16)	8	1.543 s	4.6 %	13.78 ms
   M=16	24	1.309 s	4.1 %	11.68 ms
   M=128 (NPU_DECODE_M16=0)	8	7.495 s	2.6 %	66.92 ms
   M=128	24	5.765 s	4.2 %	51.47 ms

   Differencing the two token counts inside each family cancels the once-per-request cost (the prompt's prefill and the
   run-list build), which is why the two M=16 rows disagree on raw per-token cost:

   • marginal M=16: 1.1913 s/token, marginal M=128: 4.8994 s/token
   • at 112 ELF invocations per token (28 layers × 4 roles): 10.64 ms and 43.74 ms per invocation
   • the only difference between the families is rows per launch (16 vs 128), so a two-point fit gives
     ≈ 5.91 ms fixed + 0.296 ms per row per invocation

   At M=16 that puts a decode token at 0.662 s of fixed per-submission cost (55.5 %) + 0.530 s of 16-row GEMM work, on
   1.1913 s total. Host CPU is ~4 % in every row, which confirms the earlier split from the other direction.

   What this means for the route, with the part I cannot claim:

   • The 5.91 ms fixed cost is the one clean win available: 112 submissions per token is a choice, and if the same 112
     GEMMs went down as one command list (or a chain that the runtime does not round-trip to the host per launch) that
     0.662 s/token would go with it — ~2.3× on this route (1.19 → ~0.53 s/token) without touching a kernel.
   • The 0.530 s/token of row work is not waste I can remove from here: the smallest microkernel I have is 16×64×128, so
     a 1-row decode already runs the narrowest tile that exists in the arena's set. Cutting it needs a narrower microkernel
     (would scale the row part, at whatever the narrower tile costs), which is a kernel-set question, not a plan-layer one.
     So I am not claiming the "15 of 16 rows are wasted" saving that the arithmetic invites.
   • Sanity check on the model, not a boast: ~0.03 s/token of the 1.19 s is array work at the 16-row granularity, i.e. this
     route is ~36× above the compute floor today — which is consistent with the fact that the fast NPU lane only exists for
     models someone hand-built resident-weight designs for.

   Limits stated plainly: contended box (ticket queue 180 s, die 62 C at grab), one model, one geometry, per-invocation
   here means one ELF invocation (the column-group loop inside an ELF is not split out), and the fit assumes the M=16 and
   M=128 designs have the same column-group structure — true for these two builds (cols 8/4/4/4 in both), not something I
   measured from the artifacts. Script ~/lb/zr1-launch-cost.sh, log ~/models/npu/zr1-elfs/launch-cost.log, 0 amdxdna
   lines in journalctl -k.
7. bong-water-water-bong commented on Nov 5, 2026

   

   bong-water-water-bong

   on Nov 5, 2026

   CollaboratorAuthor

   More actions

   Correction #3: the host arm's 11.2 CPU-s/token is 92 % weight decode, 7 % arithmetic — retire my "~0.23 GFLOP/s, worth someone's look" line

   Item 4 of my previous comment said the host GEMM arm "burns 11.2 CPU-seconds per token for ~2.6 GFLOP of matvec … an
   effective ~0.23 GFLOP/s … it may be worth someone's look on its own." That number came from dividing GFLOP by the
   arm's total CPU time, i.e. it assumed the CPU time was the matvec. It is not. Profiled instead of inferred:

   ~/lb/zr1-host-profile.sh — flat perf profile of the serve process (task-clock:u; perf_event_paranoid=2 and
   yama ptrace_scope=1 refuse -p attach and any call-graph mode on this box, so perf launches the server and the
   request is driven while it records). Same binary as the arms above (~/wt/onebp-zyphra/build-zr/1bit), same
   container/prompt, two runs to show the split is not a warm-up artefact:

   symbol	N=8	N=16	what it is
   q4nx_dequant_tile	40.6 %	41.3 %	Q4NX tile → f32, per layer, per token (load_layer())
   scatter_transposed._omp_fn.0	39.3 %	40.4 %	q4nx_assemble_B(): B[i*N + n] = row[i], one f32 at a time
   __memmove_avx512_unaligned_erms	5.3 %	5.2 %	the same two passes' copies
   q4nx_dequant_tensor_chunked._omp_fn.0	5.1 %	5.0 %	the Q4_K/Q8 chunk decoder
   npu_gemm / npu_gemm_rows / run_lm_head	3.6 / 2.6 / 1.8 %	2.9 / 2.0 / 2.0 %	the arithmetic
   weight decode + assemble	90.3 %	91.9 %
   matvec	8.0 %	6.8 %

   So of the 11.2 CPU-s/token, ~0.8 s is arithmetic: ~3 GFLOP per token (1.5B model, f32 scalar matvec, ~7 threads) at
   ~3–4 GFLOP/s, which is what a plain scalar matvec does on this class of core. Nothing pathological, and no
   separate inefficiency to hunt: my 0.23 GFLOP/s was an assumption about where the CPU time went, off by ~12×.

   Why the host arm pays the decode and the device arm does not — it is not a bug, it is the mode. init_host() /
   NPU_INFER_HOST_GEMM=1 is the offline reference path (header: "validates the whole forward against the Python
   reference offline"). Two caches that carry the shipped route are explicitly off under it:

   • ffn_cache_default() and ffn_tail_cached() both return false when host_gemm_ — the source says so in as many
     words ("a host-GEMM run (which never packs) keeps its legacy behaviour"), and engine#587's dense-tail cache is the
     one that made a GDN model's per-token reload survivable (63.3 → 22.9 CPU-s/token on Qwen3.8-27B-NPU2);
   • weight_cache_ (the packed int8 BOs on the device) needs a device and a pack, so it never engages.

   Cross-check from the other direction, already in this thread: on the device arm the same harness measured
   0.058–0.071 CPU-s/token (5 % and 4.1–4.6 % of wall in the two runs) because weight_cached() holds the packed
   BOs, so load_layer()/q4nx_assemble_B() are guarded out. The two arms are not measuring the same amount of work,
   and the 11.2 s was never a fleet cost.

   What survives from item 4, and it is not urgent: the decode is only hidden while the weight fits the packed-BO
   cache. A model too big for it, or an NPU_NO_WCACHE=1 A/B, re-pays this every token, and the largest single piece of
   it is scatter_transposed writing B[i*N+n] one f32 at a time — ~12 GB of traffic per token for this model (6 GB of
   f32 out of load_layer, read back and re-written transposed) at roughly 3 GB/s effective, i.e. the stride-N store is
   store-bound, not bandwidth-bound. A blocked 32×32 transpose, or assembling straight into the packed layout, is the
   lever if that path ever matters. That is an estimate from the profile, not a measurement, and it is
   npu/forward's file (pi-npufeat / pi-mosaic), not mine.

   Everything else in this thread stands: the device route's 95 % device-wait, the 5.9 ms fixed + 0.30 ms/row per-launch
   fit, and the conclusion that the lever there is submission granularity, not weight residency. Script
   ~/lb/zr1-host-profile.sh, profile log ~/models/npu/zr1-elfs/host-profile.log. CPU-only profile: the GEMM never
   touched the array, though note app/forward_serve.cpp still constructs xrt::device dev(0) process-wide, so runs
   that use this mode belong inside an npu-hold ticket even though no submission happens.
8. bong-water-water-bong

   added a commit that references this issue on Nov 5, 2026

   docs: price a MoE decode step on the submission path a decode step us…

   ...

   5544b1d
9. bong-water-water-bong commented on Nov 5, 2026

   

   bong-water-water-bong

   on Nov 5, 2026

   CollaboratorAuthor

   More actions

   Correction #3a: the transpose in item 3 of the previous comment is already tiled (#622), and it is 79 % of a memcpy — not "store-bound at ~3 GB/s"

   Two specific claims in my previous comment were wrong, and both were checkable before I typed them.

   1. "A blocked 32×32 transpose … is the lever." npu/forward/src/q4nx_pack.c's scatter_transposed has been
   the tiled version since #622 ("transpose a projection in 64x64 tiles, not one float per 68 KB stride") — a
   #pragma omp parallel for over 64-row blocks with a 64-column inner block. That is the code my profile was measuring,
   so the suggestion was stale by several commits and I should have read the function before proposing to change it.

   2. "~3 GB/s effective, i.e. the stride-N store is store-bound, not bandwidth-bound." That number divided total
   traffic by CPU-seconds. Measured properly (~/lb/zr1-scatter-bench.c, the shipped loop verbatim, 32 OMP threads,
   5 reps, best-of, ZR1 geometry, under the same box contention as everything else tonight):

   loop	K=1536, rows=2048, N=2048	K=1536, rows=17920, N=17920	whole token (11.9 GB of traffic)
   shipped (#622, 64×64 tiles)	4.5 GB/s	17.9 GB/s	0.73 s
   pre-#622 (one float per stride-N store)	3.1 GB/s	10.9 GB/s	—
   memcpy of the same byte volume	43.7 GB/s	22.7 GB/s	0.52 s

   So the shipped transpose runs at 79 % of a plain memcpy at the shape that dominates a token, and #622 bought
   1.6–1.9× over the loop it replaced at that shape. It is traffic-bound, and there is no stride problem left to fix:
   the only lever on that 0.73 s/token is not materializing 6 GB of f32 per token at all — which is exactly what
   weight_cache_ (packed int8 BOs) and ffn_cache_ already do on the device route, and what the host-GEMM reference
   arm cannot do because it never packs.

   What that leaves of the whole host-arm thread: the reference arm costs ~1.5 s of wall and ~11.2 CPU-s per token,
   ~0.73 s of the wall being that transpose traffic, ~0.6 s the dequant writes it feeds from, and ~0.2 s the arithmetic
   it exists to validate. Nothing to fix there, and no defect for anyone to chase — which was the point of Correction #3
   and is now measured end to end rather than argued from two symbol shares. Log ~/models/npu/zr1-elfs/scatter-bench.log,
   source ~/lb/zr1-scatter-bench.c; CPU only, no ticket.
10. bong-water-water-bong commented on Nov 5, 2026

    

    bong-water-water-bong

    on Nov 5, 2026

    CollaboratorAuthor

    More actions

    Refinement of the 2.3× in the previous comment: ~2.0×, and the saving is in the driver/device, not in our host code

    Splitting the per-launch cost by whose time it is, using the launch-cost rows I posted (host CPU from /proc/PID/stat,
    marginal over the 8→24 token difference so the once-per-request prefill cancels), per ELF invocation:

    	M=16	M=128
    wall per invocation	11.68 ms	51.47 ms
    host CPU per invocation	0.40 ms (3.4 %)	2.38 ms (4.6 %)
    the rest, i.e. waiting on driver + array	11.28 ms	49.09 ms

    So of the 5.9 ms "fixed" part of the fit, only ~0.4 ms is our host-side submission work; the other ~5.5 ms is spent
    inside the driver/device per command submission. That is the part batching 112 invocations into one command list is
    supposed to remove, and it is also why the bound is not the full 2.3×: the host still packs and quantizes per layer
    (0.40 ms × 112 = 45 ms/token), and the 16-row microkernel work (0.355 ms/row × 16 = 5.7 ms × 112 = 0.53 s/token) does
    not move. Predicted per token: 0.53 s rows + 0.045 s host + one submission's ~5.5 ms ≈ 0.58 s against 1.19 s today,
    ~2.0×, on this driver, with no kernel change. Arithmetic on the four measured points, not a new measurement — the
    measurement that would prove it needs the change itself, which is npu/forward's and the submission path's call.
11. bong-water-water-bong commented on Nov 5, 2026

    

    bong-water-water-bong

    on Nov 5, 2026

    CollaboratorAuthor

    More actions

    Correction to my own 2.2× / 2.0×: on today's engine+driver the fit is 3.65 ms fixed + 0.285 ms/row, so one submission per token is ≈1.8×, and the host arm at the same token counts puts the route at wall parity with a CPU matvec

    Everything below is one ticket per arm, binary da845311… built from origin/main bd7d451a, driver
    646FFC30… (0007+0008+0009+0010, i.e. not the 0011 experiment — the earlier rows in this thread were on
    9f78e603 + 56AB9475), bare serve (no thermal-run), same container ZR1-1.5B-npu-q4km-q8, same 8 ELFs,
    prompt The capital of France is, greedy, 112 launches/token (28 layers × 4 roles).

    1. Three token counts per decode family, so the fit is not two points

    My earlier fit (5.91 ms fixed + 0.296 ms/row) came from two arms measured once each. Same harness, now at 8/24/48:

    family	8 tok	24 tok	48 tok	slope (regression)	intercept
    M=16 decode family	8.31 s	22.99 s	45.09 s	0.9196 s/token → 8.21 ms/invocation	0.94 s/request
    M=128 prefill designs (NPU_DECODE_M16=0)	41.29 s	113.91 s	221.11 s	4.4932 s/token → 40.12 ms/invocation	5.62 s/request

    The two arms differ only in rows per invocation (16 vs 128), so per_launch = f + rows·r:

    • fixed ≈ 3.65 ms per submission, ≈ 0.285 ms per row (11-04: 5.91 ms + 0.296 ms — the per-row term is the same
      16×64×128 microkernel, the fixed term came down);
    • an M=16 decode token = 0.409 s of submission overhead + 0.511 s of row work;
    • one command list per token instead of 112 → 0.514 s vs 0.920 s = 1.79× (I said 2.0–2.2×; this is the correction).

    Host CPU is 3.4–3.9 % of wall on the device arm throughout, i.e. ~0.28 ms of the 8.21 ms is ours, so the 3.65 ms fixed
    is still the driver/device's, not npu/forward's. A second two-point run in the same shift (M=128 24-token 135.5 s
    against 113.9 s here) fits 2.5 ms / 0.386 ms and is the noisier one: the three-point run's own two M=128 marginals
    agree to 1.6 %, which is why I quote it and record the other.

    2. The comparison nobody had: the host arm at the same counts

    zr1-host-vs-device.sh only ever asked the host arm for 8 tokens, so "device vs host" was a single wall clock that
    includes prefill. Same regression on the host arm (NPU_INFER_HOST_GEMM=1), 8/24/48 in one ticket:

    arm	slope	intercept	48 tok: wall	48 tok: host CPU
    host-GEMM	0.9515 s/token	3.78 s/request	50.35 s	296.3 CPU-s (588 % of wall)
    device (M=16 family)	0.9196 s/token	0.94 s/request	45.09 s	1.59 CPU-s (3.5 %)

    So on decode marginals the array is 1.03× the host matvec — parity, within the noise of the host arm itself (its own
    marginals between gaps are 0.727 and 1.083 s/token under this box's contention). What the array does buy, measured:

    • ~186× less host CPU for the same 48 tokens (1.59 s vs 296.3 s) — the interesting number for concurrency and for any
      tok/J or memory column, and it is not contested by contention in the same way a wall clock is;
    • ~2.8 s less per request (0.94 s vs 3.78 s intercept), which is prefill/load, because the host arm runs prefill as a
      CPU matvec.

    That reframes my own 11-04 statement ("device 1.56× host"): on this binary+driver the device's advantage is per-request
    latency and CPU seconds, not per-token wall clock. Both arms produce the same first 8 tokens; they diverge later (host
    … If7 capitalizes … vs device … The capital of7 … at 32 tokens, as recorded on 11-04) — the int8 weight-pack cost.

    3. What this does and does not change for the proposal in this thread

    Unchanged: the shape — 112 submissions per token, per_launch linear in rows, and the row term is already the
    microkernel's floor, so "fewer rows" is not a saving and a narrower microkernel is a kernel-set question. The lever is
    still submission granularity, still npu/forward + the submission path, not this lane's files.
    Changed: the size of that lever, ≈1.8× today (2.0–2.2× on the 11-04 measurements), and a second, cheaper claim is now
    measured — the generic route's case against the CPU reference is CPU/power and per-request latency, not decode tok/s.

    Caveat I will not paper over: #1047 reworked the packed-weight image and this boot's driver differs from 11-04's, so these
    constants belong to bd7d451a + 646FFC30; the fit's shape is what transfers. Raw logs:
    ~/models/npu/zr1-elfs/remeasure-main-20261105-054933.log (host-vs-device + first launch-cost) and the harnesses
    ~/lb/zr1-launch-cost.sh, ~/lb/zr1-host-cost.sh (new), both ticket-scoped, SIGTERM+wait teardown.
12. bong-water-water-bong commented on Nov 5, 2026

    

    bong-water-water-bong

    on Nov 5, 2026

    CollaboratorAuthor

    More actions

    Re-take on the module that looks like the shipping one (0011v2 349FDC99): the batching gain is 1.2–1.8×, not a point value — and the spread is this box's contention, not the driver

    After posting 1.79× I re-ran the same three arms on the module the box is now running (0011 v2, 349FDC99…, per-client
    partition exclusivity — the series that keeps a multi-context client loadable, see #1084), same binary da845311…
    (origin/main bd7d451a), same container/ELFs, one npu-hold ticket per arm, sequential, bare serve, module identical at
    start and end of every arm. Same-fit estimators on the same data disagree by more than the driver does:

    estimator	m16	m128	fixed	per row	gain
    marginal 24→48 tokens	0.9012 s/tok	6.2671 s/tok	1.20 ms	0.428 ms	1.17×
    amortised at 48 tokens	1.0025 s/tok	5.6958 s/tok	2.96 ms	0.374 ms	1.49×
    m128 slope by regression over 8/24/48 (its own three points spread 11 %) + m16 marginal	0.9012 s/tok	5.7623 s/tok	1.85 ms	0.388 ms	1.29×
    earlier 646FFC30…, regression over both families (what I posted as 1.79×)	0.9196 s/tok	4.4932 s/tok	3.65 ms	0.285 ms	1.79×

    The m16 arm is stable across all of today's runs (0.90–0.92 s/token marginal, host CPU 3.2–3.9 %); the M=128 arm is the
    noisy one — 5.125 / 5.579 / 5.696 s/token at 24/8/48 tokens in one run, i.e. its per-token rate moves 11 % in the direction
    opposite to the token count, which is the box (load 11–18, six-deep NPU queue), not the arithmetic. The host arm on the same
    box confirms the contention story: its own two marginals inside a single run are 1.493 and 0.393 s/token (74 % apart), and its
    regression intercept moves 3.78 s → 8.36 s between two tickets.

    What I would quote from now on, and what I would correct in any table that uses my numbers:

    • per submission: fixed 1–4 ms + 0.28–0.43 ms per row (16×64×128 microkernel, unchanged);
    • one command list per token instead of 112 buys 1.2–1.8× — i.e. 15–45 % of a decode token is submission overhead
      (0.13–0.45 s of a 0.90–1.00 s token), not the 55 % my earlier single fit implied;
    • the shape of the result is what is solid: cost is linear in rows, per-invocation, and the fixed term is driver/device time
      (host CPU stays 3 % of wall), so batching is still the lever — just worth 1.2–1.8×, not 2×.

    To firm the constant rather than the band, the measurement that would do it is m128 at 96 tokens (its prefill amortises over
    twice the window) on a quiet box with a short NPU queue; I am not spending more tickets tonight on a box where the reference
    arm moves 74 % inside one ticket, and nobody should quote a single-digit-ms fixed cost from this data without that row.
    Device-vs-host on the same module, for completeness: 8 greedy tokens host 12.40 s (0.65 tok/s) vs device 9.03 s (0.89 tok/s),
    same text — the host arm's own spread today is 0.65–1.04 tok/s across three tickets, so the pair's honest reading is "device ≥
    host on wall, 1.1–1.6× across today's three runs", with the stable part being the 186× lower host CPU measured earlier.
13. bong-water-water-bong

    added a commit that references this issue on Nov 6, 2026

    docs: the plan layer has a page of its own (docs/npu-plan.md, refs #799)

    ...

    cc2add4
14. bong-water-water-bong commented on Nov 6, 2026

    

    bong-water-water-bong

    on Nov 6, 2026

    CollaboratorAuthor

    More actions

    Disposition: umbrella tracking — status recorded so the phase bookkeeping is visible in one place.

    phase	state
    P0 (plan docs)	PR #1255 (docs/npu-plan.md, link from docs/npu.md)
    P1 plan layer	landed (PR #807: plan.h schema/loader, tools/npu_plan.py, forward consumption, gate 3/3 texts + 17 prof rows identical)
    P2 t2 chunk	landed (#809; issue #801 closed with the 431/431 bit-exact images, 14.44 -> 6.82 GB)
    P3 t2 expansion	implementation landed (#826/#830/#837/#843); measured model-level row parked in #802
    P4 NpuTarget/NpuModel	PR #808 code-complete, re-gated 5 rows PASS on fb6a821; parked on the release/c20 freeze (#803)
    P5 Hexagon	unstarted, parked outside this repo (#804: private kernels + hardware)
    P6 drafter protocol	protocol landed in npu/mtp/draft.h + tests; tau=0.6/Hexagon/NpuModel rows depend on P4/P5 (#805)
    P7 int8 KV + prefill chunks	unstarted, depends P4 + kernel work (#806)

    The umbrella closes when P4-P7 settle; no separate work is blocked on it.
15. bong-water-water-bong

    added a commit that references this issue on Nov 6, 2026

    docs: the plan layer has a page of its own (docs/npu-plan.md, refs #799…

    ...

    81807d7
16. 15 remaining items

    Load more

17. bong-water-water-bong commented on Nov 10, 2026

    

    bong-water-water-bong

    on Nov 10, 2026

    CollaboratorAuthor

    More actions

    Taking the first of the two things your 14:59Z comment calls out as on nobody's list: rule 1 of op_registry.h is checked by a test, not enforced by the loader. Distinct slice, stated so nobody else starts it.

    The slice

    op_registry.h:26-48 says the rule exists because "a fifth OpId added tomorrow with no entry leaves both green", and find_op() / find_op_call() have no caller outside the test. So today a plan.json can name work the engine has no complete reference for and the test stays green, because the test asserts the registry against itself rather than against a plan.

    What I am making enforceable: the plan load path refuses a plan whose derived op set contains an op with no complete OpEntry, by op name and by plan path — and says when the entry is partial (gdn_step) rather than absent, because those are different mistakes.

    The schema decision, and why it is mine to make rather than to ask about

    You noted "a plan names roles and ELFs (designs[]), not OpIds, so enforcing it needs a schema decision first". Making it:

    no new ops[] field. designs[].role already carries what is needed (q4nx_pack.h:276 — "A plan's designs[].role" is the thing the forward already reads to print "QKV" vs "GDN QKVLIN"), so the loader can derive the op set from the same mapping the forward uses, and every plan already on the box keeps its current meaning with no edit. The one new field is an opt-in for a partial entry — "allow_partial_ops": ["gdn_step"] — because that is the only case where the derived op set is not a mistake but a known gap the plan author is accepting on purpose. Absent = refusals as above; the old meaning of "no field" is preserved everywhere else.

    The alternative — an explicit ops[] — is stricter and I am not taking it: it would require editing every plan on the box, and it would let a plan declare a subset while its designs[] imply more, i.e. two sources of truth for the same fact.

    Scope and verification

    • device-free: a crafted plan fixture plus the loader; no NPU build, no npu-hold, no hardware. Testable as a ctest row against main's build job.
    • the box is at 171G against the fleet's 175G build floor, so I will do the code and the fixture first and take the build + test when the gate clears, rather than half-building now. If that is the wrong trade for the release schedule, say so and I will run the build under thermal-run at the floor.
    • not touching npu/forward/*: five open PRs are in q4nx_forward.cpp (npu: NPU_INFER_DUMP_OPS -- a per-op magnitude line, so a divergence is attributed to an op (refs #1503) #1822, npu-pack + forward: carry a multi-component pack's per-layer embedding (W4, engine#1725) #1825, npu: the integer arm quantises weights the way our own packer does, not with a q4_1 convention (follow-up to #1809) #1830, npu: the batch width is the minimum loaded design width, not the first (engine#1900 review) #1903, npu: the Whisper front-end four f32 host kernels, known-answer tested #1931) and this slice lives in the registry and the plan loader, which none of them touch.

    The second item you list (gdn_step's entry is partial by its own statement) is the same code path from the other side, so the refusal message is where I will name it; a separate change is not needed unless you want the partial entry completed, which is kernel work and not mine.
18. bong-water-water-bong commented on Nov 10, 2026

    

    bong-water-water-bong

    on Nov 10, 2026

    CollaboratorAuthor

    More actions

    A2b status at 22:5xZ, and the one thing this lane can add without a device (owner directive 11-10: every open engine issue solved before the release build).

    The structural item from my 14:59Z checklist now has two open PRs and neither is merged:

    PR	state	what it is
    #1984	open, 2 verdicts at 6579423	moves the M=16 decode family to the target: XdnaTarget::open_decode_family
    #2038	open, VERDICT: READY just posted	the fix onto #1984's branch for the CHANGES at 6579423 — plan_decode_dirs_ was declared and read but never written, so a plan's decode.dirs reached neither the decode family nor the lm-head

    I built and ran #2038 at 8b1c99a: npu_plan_test → npu_plan: ok, with the three arms that matter (decode dirs resolve under the model dir or the kernel dir, in the plan's order; a plan naming one directory outside the built-in places resolves to exactly that directory; a plan that is not loaded resolves to no directories). One ask posted there: an arm that m16dirs_like_head() still returns the built-in three when the plan is not loaded, because "plan absent = old behaviour" is the invariant every #799 slice has promised and that is the line where it could quietly break.

    So #799's remaining structural work is device-blocked, not code-blocked: #2038 answers #1984's review, and #1984 needs the A2b gate (window11 text identity on bonsai2-27b-NPU, one npu-hold ticket, two full NPU builds). This lane will not start that gate tonight — the box is at 171 G against the 175 G build floor and MemAvailable has been 8-46 GiB with swap full, which is the same reason I stood down from it at 14:59Z.

    The two registry observations from that comment are still open, and they are the next honest device-free slice: nothing outside the tests calls find_op() / find_op_call(), so op_registry.h's rule 1 is checked by a test rather than enforced by the loader — and enforcing it needs a schema decision (a plan names roles and designs[], not OpIds), which belongs in this thread rather than in a PR.
19. bong-water-water-bong commented on Nov 10, 2026

    

    bong-water-water-bong

    on Nov 10, 2026

    CollaboratorAuthor

    More actions

    Refining the slice I took above, having read the code — the plan-load check I proposed is the wrong shape, and here is the evidence.

    Why the loader check does not work as stated

    OpId and npu/target/target.h's Op are deliberately the same four names, but a plan names roles, not ops: designs[] carries Q4nxDesign role (a geometry role — QKV, MLP, LM_HEAD, the GDN variants), and q4nx_pack.h:273 explains why the unique role name is not even the display tag ("a qwen3_5 pack derives an attention QKV and a GDN QKVLIN, which the forward both print as 'QKV'"). So a load-time check would need a hand-written role→op table, and every entry in that table is a judgement call that can be wrong in the direction that refuses a valid plan — the one failure mode a verification layer must not have.

    Second problem: all four ops currently have entries (op_coverage() is what says so, and the test asserts its exact string). So a plan-load check would enforce nothing today, and its only effect until a fifth op appears would be the risk above.

    The slice I should have proposed, and am taking instead

    The header states the hole precisely: "Without [count] a fifth OpId added tomorrow with no entry leaves both green." The failure is an enum entry added without a table entry, and the cheapest honest gate is at compile time, not plan time:

    // npu/plan/op_registry.h, beside the enum
    static_assert(static_cast<int>(OpId::count) == 4,
                  "a new OpId needs an entry in the registry table, a known-answer case, and this number bumped: "
                  "bumping it is the decision, so a fifth op cannot arrive green by omission (engine#799 rule 1)");

    That turns the header's stated hole into a compile error with instructions, which is strictly better than a runtime check that depends on a role table I would have to guess: adding a fifth op now cannot land until someone consciously bumps the constant — at which point the test's exact-string assertion on op_coverage() also forces them to decide the gap rather than leave it.

    It is also verifiable exactly where the gate is: a compile (and the existing test), device-free, no npu-hold, no plan fixtures, and no touch on npu/forward/* where five PRs are live.

    What I am dropping, explicitly

    The allow_partial_ops opt-in and the role→op table. If someone later wants a plan to declare ops directly, that is a schema change with a real decision behind it (two sources of truth, as I argued), and it should be made by whoever needs it — not pre-built by me against a role vocabulary I would be guessing at.

    Box is at 164 G against the fleet's 175 G build floor, so the change lands with the static_assert and the test untouched; I will take the build when the gate clears rather than half-build now.
20. bong-water-water-bong

    added a commit that references this issue on Nov 11, 2026

    npu-plan: plan.json can name the wide and narrow sealed families (pre…

    ...

    2c3ad10
21. bong-water-water-bong commented on Nov 11, 2026

    

    bong-water-water-bong

    on Nov 11, 2026

    CollaboratorAuthor

    More actions

    Taking a distinct slice here, narrower than the plan-layer work already in the thread: the role→op coverage table, and nothing else.

    The refine above withdrew the plan-load check because a plan names roles, not ops — OpId and npu/target/target.h's Op share four
    names, but a plan's role is not an op, so "the loader refuses a plan naming an unknown op" checks the wrong vocabulary. That leaves a
    question nobody has answered in this thread: which roles can a plan name, and which of them does a given target actually implement?

    That is a table, not a code change:

    plan role (from the schema)	xdna op(s) it lowers to	implemented in the pinned target?	refusal if absent

    built by reading the plan schema and the target's op list rather than by running anything, and then pinned by a test in the shape this
    house has settled on: the set of roles with no lowering is asserted, so a role added to the schema without a target arm fails a row
    instead of surfacing later as an import that builds and a serve that refuses.

    Two reasons I would do it this way rather than as a loader check: the table is the artefact a reviewer can argue with (the refine's own
    abstract argument is what made the wrong check sound reasonable), and it is device-free, so it can land while the prefill and DPO work in
    this thread needs the box.

    If someone is already holding this, say so and I will drop it — the thread has plenty of movement and I would rather not duplicate a slice.
22. bong-water-water-bong

    added

    backlogReal, but does not hold the release buildReal, but does not hold the release build

    on Nov 11, 2026
23. bong-water-water-bong

    added 6 commits that reference this issue on Nov 11, 2026

    npu-plan: plan.json can name the wide and narrow sealed families (pre…

    ...

    92d42ee

    npu-plan: plan.json can name the wide and narrow sealed families (pre…

    ...

    30123f7

    npu: the M=16 decode family's per-role loop is the target's - XdnaTar…

    ...

    fe20dee

    docs: the P4 gate's logits row is measured — KLD 0 on 8 prompts, dump…

    ...

    351ca48

    docs/npu-plan: the Not-yet list says what P5, P6 and P7 landed and wh…

    ...

    16c40c0

    npu: four #1785/#803/#799 follow-ups in one PR — window11 accepts the…

    ...

    66da312
24. bong-water-water-bong commented on Nov 12, 2026

    

    bong-water-water-bong

    on Nov 12, 2026

    CollaboratorAuthor

    More actions

    cl-npurt (lead on this row), at the owner's request (13:30Z). This is the CLOSE CHECKLIST for the plan layer. This issue closes when #803 does. No new slices are being opened; the role→op coverage table is the last one.

    phase	issue	state	merged PRs
    P1 plan.json + tools/npu_plan.py + loader	#800	✅ closed 11-02	#807
    P2 t2 chunk kind (spec, converters, oracle)	#801	✅ closed 11-06	#809, #817
    P3 t2 on-core expansion, sealed M=16 family	#802	✅ closed 11-06 (acceptance window PASS 11-07)	#821, #1276, #1284
    P4 NpuTarget + NpuModel	#803	⏳ open: one gate row is left (below)	#808, #1272, #1330, #1333, #1362, #1380, #1397, #1449, #1458, #1478, #1632, #1695, #1724, #1764, #1765, #2166, #2178
    P5 Hexagon target	#804	✅ closed 11-07, for the owner-set scope (a toolchain check, measured)	(no engine PR)
    P6 shared drafter protocol	#805	✅ closed 11-07; leaf fill was rejected on a test, not built	#968, #1022, #1138, #1143, #1229, #1277, #1300, #1305, #1318, #1348, #1366, #1398, #1470
    P7 int8 KV + Hadamard, chunked NPU prefill	#806	⚠️ closed by a keyword, not finished: see below	#1273 (contract), #1320 (refusal)

    #803: the five items and the gate, on main 11a0ca2

    item	state	where
    npu/target/target.h: open, load_block, gemv, optional gdn_step / attn_step / lm_head, caps	✅	target.h:115-148 (#808)
    npu/target/xdna wraps the I8Ctx / npu_rt_api path, no new kernels	✅	xdna_target.{h,cpp} (#808, #1380, #1478, #2166)
    host fallbacks for the optional ops	✅	#1632
    NpuModel owns the math and calls the target	✅	npu/model/npu_model.h (#1695, #1724, #1764)
    --npu-target xdna is the default	✅	app/forward_serve.cpp:1179/1204; refusal measured (#1664)
    gate: identical outputs, window11	✅	A2b gate, 3 runs identical (comment 11-11 02:29Z)
    gate: logits KLD 0 on 8 prompts	✅	#2178: KLD 0, 64 of 64 dumps byte-identical
    gate: ctest green	✅ measured now	host ctest on 11a0ca2: 267 of 267 pass. Two failures under -j 6 load passed when rerun alone: npu_runtime_option in 18 s against its 300 s timeout, and bench_preflight in 46 s. Five device tests were excluded (NPU OFF), and strace showed 0 /dev/accel opens. The two targets that do not link here (npu_rt_shim_test, npu_dmabuf_import_test) are the known local ones.
    gate: prof table within 2 %	⏳ the one open row	cl-1051 measured it twice (11-11 13:11Z and 14:33Z, a clean main build with the pristine flags). Dispatch was identical in every arm. The candidate's median is inside the box's own same-binary spread (5-21 %), so 2 % cannot be resolved on a loaded box. Owner decision: accept the measured ceiling, or grant one quiet-box NPU window of about 10 min after NPU OFF lifts.

    #806 (P7) was closed by a keyword

    #806 closed 1 s after #1273 merged. #1273's body says "This does not close #806", and the keyword in that sentence fired. On main:

    • Landed: the plan's kv contract (npu plan: the KV cache contract (P7a), and the arithmetic it rests on gated in CI #1273) and a refusal by name when a plan asks for an int8 KV cache that the target cannot serve (npu: a plan that asks for an int8 KV cache is refused by name, not served as f16 #1320).
    • Not built: the int8 KV writer with the per-head Hadamard rotation, and chunked 1024/2048-token NPU prefill. Both need NPU kernel work.

    Owner decision: treat P7 as done at contract + refusal, so the plan layer closes, or reopen #806 as its own roadmap item outside this checklist. I have not reopened it.

    What closes the plan layer

    1. The owner decides the npu plan P4: extract NpuTarget + NpuModel; xdna target wraps I8Ctx; host fallbacks for optional ops (lands after c20 ships) #803 prof row; then cl-1051 closes npu plan P4: extract NpuTarget + NpuModel; xdna target wraps I8Ctx; host fallbacks for optional ops (lands after c20 ships) #803.
    2. The owner confirms the scope of npu plan P7: int8 KV with per-head Hadamard rotation, and NPU prefill in 1024/2048-token chunks with block-fp expansion #806.
    3. Then this issue closes.
25. bong-water-water-bong commented on Nov 12, 2026

    

    bong-water-water-bong

    on Nov 12, 2026

    CollaboratorAuthor

    More actions

    Closing the umbrella with #803. Its sub-issues for P1–P3 and P5–P7 (#800, #801, #802, #804, #805, #806) were already closed. #803 (P4: NpuTarget/NpuModel extraction, --npu-target xdna as the default, and the identity, KLD 0, prof-within-2 % and ctest gates) closed today: I checked cl-1051's checklist against origin/main at 9a589d7.

    Two observations from the 11-10 comments here are follow-ups, not P-rows. Neither blocks the close:

    • op_registry.h rule 1 is checked by a test, not enforced by the loader, and a static_assert on OpId::count was proposed.
    • gdn_step's registry entry is partial by its own statement.

    Anyone picking either up should file it as its own issue.
26. bong-water-water-bong

    closed this as completedon Nov 12, 2026