npu: load Ryzen AI OGA NPU model packs on Linux through the engine (pack import -> Q4NX -> our kernels)

Goal (owner, 2026-10-30): a Ryzen AI OGA NPU model pack (*_rai_*_npu_*: model.onnx + external data, genai_config.json, descriptors) dropped on a Linux box runs on the NPU through 1bit serve with the engine's own runtime and kernels. The pack supplies the weights, tokenizer and config; the engine supplies everything else. No vendor wheel or shared object is linked, shipped or required, and no vendor-compiled NPU program is executed.

Plan:

1bit npu-pack import <pack-dir> -o <dir>: parse the pack's quantized weights (MatMulNBits tensors, documented com.microsoft contrib op) into the Q4NX container with a layout-version field; self-check dequantises sampled tensors against the reference values bit for bit.
Per-shape full ELFs for the pack model's geometry from our generator (decode M=16 twins, prefill M=128), bit-exact against the CPU reference.
Gates on the first pack (Qwen2.5-0.5B-Instruct npu_4K): greedy output identical to the vendor stack on three prompts, 12-chunk KLD of our logits vs the vendor stack, decode tok/s at ctx 128/512/2048 and prefill at 128/2048 side by side with the vendor stack measured the same hour, medians of 3.
A second pack of a different family from the public collection; unsupported pack features filed as issues.
Later: the ORT execution provider (onnx: ONNX Runtime in-process behind -DONEBIT_ORT, with our plugin EP #647) claims the pack graph; out of scope here.

Related: #647 (ORT EP), #640 (two NPU engines), #650-#653 (NPU feature gaps).

🤖 Generated with Claude Code

Triage: enhancement. Item 1 is genuinely not started at origin/main 31d52e1:

$ git grep -ln 'genai_config\|MatMulNBits\|npu-pack' origin/main -- app/ npu/ tools/ docs/   -> no file

Lane: pi-ogaload. Two facts that bound the first PR:

Items 1 and 2 are separable, and item 1 is the one that can land from this box. The importer reads the
pack's model.onnx initializers plus genai_config.json and writes Q4NX with a layout-version field; the
self-check (sample tensors dequantised against the documented com.microsoft MatMulNBits formula,
bit-exact) is the test, and it needs no device. Item 2's per-shape full ELFs (decode M=16 twins, prefill
M=128) are device work and follow.
onnx: ONNX Runtime in-process behind -DONEBIT_ORT, with our plugin EP #647 is not a substitute for this. 8b6895e lands ONNX Runtime in-process behind -DONEBIT_ORT with
our plugin EP — that claims the pack graph. This issue claims the weights by re-packing them into our
container, which is what makes a pack runnable on the engine's own kernels today, with no vendor wheel,
shared object or vendor-compiled program anywhere in the path. Keeping them separate is what keeps that
constraint checkable.

One thing I would add to the plan while it is still cheap: a pack that the importer cannot represent
(supported pack features are a moving target) should exit with a named refusal listing the unsupported
attribute, not a partial conversion — item 4 says unsupported features become issues, and that only works if
the importer names them.

npu-pack: say when there is no reference, instead of checking against…

Item 4, first pass (second family). amd/Llama-3.2-1B-Instruct_rai_1.8.0_npu_4K — a different
architecture (llama, hidden 2048, intermediate 8192, 32 heads / 8 KV heads, head_dim 64, vocab
128256, rope_theta 500000, 16 layers) — on branch 1bit/671-oga-pack-import:

$ 1bit npu-pack info ~/models/oga/Llama-3.2-1B-Instruct_rai_1.8.0_npu_4K
  dims: {"hidden_size": 2048, "intermediate_size": 8192, "num_attention_heads": 32,
         "num_key_value_heads": 8, "head_dim": 64, "vocab_size": 128256, "rope_theta": 500000.0,
         "model_type": "llama", "num_hidden_layers": 16}
  tensors: 147   incomplete: 145   (2 x model.onnx.data, 145 x fconst)

$ 1bit npu-pack import ... -o /tmp/llama-import --allow-incomplete
imported 2 tensors, 32064 q4_1 tiles into /tmp/llama-import
  self-check: lm_head vs the pack's fp16 embedding  corr 0.999999  rel_err 0.000701
  controls:   nibble-swapped corr 0.0004, no zero-point corr 0.8586 (both must fail)
  container:  read back through q4nx::dequant_tile  rel_err 3.343e-03  max |diff| 4.883e-04
  parameters: bf16 scale/zero exact for 43.09% of 32-column groups (format limit)
  incomplete: 145 tensors, listed in /tmp/llama-import/model.q4nx

What that establishes:

The reader is family-generic: geometry comes from config.json, the head MatMulNBits (N=128256,
K=2048, bits 4, block 128) decodes to the same fidelity as on the Qwen2.5-0.5B pack — corr 0.999999,
rel 7.0e-4, controls rejected — so the decoder and the tiling are not tuned to one pack. 2,048 is a
whole number of 256-column tiles, so this pack needs no K padding at all: 4008 x 8 = 32,064 tiles
exactly, 658 MB.
The container's bf16-parameter limit reproduces across packs (43.1 % of groups exact vs 43.5 % on the
0.5B, rel 3.34e-3 for both), i.e. it is a property of the quantiser + chunk format, not of one model.
145 refusals, of which 33 are a new reason this pack exposed: the .fconst layer-block stride and
both norm offsets are hidden-dependent (the attention payload is sum of N/64 tiles), so the constants
measured on hidden 896 do not carry over. Filed as npu-pack: layer norms (and the final norm) are located only for hidden 896 #734 with the derivation and the per-op N/K that
the pack's own cache/*_meta.json declares.
Untied heads: the self-check only means something when tie_word_embeddings is true (it is, in every
1.8.0 pack checked). The importer now reads that flag and prints "skipped — no independent reference"
rather than a correlation against an unrelated tensor.

Still blocked for items 2-3: the .fconst weight payload is requantised, so it cannot be matched against
the pack's ONNX tensors under any packing (amd-oss/npu-re docs/FACTS-oga-pack-weight-layout.md §7.5,
§7.6 — the latter documents the vendor's own tile machinery, incl. 2,176-byte subvolumes and the
wh_w8_index code order, read from the C++ headers the wheel ships). Reproducing the payload therefore
needs the writer to run on weights we choose, not more reading.

Item 4's "repeat 3-5" (kernels + gates on a second family) still waits on items 2-3.

npu-pack: import a Ryzen AI OGA NPU pack into a Q4NX model directory (#…

npu-pack: say when there is no reference, instead of checking against…

npu-pack: import a Ryzen AI OGA NPU pack into our container (#671 ite…

npu: derive GEMM K and N on the container's and the kernel's grids (e…

npu: derive GEMM K and N on the container's and the kernel's grids, s…

npu forward: drive the forge q4_1 GEMV ELF family (the pack's own ker…

npu forward: read the batched prefill's GEMM outputs at the kernel's …

Step 5 progress: the OGA 0.5B pack runs on the array through our kernels — with measured gate numbers

This is the first run of an imported Ryzen AI OGA pack on the NPU without any AMD kernel, xclbin, TXN blob or DX4m
tile: standard engine designs, emitted from source on this box, @seq instance, our runtime.

The move that unblocked it: a different kernel family, not a different licence answer

The gate-(a) blocker written up on #855 was the forge gemv family: its kernel wants DX4m tiles, so the launcher's
container bytes are unreadable to it, and the tile format lives in the private (Binary-Licence) tree. But the container
itself declares q4nx_weight_layout: rm ([forward] weight layout: rm row-major), and the engine has a device-proven
family for that layout — the one every other NPU model here runs. Emitting the pack's shapes in it
(~/ogaload/oga4/emit_i8_elfs.sh, QKV 1536×1024, O 1024×1024, GU 9728×1024, D 1024×4864, M=128 + _m16) took the
DX4m question out of the correctness path entirely. No LMHEAD ELF is needed either: on this family run_lm_head()
reuses the widest K-grid design and tiles the vocabulary (ntiles = ceil(NV/T) = 16).

(a) greedy agreement vs AMD's stack on the same pack

Reference = ~/ogaload/window/oracle.json, captured 11-02 through OGA on AMD's own stack. Same chat template on both
sides, greedy, one token budget per oracle row.

prompt	AMD's stack	our array + int8 kernels	our container, host f32 (no kernel)
capital of France	`The capital of France is Paris.`	identical	identical
first three primes	`… are:\n\n- 2\n- 3\n- 5`	`… are 2, 3, and 5.`	identical
why the sky is blue	`… scattering of sunlight by tiny water droplets …`	different (fluent, on-topic)	different (fluent, on-topic)

The container is right: on the host in f32 it reproduces AMD byte for byte on 2 of 3, including the exact list
formatting — weights, dequantisation, rope, chat template, EOS and tokeniser are correct. The array's extra mismatch is
the int8 GEMM's per-row activation quantisation (AMD's program is gemmbfp16), i.e. the one thing the gemv family has
in bf16. Prompt 3 diverges in f32 too, at a genuine near-tie.

(c) speed, and one lever already taken

Prefix cache off (--prefix-slots 0), prefill/decode separated by a max_tokens=1 vs max_tokens=33 pair, medians of 3:

prompt	decode (rm)	decode (b8)	prefill (rm)	prefill (b8)
~147 tok	2.73 tok/s	6.13 tok/s	94 tok/s	146 tok/s
~544 tok	2.38 tok/s	6.02 tok/s	116 tok/s	155 tok/s
~2 133 tok	2.48 tok/s	2.46 tok/s	79 tok/s	57 tok/s

AMD's recorded figures for this pack are 72-77 tok/s decode / 2.5-5.6k tok/s prefill (11-02, not re-measured here).
b8 is the blocked-B family (_b8/_b8_m16 ELFs + q4nx_weight_layout: "b8", no re-import: packB writes the image
host-side) and it lands where the engine's own numbers predict, 3.3 → 8.25 GB/s. At 2 133 tokens the two layouts are
identical, so the long-context ceiling is attention/KV, not the weight stream.

What this leaves, in the order I would take it

gemv family — the only one with AMD's bf16 activation contract and the per-column stream built for these packs
(55-59 GB/s on the arena probe). It needs the DX4m tile relayout, whose implementation exists only in the private tree:
that is the licence question on my state file, and it now has a measured price (1/3 → possibly 3/3 token agreement,
6 → nearer 70 tok/s).
b8 for OGA packs by default — the importer has no q4nx_weight_layout option today; declaring b8 there is a
small change in this lane, and it is worth 2.4× on decode for every pack.
M=16 core count — the arena finding F1: at M=16 the design can place only 8 of 32 cores; partitioning N across core
rows is a generator change.

Device discipline for all of the above: every run went through npu-hold tickets, alone on /dev/accel/accel0, with
0 amdxdna errors and 0 IO_PAGE_FAULT in every ticket, and :18097 untouched throughout. Gate (b) (KLD over 12
chunks) is next and does not need the array — the f32 reference above is its vehicle.

The first speed row in AMD's own units — the 2026-10-30 decision finally has an instrument

On 2026-10-30 this lane adopted a reporting rule for step 5: "speed numbers to be reported with ort_run's AMD
model_benchmark protocol definitions (engine #683): prompt processing / token generation / E2E / peak RSS, -l 2048..128,
-g 128, -r 5, so our NPU rows sit in the same units as AMD's published rows." Every row since was taken with this
lane's delta-method instrument instead — sound numbers, wrong units for a side-by-side.

Why it never happened: one missing field. onnx/bench_npu_serve.py is the harness that reproduces AMD's layout, and
it reads timings.prompt_ms from the reply. app/unified.cpp emits that; app/forward_serve.cpp — the route every
imported pack container runs on — did not, so --mode json stopped with SystemExit: no timings in the reply.

PR #1113 (72f2c00) adds it: CallStats carries prompt_ms/predicted_ms, the clock closes where the first token's
distribution exists and again at the end of generation, and the reply carries timings {prompt_ms, predicted_ms, predicted_per_second} in the same shape unified.cpp uses. Additive only — sampling, stops, logprobs, prefix cache and
the existing usage object are untouched.

The row, array, ogaload-amdproto2 (-t 40, queued 4 s, held 408 s, rc 0, die 66 C, 8/8 b8 ELFs, amdxdna ERROR 0 -> 0, --prefix-slots 0):

prompt tokens	prompt processing tok/s	token generation tok/s	E2E ms
2678	78.606	5.1932	37149.67
1335	118.550	5.5143	14162.65
672	150.811	5.7255	7250.42
332	172.050	5.8378	4670.45
162	139.929	5.8966	3871.16

stddev 0.08 % on prompt processing and ≤1.2 ms/token on generation, so the five repetitions agree. Prompt tokens are
~1.3x the --lens value (the harness estimates words per token for its filler sentence); the protocol reads the real
count from usage, so the table is keyed on tokens.

Against AMD's recorded figures for the same pack (72-77 tok/s generation, 2.5-5.6k prompt processing — recorded
11-02, not same-hour): generation 12.2-14.8x off, prompt processing 14.5-71x off. And the two instruments agree
where they overlap (5.73-5.90 here vs 5.56-6.03 delta-method), so the gap is the engine's.

Two process notes for anyone repeating this: the box had recovered to MemAvailable 61.7 GB / Cached 47.5 GB for this
row (the 4B row earlier today was taken with 12 GB cache and is not comparable), and --prefix-slots 0 matters — a prefix
hit makes a prefill look instant, which is exactly the fault that wasted this lane's first gate (c) attempt.
Hub: docs/FACTS-npu-decode-dispatch-latency.md (5d67ba3).

Gate (c) for the 4B, finally with numbers that can be quoted — and the row that still cannot

The 4B row has been the one number this lane could not stand behind: measured as 1.55 / 1.38 / 1.00 / 0.84 / 0.42 tok/s across
five attempts in 20 hours on identical artifacts. It is now measured one context per ticket, short enough for its 3.07 GiB
container to stay resident, with the instrument recording cache state and the effective weight rate so residency is evidence
rather than assumption:

ctx	prefill tok/s	decode tok/s	weight rate	spread (prefill / decode)	quotable
128	53.5	1.51	4.63 GB/s	21 % / 3 %	yes (median of 3)
512	35.3	1.45	4.43 GB/s	2 % / 0.2 %	yes (median of 3)
2048	6.4	0.81	2.00 GB/s	26 % / 68 %	no — partial residency

The weight rate is decode tok/s × container GiB. 4.4-4.7 GB/s is cache; 1.3 GB/s is this box's measured disk speed (a
3 GiB file reads in 2.35 s); 2.0 GB/s is a container that is half reclaimed. That single number explains all five earlier
readings, and it is why the two short rows are quotable and the long one is not — nothing about the model or the kernels
changed.

Per-dispatch from the quotable rows: 4.60 ms and 4.79 ms (144 dispatches/token), consistent with the 4.2-4.48 ms this
model has shown since 2026-11-04, and with the 2026-11-04 two-run row (57.3 / 1.55) within noise.

Two instrument defects found while getting these, both now fixed in the lane's script:

The wall-clock delta method has no signal at long context. It derives the rate as (n-1)/(t_32tok − t_1tok), i.e. the
difference of two 200-320 s numbers; at ctx 2048 one attempt produced −9.56 tok/s (the 32-token request finished
before its own 1-token request) and another disagreed with itself (274 s vs 461 s for the same request). The fix is the
field that merged today: forward-serve now reports timings.prompt_ms/predicted_ms, measured inside one request, so
the script reads them and prints both readings — server stages primary for long contexts, delta method as the cross-check
(they agree to 0.7 % at ctx 64 on a quiet box).
A row with no weight rate cannot distinguish "not resident" from "slow", which is exactly how five contradictory 4B
numbers got recorded. Every row now carries container size, MemAvailable/Cached/SwapUsed, serve RSS and the weight
rate.

Still open: ctx 2048 needs the box quiet for its whole duration (~20 min at resident speed). Its four attempts never got
there (weight rates 1.30 and 2.00 GB/s), so the row is unmeasured rather than bad — it is with the owner as a scheduling
item, and it will state its own residency when taken.

Hub: docs/FACTS-npu-decode-dispatch-latency.md (d11c762).

Where the time actually goes: per-role profile of the Qwen3-4B b8 path (2026-11-05)

Brief step 4's last open item was per-ELF rows for the 4B set. forge_bench cannot produce them — it accepts Q4NX 32x256
tiles and repacks to DX4/DX4m only (forge_bench.cpp:9-19,215; no b8/blocked anywhere in the file), while a _b8 ELF
walks 8 KiB blocked microtiles (npu/forward/include/q4nx_weight_layout.h). So I used the engine's own attribution instead,
which measures the real call pattern: NPU_PROFILE=1 NPU_PROFILE_TOKEN=1 + NPU_PHASE_PROF=1
(instrument: ~/ogaload/oga11/prof_roles.sh, host-mode dry run included).

Decode (M=16), median of per-token dumps, load1 21, ctx 128

role	N x K	launches/token	ms/launch (device)	MiB	effective GB/s
QKV	6144 x 2560	36	2.642	8.0	2.8
O	2560 x 4096	36	1.890	5.3	2.6
GU	19456 x 2560	36	6.568	25.2	3.6
D	2560 x 9728	36	3.610	12.6	3.2
lm_head	151936 x 2560	1	5.846	197.1	31.3

Device budget 559 ms/token vs a measured 626-635 ms wall ⇒ the host path is ~70 ms (11 %). Two consequences:

the 4B is device-bound at decode; "96 host-side submissions per token" is a load-sensitivity mechanism, not the
throughput ceiling (corrects an attribution we had been carrying).
lm_head streams 197 MiB at 31.3 GB/s where the four layer GEMMs do 2.6-3.6 GB/s — same token, same layout, same M.
At the head's rate the 1.84 GiB of layer weights would take ~60 ms, not 553. That 9-12x is the per-launch weight-stream
gap, and the existence proof that the number is reachable is inside our own stack.

Prefill at ctx 2048, per 128-token chunk ([bs] decomposition; 16 chunks = the measured 210.7 s):

chunk start	wall s	device `wait` s	host s	`[lh]`
1408	22.18	2.20	19.59	`attention 98.9 % of gemm-free host`
1536	27.79	2.10	25.46
1664	30.39	2.14	28.04
1792	27.82	2.17	25.43
1920	33.58	2.08	31.26	`attention 30920 ms / ACCOUNTED 28870 / UNATTRIB 2393`

98.9 % of the host time is attention, i.e. an MHA running on the CPU. Not inferred from the label: the device saw
n=144 launches per chunk and per token — exactly 36 layers x 4 GEMMs. If attention were dispatched, n would be 180.

Structural reason: every b8 set on this box (Qwen2.5-0.5B oga4/kernels-i8-b8, Llama-3.2-1B oga13/kernels-llama-b8,
Qwen3-4B oga11/kernels-4b-b8) ships four roles x 2 M-tiles = 8 ELFs — QKV, O, GU, D. There is no attention kernel in
any of them, so MHA/rope/rmsnorm/silu/residual run host-side by construction, while the pack we import fuses
MLADFMHAROPE/FLATMHA/SILU into its dispatched unit. Our GEMM-only sets were sufficient for the loader and for
text-identical output; they are the whole of the speed gap.

Levers, largest measured headroom first:

an NPU attention kernel for the b8 families — M=128 prefill and M=16 decode twins — in private npu-kernels (brief step 4's
territory; this is the 98.9 %);
the layer GEMMs' per-launch stream efficiency (3.6 → 31.3 GB/s is what one launch of lm_head already achieves);
nothing about the host submission path as the primary lever — it is 11 % of a decode token at ctx 128.

Hub: FACTS-npu-decode-dispatch-latency.md, new section plus the 2026-11-05 CORRECTION 7 (d41f851, 31b8c72), evidence
under evidence/npu-dispatch-latency/2026-11-05-{longctx-stages,role-profile}/. Nothing merged, tagged or published; three
npu-hold tickets for this item, all rc 0, amdxdna ERROR 0 -> 0.

CORRECTION to my own comment above (2026-11-05, the per-role profile): the "lm_head streams at 31.3 GB/s / 9-12x per-byte gap" line was an arithmetic error, and I am withdrawing it.

lm_head issues 8 launches per token (one per eighth of the 151,936 vocab — the dump's own decode dispatches=152 = 36 layers x 4 roles + 8). Last night I divided its per-token 206.6 MB by one launch's 5.85 ms and published 31.3 GB/s, then derived from it that "the number is reachable inside our own stack" and that burst length/launch width was the mechanism. The correct division is bytes-per-launch / seconds-per-launch: 25.83 MB / 5.887 ms = 4.39 GB/s.

Recomputed from the same committed log (no new device time), medians over the run's stable phase:

role	N x K	launches/token	ms/launch	MB/launch	GB/s	MB/token
QKV	6144 x 2560	36	2.654	8.36	3.15	300.8
O	2560 x 4096	36	1.882	5.57	2.96	200.5
GU	19456 x 2560	36	6.557	26.46	4.04	952.6
D	2560 x 9728	36	3.613	13.23	3.66	476.3
lm_head	151936 x 2560	8	5.887	25.83	4.39	206.6
total		152			3.71	2136.8

The rows now reproduce the run's own per-token [ph] totals exactly — layers 529.5 ms vs 529.3 measured (0.0 %), head 47.1 vs 46.8 (0.6 %), 152/152 launches accounted for — and that check is now executable rather than a claim: recompute_rates.py in the evidence directory exits non-zero if a row's launches x ms/launch does not reproduce the total. Decode budget is therefore 577 ms array / ~50-58 ms host (8-9 %), not 559/70.

What this changes:

The width axis is measured as weak: 2.96 GB/s at N=2560, 4.04 at N=19456, 4.39 at N=151936 — 59x of width buys 1.5x. So "stream the weights through one wider launch" is not a 9-12x lever, and the "one shape at two N widths" experiment I proposed on the strength of the bad table is deprioritised (the five existing shapes already answer it).
No headroom inside our GEMM path is demonstrated. The layer roles are within 20 % of the best per-byte rate anywhere in our stack. The "existence proof" claim is withdrawn.
Unchanged, and now the whole story: the long-context prefill finding in that same comment — 98.9 % of the host time is a CPU attention because no b8 set contains an attention kernel. Different instrument ([lh]), untouched by this correction. That one still stands as the largest measured lever.
What the corrected numbers point at instead: a uniform ~3.7 GB/s aggregate against the ~15.5 GB/s implied by AMD's recorded 72-77 tok/s on the 0.5B (206 MB/token) — a 3.5-5x per-byte deficit with no measured mechanism yet. The two candidate axes are the b8 burst-within-a-launch (one 8 KiB burst per tile) and host memory pressure during the measurement (load1 21.07 in this very run; 10 of its 45 decode windows sit 1.6x higher than the other 35). Neither is established; I am recording them as open questions, not as a diagnosis.

Hub: npu-re @ e2879c2 (docs/FACTS-npu-decode-dispatch-latency.md, CORRECTION 8) with evidence/npu-dispatch-latency/2026-11-05-role-profile/{recompute_rates.py, recompute-2026-11-05.txt}.

The check that would have caught it, now a rule for every row I publish: a per-launch rate needs its launch count in the same row, and the row's product must reproduce the per-token total — a table that mixes a per-launch time with a per-token byte count under a GB/s heading is dimensionally invalid, which is what this one was.

The decode dispatch law, measured instead of described two different ways: t = 0.76 ms + MB / 4.76 GB/s, over ten launches on two models — and the two levers priced apart.

The per-role profile in my comment above existed for the 4B only; the 0.5B is the model whose shapes are 3-10x smaller at the same M=16 twins and the same b8 design, so its row makes the primitive measurable. Ticket ogaload-prof05b (held 10 s, rc 0, 8/8 b8 ELFs, amdxdna ERROR 0 -> 0, load1 21.96), 112 launches/token (24 layers x 4 + 16 head launches):

model	role	MB/launch	ms/launch	GB/s
Qwen3-4B	QKV / O / GU / D / HEAD	8.36 / 5.57 / 26.46 / 13.23 / 25.83	2.658 / 1.886 / 6.563 / 3.619 / 5.888	3.14 / 2.95 / 4.03 / 3.66 / 4.39
Qwen2.5-0.5B	QKV / O / GU / D / HEAD	0.84 / 0.56 / 5.29 / 2.65 / 5.17	0.856 / 0.976 / 1.863 / 1.414 / 1.596	0.98 / 0.57 / 2.84 / 1.87 / 3.24

Least squares over those ten points: t_ms = 0.760 + MB / 4.76, worst residual 15.6 % (the 0.5B head is the outlier at -16 %; the other nine are within 10 %). Two models, 120x range in bytes per launch, one line. The "GB/s per role" spread from 0.57 to 4.39 is not a rate difference — it is a fixed floor divided by a small byte count.

Per-token budgets, same fit:

model	launches/token	MB/token	floor ms	stream ms	total	floor share
Qwen3-4B	152	2136.8	115.6	448.5	564.0 (measured 576.6)	20 %
Qwen2.5-0.5B	112	306.6	85.2	64.3	149.5 (measured 146-171 wall)	57 %

Why this matters more than either previous column: it prices the two candidate fixes separately, on the model where AMD's own number is on record. Against AMD's recorded 72-77 tok/s on the 0.5B:

what	ms/token	tok/s
measured	149.5	6.7 (gate row 5.83-6.03)
floor removed (112 launches -> 1, same per-byte rate)	65.1	15.4 (2.3x)
stream at AMD's implied 23 GB/s (floor kept)	98.5	10.2 (1.5x)
both	14.1	71.0 ≈ AMD's 72-77

So both terms are real and neither substitutes for the other: our per-byte stream is 4.8x below AMD's implied rate, and our floor is ~112x theirs, because AMD's recorded contract is one EXEC per decode token. Fusing launches is worth 2.3x on the 0.5B at an unchanged rate; a 4.8x better stream is worth 1.5x while 112 launches remain.

This also explains a thing that has bitten this issue twice: the "constant ms/dispatch" law drifted (1.8 -> 4.5 ms between models and layouts) because the constant is a floor, and it only dominates when the bytes per launch are small. One number stops being a "design property" and becomes a floor plus a rate.

It does not touch the long-context prefill finding (98.9 % host attention) — decode and prefill have different dominant terms, which is exactly why one "ms/dispatch" number kept misleading us.

Caveat: medians include each run's contended tail (spread shown in the hub table), load1 was 21.07 and 21.96. A quiet box would move the stream term; the floor is per-launch work, not memory traffic.

Hub: npu-re @ 17a50f0 — docs/FACTS-npu-decode-dispatch-latency.md (new section) + evidence/npu-dispatch-latency/2026-11-05-05b-roles/{instrument-log-05b-ctx128.log, rows-05b-ctx128.txt, fit_floor.py, fit-2026-11-05.txt}.

Third family out of sample: the floor+stream law holds — t = 0.794 ms + MB / 4.83 GB/s over 15 launches on 3 models, and it explains the "4 ms per dispatch" row this doc/issue has been quoting for the 1B.

Ticket ogaload-profl1b (held 13 s, rc 0, 8/8 b8 ELFs, amdxdna ERROR 0 -> 0, load1 12.41). Llama-3.2-1B: 72 dispatches/token = 16 layers x 4 + 8 head launches.

role	MB/launch	ms/launch (spread)	GB/s	law says	resid
QKV	3.34	1.571 (1.557-1.653)	2.13	1.486	+5.4 %
O	2.23	1.325	1.68	1.256	+5.3 %
GU	17.83	4.703	3.79	4.485	+4.6 %
D	8.91	2.729	3.27	2.640	+3.3 %
HEAD	17.44	4.020	4.34	4.406	-9.6 %

Refit on all fifteen launches: t = 0.794 + MB / 4.83 GB/s, worst residual 16.8 % (4B all within 5 %, 1B all within 10 %, 0.5B within 17 %; ten of fifteen under 7 %). The heads are consistently faster than the single line in every container, which is the width effect showing through one slope.

Per-token budgets, same fit:

model	launches/token	MB/token	floor	stream	total (measured device)	floor share
Qwen3-4B	152	2136.8	120.7	442.4	563.2 (576.6)	21 %
Qwen2.5-0.5B	112	306.6	89.0	63.5	152.4 (146.3)	58 %
Llama-3.2-1B	72	656.5	57.2	135.9	193.1 (196.9 device / 212 wall)	30 %

The 1B "4.02 ms/dispatch" is explained: it is wall time divided by 64 layer launches — it excludes the head's 8 — whereas the device figure is 2.73 ms over 72 launches, and the model's device share is 197 of 212 ms (7 % host). So there is nothing 1B-specific about it; it was the usual accounting asymmetry, in the same family as the lm_head error I corrected earlier today.

My own fault on the way here, recorded rather than quietly dropped: I pre-registered the 1B prediction before the run and predicted 111 ms device / 7.7-8.1 tok/s, which would have made the measurement a 1.6x falsification. I had halved the per-launch byte counts (QKV written as 1.671 MB where 3072 x 2048 int4 + 6.25 % scales is 3.34 MB). With correct bytes the law predicts 193.0 ms — 2 % from the measured 196.9 — and all five roles land within 10 %. Same lesson as the earlier correction, one level down: a byte count inside a prediction is a derived number and needs its own control (that N x K x 0.5 x 1.0625 reproduces the ELF's tile count).

One comparability caveat worth flagging, not resolved: this run's 1B wall is 212 ms/token (4.71 tok/s) where the recorded gate (c) row for the same family is 257 ms (3.89 tok/s) — 21 % apart, with the same container and kernel set but a different engine binary (~/wt/oga-fwd-timings/build/1bit here, ~/wt/i671kld/build-npu/1bit there). Either the binary or the box accounts for it; until someone runs both to a clean log, the two numbers should not be mixed in one table.

Hub npu-re @ c0055eb — docs/FACTS-npu-decode-dispatch-latency.md (fit section extended, the title/law notes now point at the three-model constants) + evidence/npu-dispatch-latency/2026-11-05-1b-roles/ and fit-3models-2026-11-05.txt.

Follow-up to my "binary or box?" caveat above — it is neither. Interleaved A/B in one ticket says the two engine binaries are equivalent, so no gate row understates main, and the 21 % was a 4-token vs 32-token measurement window.

Instrument ~/ogaload/oga13/binab.sh (dryrun-validated; prints the two shas and exits without touching the device), ticket ogaload-binab (held 159 s, rc 0, amdxdna ERROR 0 -> 0, same container + kernel set, order alternated between rounds so box drift cannot land on one arm):

round	arm	binary	decode tok/s	how
1	A	`41eb6129` 11-03 (the gate row's build)	3.49	wall-delta (predates the `timings` key)
1	B	`e670036` 11-05 (the profile's build)	3.97	stages
2	B	`e670036`	3.96	stages
2	A	`41eb6129`	3.99	wall-delta

A's two readings straddle B's; B reproduces to 0.3 %. Both reproduce the recorded gate (c) row (3.89 tok/s) within 2 %. So:

The binary is not the cause — and, importantly, the recorded gate rows do not understate current main by ~20 %. That was the version of this caveat that would have invalidated every AMD comparison in the issue; it is closed rather than left hanging.
The box is not the cause either in the way I framed it: this ticket ran at load1 7.4-9.0, lower than the profile run's 12.41, and still produced ~4.0 tok/s.
What it was: the profile run's [ph] line is a 4-token generation (gen=4); the gate row and this A/B are 32-token ones. Same binary (e670036), same container, same 128-token prompt: ≈212 ms/token vs ≈252 ms/token. So short generations read 15-20 % fast here.

Practical rule this adds (and the reason I am not just deleting the number): a decode rate must name its generation length. The per-role ms/launch numbers behind the floor+stream law are unaffected — they are per-launch device waits inside a token and reconcile with [ph] to 0.1 % — but the 4.71 tok/s cell I quoted for the 1B should not be used as a decode row, and won't be again.

Not explained, and deliberately not guessed: why the first ~3 tokens are 18 % cheaper than a 31-token average (position in the generation, KV/L2 warmth, and [ph] bookkeeping are all candidates). Separating them needs a per-token wall dump across a 32-token generation — one ticket; not spent here.

Hub npu-re @ d31334f — docs/FACTS-npu-decode-dispatch-latency.md (new section) + evidence/npu-dispatch-latency/2026-11-05-binaab/{binab.sh, binab_array.txt}.

Brief step 6, third pack: the loader imports a fourth pack (Qwen3-1.7B, a new hidden size) with no loader change — and the pack is coherent end-to-end device-free; what it lacks is kernel shapes, which now have a predicted rate from the law.

Qwen3-1.7B_rai_1.8.0_npu_4K was sitting on the box unimported. It is a genuinely different shape from the three we run: hidden 2048, 28 layers, 16/8 heads, head_dim 128, qk_norm, rope_theta 1e6 (1bit npu-pack info output, incomplete: []).

Import (1bit npu-pack import, 10 s, device-free): 311 tensors, 210,016 q4_1 tiles, 1.70 GB container, and every one of the importer's own checks passes:

self-check: lm_head vs the pack's fp16 embedding  corr 0.999999  rel_err 0.000764
controls:   nibble-swapped corr 0.0017, no zero-point corr 0.8767 (both must fail)
container:  read back through q4nx::dequant_tile  rel_err 2.842e-03  max |diff| 4.883e-04
attention:  2688/2688 .fconst tiles survive decode->encode byte for byte; 112 projections
MLP:        3584/3584 .fconst groups survive decode->encode byte for byte; 84 projections
fconst:     stride 27931456, second header 4739136; 56/56 RMSNorm anchors found (28 layers x 2)
text:       8 of 8 tokenizer/chat files copied

End-to-end, host GEMM, device-free (NPU_INFER_HOST_GEMM=1, 24 tokens): "The capital of France is" → " Paris. The capital of the United States is Washington, D.C. The capital of Canada is Ottawa.", and "2 + 2 =" → " 4, 4 + 2 = 6, 6 + 2 = 8, 8 +". Coherent and correct, so the import is not just structurally complete.

What is missing is kernels, and the law prices it. Its five GEMMs (QKV 4096x2048, O 2048x2048, GU 12288x2048, D 2048x6144, head 151936x2048) have no _b8 ELFs, so the NPU lane cannot run it. Applying today's measured law (t = 0.792 ms + MB/4.83 GB/s, 15 launches / 3 models): 914.0 MB/token over 120 launches (112 + 8 head) ⇒ 95.0 ms floor + 189.2 ms stream = 284 ms/token, i.e. a 3.52 tok/s device ceiling (floor share 33 %). A new shape set for this pack is therefore a ~3.5 tok/s proposition before host cost, which is the number an owner sizing call needs — it is the same class as the attention-kernel ask, not a bigger one.

Generality measured across all four packs (all on the box, all device-free, ~40 s total) — this is brief step 6's "repeat 3-5" with the loader's own evidence as the metric, now in FACTS-oga-pack-weight-layout.md:

pack	layers	lm_head corr / rel_err	controls (must fail)	attention payload	norm anchors	container round trip
Qwen2.5-0.5B	24	0.999999 / 7.00e-4	-0.0008 / 0.8788	768/768 byte-exact	48/48	3.34e-3
Llama-3.2-1B	16	0.999999 / 7.01e-4	0.0004 / 0.8586	1280/1280	32/32	3.34e-3
Qwen3-1.7B	28	0.999999 / 7.64e-4	0.0017 / 0.8767	2688/2688	56/56	2.84e-3
Qwen3-4B	36	0.999999 / 7.37e-4	-0.0020 / 0.8645	4896/4896	72/72	2.75e-3

That also closes #734 by measurement: norm anchors are found at eps-derived strides on hidden 1024 / 2048 / 2560, not just 896.

Two things worth naming as facts rather than leaving implicit:

The fidelity floor of any imported container is bf16 scales. Our container keeps scales/zero-points as bf16 where the pack has fp32, so only 43-48 % of 32-column groups survive bit-equal and the round-trip is 2.7-3.3e-3 relative. That is the bf16 mantissa, not a parse error — but any "weight-exact" claim about the loader should carry it.
One pack is text-incomplete: Llama-3.2-1B ships 5 of 8 text artifacts — no vocab.json, merges.txt, added_tokens.json (the other three packs ship all 8). The importer reports it rather than proceeding silently, and that container serves and passes gates on tokenizer.json alone, so it is a recorded pack difference on exactly one family out of four.

Hub npu-re @ 0f9dd0a (+ evidence/oga/pack-layout-2026-11-05/import-4packs.txt, with the reproduce commands).

Retracting my own explanation from an hour ago: the "4-token vs 32-token window" was not the cause of the 21 %. Per-token decode cost is flat in position, and the spread is between-run.

I ran the experiment my previous comment said was missing — a 32-token generation with NPU_PHASE_PROF=1, so [ph] tok gives the wall per token (ticket ogaload-wallgrowth, both logs in the hub):

model	per-token wall, pos 128→158	first 4	last 4	growth	mean over 31 tokens
Llama-3.2-1B	239, 245, then 214-220 flat	229.4 ms	217.0 ms	−5.4 %	218.1 ms = 4.59 tok/s
Qwen2.5-0.5B	154-167 flat, then a step at pos 153 to 183-206	157.1 ms	197.6 ms	+25.8 %	165.6 ms = 6.04 tok/s

So:

Generation length: falsified. The 1B's 32-token mean is 218.1 ms against the 4-token run's 212.4 — 2.7 %, and the short window is if anything slower per token (first two tokens pay ~10 % setup). My claim that "short generations read 15-20 % fast" was wrong in both magnitude and sign.
The 0.5B's +25.8 % is a step, not a slope — flat for 25 tokens, then a jump at pos 153, while the box's load went from ~7 to 39 during the run. That is external load arriving mid-generation; a [ph] tok series is a good load tripwire and says nothing about position.
load1 is excluded in the opposite direction to expectation: the slower runs had the lower load (A/B at load1 7.4 → 3.96 tok/s; this run at a rising 7→39 → 4.59 tok/s).
What remains is between-run variance of 15-21 % on identical settings: same binary e670036, same container, same 128-token prompt, 32 tokens — 251-253 ms in the A/B, 218 ms here. Candidates I am not going to pick between from this data: the profiling env (NPU_PROFILE/NPU_PROFILE_TOKEN/NPU_PHASE_PROF) and --prefix-slots 0 (this run had both, the A/B had neither), and container page-cache residency (the 4B rows already showed 1.55 → 1.05 tok/s with a cold container).

What this changes about the earlier comment's practical conclusion: the part that stands is that the two binaries are equivalent and the gate rows are not invalidated (that came from the interleaved A/B, not from the window theory). The part that falls is the mechanism I attributed it to. And the number now attached to this lane's oldest rule: a single decode run on this box carries ±15-20 %, so a decode row is a median of interleaved runs or it is not a row — the gate (c) rows satisfy that, and my own 4.71 / 4.59 single-run figures do not (read them as "≈4.6 tok/s under that run's conditions").

Next experiment, for whoever takes it: apply the A/B pattern to flags rather than binaries — profiling env on/off × --prefix-slots 0/default, order alternated, 32-token generations, plus a container-residency read before each arm (the instrument already exists in gate_c_llama.sh). One ticket separates the remaining candidates.

Hub npu-re @ 9100522 — docs/FACTS-npu-decode-dispatch-latency.md CORRECTION 9 + evidence/npu-dispatch-latency/2026-11-05-wallgrowth/{per-token-wall-1b-32tok.log, per-token-wall-05b-32tok.log, walltrend.py}.

That is the third number today that survived one check and failed the next, all for the same structural reason — an implicit selection (which window, which launch count, which run). The tooling added today (recompute_rates.py's reconciliation, the spread guard, walltrend.py's coverage/control refusals) exists to make that class detectable without me remembering it.

npu-pack: an imported container keeps the pack's fp32 scale/zero and …

Status (recorded so the tracker's map is complete): this is the owner's flagship OGA-pack program, not a parked defect — the lane (pi-npufeat / pi-mosaic) is working it in rungs, and the thread's latest work is performance attribution on imported packs.

Where it stands per the thread: the route works (packs import through the engine's own importer and serve), and the latest measured state is the per-token wall analysis on Llama-3.2-1B (218.1 ms = 4.59 tok/s mean over 31 tokens, flat in position — the "4-token vs 32-token window" explanation was retracted) and Qwen2.5-0.5B (165.6 ms = 6.04 tok/s, with a step at pos 153). The program's own acceptance (vendor-stack parity gates: greedy output on three prompts, 12-chunk KLD, decode/prefill rows side by side, a second family) is the closing bar.

No action is blocked on this entry: it closes when the lane's vendor-parity rows land.

Sweep (pi-verify) at `918f731a` — plan step 1 is in the tree and tested; the gates are not

1bit npu-pack <info|import> exists and dispatches to the importer (app/main.cpp:95,267, app/pack_import.h), with
tests/npu_pack_import_test.cpp and tools/oga_pack_tiles_golden.py covering it. So step 1 is in better shape than "exists":
the import half is in the tree and gated.

What the issue still asks for and the tree does not contain: the per-shape ELF set for a pack's geometry (generator side lives
outside this repo), the three gates against the vendor stack on the first pack, and a second pack of another family. Still open,
unchanged in shape.

Ledger: what has landed, what remains, and the two bar items that cannot be re-measured

Lane pi-ogaload, 2026-11-06. Posted because this is a program tracker: its own closing bar is
measurement, so here is every item of the plan with its evidence, then the three things genuinely open.

plan item	state	evidence
1. `1bit npu-pack import`	done	6 packs on this box import with no refusals (Qwen2.5-0.5B, Llama-3.2-1B, Phi-3.5-mini, Qwen3-4B, Qwen3-1.7B; the `_hybrid` pack is correctly reported as not an OGA NPU pack); #780, #1018, #1185, #1217
2. per-shape ELFs from our generator	4 of 5 packs	0.5B, Llama-1B, Qwen3-4B and Phi-3.5 have b8 sets (`full_i8__b8[_m16].elf`), validated by a full engine forward on the array. Qwen3-1.7B has none* (import-only, host-GEMM route) — private artifact work, not engine
3. gates on the first pack	closed on the restated bar	(a) greedy: 2 of 3 prompts identical to AMD's own stack through EOS, the third is AMD's own 0.057 near-tie; (b) engine fidelity over 12 chunks KL 3.1e-5; (c) b8 rows, medians of 5: prefill 220.7 / 180.4 / 93.8 tok/s, decode 5.83 / 6.03 / 5.56 tok/s at ctx 128 / 512 / 2048
4. a second pack, different family	done, four times over	Llama-3.2-1B (llama), Phi-3.5-mini (phi3, MHA 32/32 at head_dim 96), Qwen3-4B (qwen3, qk_norm), Qwen3-1.7B (import only). Unsupported features were filed as issues, not silenced
5. ORT path	written up, not started (as scoped)	the pack's `model.onnx` is one `com.ryzenai.DynamicDispatch` node, so there is no graph for an EP to partition; write-up on this thread

Landed today, both measured: #1243 (host attention: AVX2+FMA score dot + head-parallel, so a
2048-token prefill went 208.1 s → 52.7 s, 3.94x, same protocol as the gate (c) rows) and #1290
(npu-pack import --kernels, which links the kernel set into the container so 1bit serve -m <dir> --device npu needs no hand-made symlink — that was the last manual step in this issue's route).

What remains, in order

The dispatch gap, which is now the whole decode difference. Decode is 5.83-6.03 tok/s against
AMD's recorded 72-77, and the measured law over 20 launches on three models is
t = 0.76 ms + MB / 4.76 GB/s for 112-128 launches per token where AMD issues one EXEC per
token. Faster kernels do not move this; fewer submissions do (the fused-lane artifact family:
layer_ctx{1,2,17}.elf + lmhead.elf + layer.pdi). Artifact work in the private repo.
A kernel set for Qwen3-1.7B (item 2's remainder), then its gate rows.
The Qwen3-4B gate (c) row is quotable only at ctx 128: its 3.07 GiB container's rows drift with
page-cache state (4.76 → 3.22 GB/s effective weight rate across three runs of the same artifacts),
which is a box-state fact, not a model property.

Two bar items that cannot be re-measured, and why

The bar says "12-chunk KLD of our logits vs the vendor stack" and "side by side with the vendor stack
measured the same hour". AMD's 1.8.0 Linux stack was deleted on 2026-10-23 on owner approval and its
wheels are in no backup, so neither can be taken again: the vendor-parity rows here are against the
captured vendor logits (3 positions per prompt, KL 0.0009 / 0.0006 / 0.0054) and against AMD's recorded
speed figures, which is the restatement the orchestrator issued on 2026-10-30. Everything else in the bar
is measured and linked above.

So: closing this tracker is an owner/queue call on the restated bar, not something I will do by hand.

Closing: the engine side of this program is complete, and the one remaining gap is now #1309

Re-checked against this issue's own plan, item by item, with the evidence. Nothing here is inferred from a
previous comment: each row is a measurement that exists.

item	state	where the proof is
1. `1bit npu-pack import`	done	5 OGA packs import with no refusals on this box (Qwen2.5-0.5B, Llama-3.2-1B, Phi-3.5-mini, Qwen3-4B, Qwen3-1.7B), and the `_hybrid` flavour is correctly reported as not an OGA NPU pack. Weights byte-exact through decode->encode on every one (e.g. Qwen3-1.7B: 2,688/2,688 attention tiles, 3,584/3,584 MLP groups, 56/56 norm anchors, fp32 scale/zero exact for 100.00 % of groups)
2. per-shape ELFs from our generator	done, 5 of 5 imported packs	0.5B, Llama-1B, Qwen3-4B, Phi-3.5-mini, and Qwen3-1.7B closed today (private npu-kernels #28): 8 ELFs, host-only emission, hidden 2048 / head_dim 128 / GQA 16-8
3. gates on the first pack	closed on the restated bar	(a) 2 of 3 prompts identical to the vendor's own output through EOS, the third the vendor's own 0.057 near-tie; (b) engine fidelity over 12 chunks, KL 3.1e-5; (c) b8 rows, medians of 5, ctx 128/512/2048: prefill 220.7 / 180.4 / 93.8 tok/s, decode 5.83 / 6.03 / 5.56 tok/s
4. a second pack, different family	done four times over	Llama-3.2-1B (llama), Phi-3.5-mini (phi3, 32/32 MHA at head_dim 96), Qwen3-4B (qwen3, qk_norm), Qwen3-1.7B (today). Each one's unsupported features were filed, not silenced
5. ORT EP claims the graph	written up, out of scope here by this issue's own text	the pack's `model.onnx` is `Shape -> Constant -> Equal -> If` with the whole body in one `com.ryzenai.DynamicDispatch` node, so there is no unfused graph to partition

The last item, closed today

Qwen3-1.7B imported but had no kernel set, so it could only run the host-GEMM route. Now:

1bit npu-pack import <Qwen3-1.7B_rai_1.8.0_npu_4K> -o ctr-17b-b8 --weight-layout b8 --kernels kernels-qwen17-b8
  kernels:    ctr-17b-b8/npu -> .../kernels-qwen17-b8  (8 full_i8 ELFs: 4 M=16 twins, 8 of the _b8 family)

npu-hold ogaload-671-qwen17: queued 12 s, held 17 s, rc 0, amdxdna ERROR 0, 8/8 distinct uuids loaded
  [forward] weight layout: b8 blocked 8x8 microtiles, tile-n-major

and the acceptance for a new set, which is that the device route agrees with the container's own f32
reference path on the same container:

prompt	host-GEMM (f32 reference)	NPU, these ELFs
The capital of France is	`Paris. The capital of the United`	`Paris. The capital of the United`
2 + 2 =	`4, 4 + 2`	`4, 4 + 2`

That also exercised the merged #1290 (import --kernels), so the route no longer needs a hand-made
symlink: the kernel dir the serve loads is the one the import wrote.

What is being carried forward, and what cannot be re-measured

npu: an imported OGA pack issues 112-128 launches per decode token where the vendor issues one EXEC, and that is the whole decode gap #1309 — the dispatch gap: 112-128 launches per decode token against the vendor's one EXEC per token,
with the measured law t = 0.76 ms + MB / 4.76 GB/s over 20 launches on three models. That is artifact
work (a fused-lane set), not host code, and after engine#1243 removed prefill's quadratic term it is the
entire remaining difference. Filed rather than left in this thread.
The Qwen3-4B speed row is quotable only at ctx 128: its 3.07 GiB container's rows drift with
page-cache state (4.76 -> 3.22 GB/s effective weight rate across three runs of identical artifacts), a
box-state fact recorded in the ledger comment above.
Two bar items cannot be taken again at all: "12-chunk KLD vs the vendor stack" and "speed rows
measured the same hour" need AMD's 1.8.0 Linux stack, which was deleted on owner approval 2026-10-23 with
no backup. The rows here use the captured vendor logits and the vendor's recorded figures, which is the
2026-10-30 restatement.

Closing on that basis: the program's stated goal — a pack dropped on a Linux box runs on the NPU through
1bit serve with our runtime and our kernels, five packs, four architectures — is measured and reproducible
from the commands in docs/npu.md.

Navigation Menu

Repository navigation

npu: load Ryzen AI OGA NPU model packs on Linux through the engine (pack import -> Q4NX -> our kernels) #671

Step 5 progress: the OGA 0.5B pack runs on the array through our kernels — with measured gate numbers

The move that unblocked it: a different kernel family, not a different licence answer

(a) greedy agreement vs AMD's stack on the same pack

(c) speed, and one lever already taken

What this leaves, in the order I would take it

19 remaining items

The first speed row in AMD's own units — the 2026-10-30 decision finally has an instrument

Gate (c) for the 4B, finally with numbers that can be quoted — and the row that still cannot

Where the time actually goes: per-role profile of the Qwen3-4B b8 path (2026-11-05)

Sweep (pi-verify) at `918f731a` — plan step 1 is in the tree and tested; the gates are not

Ledger: what has landed, what remains, and the two bar items that cannot be re-measured

What remains, in order

Two bar items that cannot be re-measured, and why

Closing: the engine side of this program is complete, and the one remaining gap is now #1309

The last item, closed today

What is being carried forward, and what cannot be re-measured

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

npu: load Ryzen AI OGA NPU model packs on Linux through the engine (pack import -> Q4NX -> our kernels) #671

Description

Activity

bong-water-water-bong commented on Oct 31, 2026

bong-water-water-bong commented on Nov 1, 2026

bong-water-water-bong commented on Nov 4, 2026

Step 5 progress: the OGA 0.5B pack runs on the array through our kernels — with measured gate numbers

The move that unblocked it: a different kernel family, not a different licence answer

(a) greedy agreement vs AMD's stack on the same pack

(c) speed, and one lever already taken

What this leaves, in the order I would take it

19 remaining items

bong-water-water-bong commented on Nov 5, 2026

The first speed row in AMD's own units — the 2026-10-30 decision finally has an instrument

bong-water-water-bong commented on Nov 5, 2026

Gate (c) for the 4B, finally with numbers that can be quoted — and the row that still cannot

bong-water-water-bong commented on Nov 5, 2026

Where the time actually goes: per-role profile of the Qwen3-4B b8 path (2026-11-05)

bong-water-water-bong commented on Nov 5, 2026

bong-water-water-bong commented on Nov 5, 2026

bong-water-water-bong commented on Nov 5, 2026

bong-water-water-bong commented on Nov 5, 2026

bong-water-water-bong commented on Nov 5, 2026

bong-water-water-bong commented on Nov 5, 2026

bong-water-water-bong commented on Nov 6, 2026

bong-water-water-bong commented on Nov 6, 2026

Sweep (pi-verify) at 918f731a — plan step 1 is in the tree and tested; the gates are not

bong-water-water-bong commented on Nov 6, 2026

Ledger: what has landed, what remains, and the two bar items that cannot be re-measured

What remains, in order

Two bar items that cannot be re-measured, and why

bong-water-water-bong commented on Nov 7, 2026

Closing: the engine side of this program is complete, and the one remaining gap is now #1309

The last item, closed today

What is being carried forward, and what cannot be re-measured

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

Sweep (pi-verify) at `918f731a` — plan step 1 is in the tree and tested; the gates are not