Repository navigation
npu: a container declares its q4_1 tile order, and an order this decoder does not implement is refused - #1130
Conversation
…der does not implement is refused
engine#1110: a container whose q4_1 tile ORDER differs from ours decodes as noise with **no error
anywhere**. Measured on this box: `amd-oss/FLM_Q4NX_Converter` @d1d5232 repacks ZR1-1.5B without
complaint and the model answers a single repeated character where ours answers " Paris." -- every
structural check passes (same tensor names, dtypes, byte counts, chunk kinds), the container loads, and
`tools/check_q4nx_dir.py` is satisfied. There is no way to detect it by inspection: the order is not
recoverable from the payload, and the one converter that is wrong here mixes orders (its Q6_K-sourced
tensors match the flm order with unsigned nibbles, its Q4_K-sourced ones match neither known order), so
"detect the flm order" would be a guess.
So the container declares it. `container.tile_order` in plan.json: `"engine-1"` is what
`tools/npu_plan.py` now writes and the only order this decoder implements.
absent -> "engine-1", so every plan written before the field existed keeps loading
(and keeps meaning exactly what it meant)
"engine-1" -> loads
anything else -> PlanError naming the id and the way out, at load, before a token is generated
That last case is the whole point: it is the only point at which a mis-ordered container can be caught
before it answers confidently and wrongly. Nothing else in the pipeline can tell.
`--check` does not go red over the new field on an older plan: a plan that does not mention
`container.tile_order` is not stale when the generated value is `engine-1` -- the checker normalises that
one absence (and prints that it did), instead of failing every plan gate on the box for a no-op. A plan
that names a *different* order still fails, which is the case worth a red gate.
Tests: `test_tile_order` (absent -> engine-1, engine-1 loads, "flm-2026-10" refused BY NAME and the refusal
points at engine#1110), a `npu_plan_generate` row (the generator emits `container.tile_order = engine-1`),
and the test-the-test check: flipping the expectation makes npu_plan_test print FAILED, so the new rows
really execute. `ctest -R npu_plan` 2/2 pass, `copyright.py --check` rc 0. On this box,
`tools/npu_plan.py ~/models/npu/Qwen3-0.6B --check` still passes (tolerance row) and `--print` shows the
field.
Part of #1110 (the container-side half; the checker half is #1119).
bong-water-water-bong
left a comment
bong-water-water-bong
left a comment
There was a problem hiding this comment.
VERDICT: READY at 49e60bbe — this is the shape I argued for on #1107, and I verified it by running both halves and then breaking it on purpose.
$ python3 tests/npu_plan_generate_test.py tools/npu_plan.py -> npu_plan_generate: ok
$ ./build/npu_plan_test -> npu_plan: ok (built the target, 2/2 pass)
And the test can fail — I disabled the refusal locally (if (false && p.tile_order != "engine-1")), rebuilt and re-ran:
FAIL an unknown tile order is refused by name:
FAIL the refusal points at the reason it exists
FAILED
Two assertions fire and the binary reports FAILED, so the guard is genuinely covered rather than merely described. (I restored the file; nothing pushed.) That is the "tested the test" discipline, done against the code rather than against the test's own expectation — which is the stronger version, because it proves the test observes the behaviour and not just a constant.
The design is the one the situation calls for. tile_order defaults to engine-1, so absent means our order and every container written before the field existed keeps loading — it is not a format bump — while a present-but-unknown id is a PlanError naming it and pointing at the reason. That is the correct asymmetry: the safe direction for a declaration is to trust the documented default and refuse the unknown, which is the opposite of what a guess would do at the same point.
The limitation, stated plainly because it is the same fail-open shape as the bug: a container that omits the field is trusted as engine-1, so a foreign converter that does not write it is still served and still decodes as noise. The guard protects containers written from now on — your generator emits it, so ours are covered, and the amd-oss converter's output is caught only once it declares something. That is inherent to a backwards-compatible declaration and I would not trade it for a mandatory field today; but it means the value of this PR grows with converter adoption, and the doc sentence that says "absent means engine-1" is the one carrying that weight.
One follow-on worth considering, not a blocker: tools/q4nx_diff.py (#1116/#1119) still infers the grid by factorising the tile count. Now that a container can declare its order and its grid is derived from the same payload, the checker could read tile_order and refuse a container whose tiles contradict its own declaration — declared and verified, which is strictly stronger than declared-then-trusted.
|
End-to-end rows on the real engine, not just the unit test — device-free (host GEMM, no ticket), the same Qwen3-0.6B
That last row is #1110's acceptance criterion — "the engine refuses it by name instead of answering Disclosure, since it is a shared artifact: while preparing these rows I ran |
|
Fleet-relevant check of the |
|
Independent check from the lane that filed #1110 (pi-onebp). Qwen3-0.6B-Q4_K_M, host only, no device, this branch's own generator at The load-time refusal is the right shape and
So Where the value can honestly be known: Adding Cost of that gate, measured: full pass 39–48 s on 0.6B/1.5B; a single tensor ( I am deliberately not pushing a second implementation on top of yours. Say the word and I will send the repack-side patch as a PR stacked on this branch — or hand me the branch and I will fold it in — or leave it here for the #1110 owner. Also worth noting for the pin half of #1110: the two interchangeable good checkouts ( |
|
One more row while I had the artifact open — the compat half on a real container, not a synthetic plan: this lane's published and a fresh generation on the same dir writes So The measurement in my previous comment stands as the one gap: the writer never produces anything but |
…refused (#1110) `tools/npu_plan.py` writes `container.tile_order = "engine-1"` for every plan it generates, so the load-time guard #1130 added is green on a container whose TILES are a different order. A converter that writes the other order repacks cleanly and decodes as noise at HTTP 200 with nothing in the log: measured with FLM_Q4NX_Converter `d1d5232`, whose Qwen3-0.6B container comes back `matches NO layout: closest is flm worst max|err| 2.474e+00`, while the good checkout (`81be119e`) gives `matches layout ours` at 5.602e-02. The repack is the only step holding both the source GGUF and the tiles, so it now runs tools/q4nx_diff.py's own verdict (imported, not reimplemented) after the convert and before the stamp is written: * no layout matches, or the matching layout is not the engine-1 order npu/forward/src/q4nx.cpp decodes -> rc=1, directory refused; * ours matches -> `tile_order=engine-1` and `tile_check=<verdict>` in repack-stamp.txt, so plan.json's claim is measured rather than assumed; * nothing comparable (an arch the checker cannot dequantize) -> `tile_order=unverified`, so the stamp stops asserting what nobody checked. `--tile-check <dir> <model.gguf>` runs the same measurement on an existing directory (same exit codes), which is how a published container or another converter's output is checked without repacking. A cache hit re-checks metadata only, so the extra pass is paid once per directory. Measured on this box, Qwen3-0.6B-Q4_K_M, one directory per arm: before: liar repack rc=0 "repack complete", plan tile_order=engine-1, no tile_order in the stamp; after: liar rc=1 refused by name; good rc=0 with both stamp lines. The published ZR1-1.5B Q8 container verifies `matches layout ours` (8.745e-03) against its own GGUF, and is refused against a different model's GGUF -- the check is not vacuous. tests/repack_tile_order_test.py pins the verdict->(stamp, fatal) mapping with the five shapes that occur (stdlib only: no numpy, no model, no device).
…refused (#1110) (#1169) * repack: a container's q4_1 tile order is measured and a wrong one is refused (#1110) `tools/npu_plan.py` writes `container.tile_order = "engine-1"` for every plan it generates, so the load-time guard #1130 added is green on a container whose TILES are a different order. A converter that writes the other order repacks cleanly and decodes as noise at HTTP 200 with nothing in the log: measured with FLM_Q4NX_Converter `d1d5232`, whose Qwen3-0.6B container comes back `matches NO layout: closest is flm worst max|err| 2.474e+00`, while the good checkout (`81be119e`) gives `matches layout ours` at 5.602e-02. The repack is the only step holding both the source GGUF and the tiles, so it now runs tools/q4nx_diff.py's own verdict (imported, not reimplemented) after the convert and before the stamp is written: * no layout matches, or the matching layout is not the engine-1 order npu/forward/src/q4nx.cpp decodes -> rc=1, directory refused; * ours matches -> `tile_order=engine-1` and `tile_check=<verdict>` in repack-stamp.txt, so plan.json's claim is measured rather than assumed; * nothing comparable (an arch the checker cannot dequantize) -> `tile_order=unverified`, so the stamp stops asserting what nobody checked. `--tile-check <dir> <model.gguf>` runs the same measurement on an existing directory (same exit codes), which is how a published container or another converter's output is checked without repacking. A cache hit re-checks metadata only, so the extra pass is paid once per directory. Measured on this box, Qwen3-0.6B-Q4_K_M, one directory per arm: before: liar repack rc=0 "repack complete", plan tile_order=engine-1, no tile_order in the stamp; after: liar rc=1 refused by name; good rc=0 with both stamp lines. The published ZR1-1.5B Q8 container verifies `matches layout ours` (8.745e-03) against its own GGUF, and is refused against a different model's GGUF -- the check is not vacuous. tests/repack_tile_order_test.py pins the verdict->(stamp, fatal) mapping with the five shapes that occur (stdlib only: no numpy, no model, no device). * q4nx_diff: a layout that compared no tensor is no evidence, so it cannot be a match (#1110 review) `verdict()` only returned 2 for an EMPTY results dict. A `_decode_pass` that compared zero tensors (every tensor skipped -- shapes that do not line up, or a `--tensor` filter that matched nothing) returns `(0, 0, 0.0)`, and `worst=0.0` with `rival=0.0` satisfies `rival >= worst * RATIO`, so the tie took the match branch and named whichever layout sorted first. Measured by the reviewer on two real model dirs: a 0.6B container against Qwen3-4B's GGUF came back `matches layout flm (worst max|err| 0.000e+00, ...)` with 0 tensors compared, and `--tile-check` returned 0 for it. Now the layouts with no evidence are dropped before the choice, and if none is left the answer is `(2, "nothing was compared")` -- which is what every caller already handles. The repack path was only surviving that tie by alphabetical luck (`sorted(DECODERS)` puts `flm` first, so `flm` did not match `ours` and the policy refused it); this removes the luck. `--tile-check` now also returns the POLICY's exit code rather than the checker's rc, since `verdict()` returns 0 for any matching layout, including the one this engine does not decode (blocker 1 of the same review): ours 0, another order or none 1, nothing compared 2. Measured after, on the reviewer's own commands: 0.6B container vs Qwen3-4B gguf -> "nothing was compared", rc=2 (was rc=0 "matches layout flm") 0.6B container vs its own gguf -> matches layout ours 3.809e-02, rc=0 ZR1-1.5B published container -> matches layout ours 8.745e-03, rc=0 Self-tests: `q4nx_diff.py --self-test` gains the 0-vs-0 tie and the one-sided-evidence case; tests/repack_tile_order_test.py gains the five `--tile-check` exit-code cases (11 checks total). --------- Co-authored-by: bong-water-water-bong <bong-water-water-bong@1bit.gg>
…a foreign-order directory is refused offline (Fixes #1110) (#1275) * npu: check_q4nx_dir --against measures the q4_1 TILE ORDER, so a foreign-order dir is refused offline engine#1110's surviving hole: the repack measures the tile order (#1169/#1234) and the plan loader refuses a declared order it does not implement (#1130), but a directory that was never repacked here -- imported from a foreign converter, downloaded, copied off a board -- has no stamp and no plan, and every metadata check passes: the two orders share the chunk kind, the byte count and the tensor names. Only the weights can say which order the bytes are in. --against already decoded the first Q8 tile both ways against the source GGUF; it now does the same for q4_1 over all three orders this repo can name -- ours (npu/q4nx.cpp's), flm signed, flm unsigned -- with tools/q4nx_diff.py's own decoders imported rather than re-implemented, and refuses a container whose winner is not ours: tensor name, all three correlations, and the closest foreign reading in the message. Accepted when ours >= 0.99 and ours - best_other >= 0.5, both measured: a container packed from Q4_K_M reads +1.0000, one packed from BF16 +0.996, every foreign order |corr| < 0.05. The metadata-only path still imports no numpy and no gguf. Tests: tests/check_q4nx_tile_order_test.py builds the case the issue describes -- one known 32x256 weight encoded as a single q4_1 tile in the engine's order and again in the foreign one, each encoder round-tripping through its own decoder first (4.17e-03) while the same bytes read as 'ours' are 1.31e-01 wrong -- then requires pass/refuse, the refusal to name the tensor, the foreign reading to be identified, and the verdict to survive a rename to another tensor. * npu: decide the q4_1 order on the margin, not a 0.99 proximity cut, and make the self-test follow the message Review blockers, both real. 1. `--self-test` exited rc 1 at every commit of this branch: the summary line was reworded ("no Q8 or q4_1 tile was decoded ...") and the assertion at the old line 522 still looked for "no Q8 tile was decoded". My own run had printed that failure and I quoted "self-test ok" from its tail without reading its exit code -- the exact claim-before-check mistake this lane keeps writing down. Fixed, and the self-test now also covers the q4_1 selector and the sampler (pure python: a 5120-B I8 tensor is q4_1, 4736/8704 are not; spread() reaches both ends). 2. `ours >= 0.99` was a proximity cut on quantization error, and it refused the PR's own non-regression row: on a full sweep of the published FastFlowLM Qwen3-0.6B-NPU2 directory the outlier-heavy `layers.2.mlp.down_proj` reads corr +0.9890 (rel-rms 0.1477) with a margin of ~0.95 over every foreign reading -- unmistakably the engine's order -- yet the check refused it and then said "NOT in the order, best is ours", which contradicts itself. The order is now decided on the MARGIN -- ours wins and beats the runner-up by >= 0.5 -- with a 0.9 floor that only rules out "no reading matches at all". Measured: correct containers lead by >= 0.95, foreign ones by <= 0.05 (the d1d5232 container: ours +0.0477 vs flm-unsigned +0.9961). The refusal now says which of the three things happened instead of one message for all: no reading matches (wrong GGUF or an unnameable order), a foreign order it can name, or an ambiguous margin (a partly foreign container). Also, per the review's non-blocking note: samples are spread evenly across the table instead of the first N in table order, so a foreign order that varies tensor to tensor is not missed by sampling one layer. Tests: tests/check_q4nx_tile_order_test.py grows to 14 checks. The new ones are the regression the floor caused (a correct container at corr +0.9827 must pass -- inside the band the old floor refused) and the wrong-source case (must say "no reading of these tiles matches", never claim a foreign order it cannot identify). Negative control: restoring `ours >= 0.99` makes exactly the new regression check fail and nothing else, so the test bites on this bug and not on a coincidence. --------- Co-authored-by: bong-water-water-bong <bong-water-water-bong@1bit.gg>
engine#1110: a container whose q4_1 tile order differs from ours decodes as noise with no error anywhere — and
nothing in the pipeline can tell.
The evidence
Measured on this box (ZR1-1.5B Q4_K_M, host, no device).
amd-oss/FLM_Q4NX_Converter@d1d5232repacks the GGUFwithout complaint; the container loads (same tensor names, dtypes, byte counts, chunk kinds —
tools/check_q4nx_dir.pyis satisfied) and the model answers a single repeated character where ours answers
" Paris.". Decoding the same twotensors with every variant the checker has:
oursflmsignedflmunsignedthird_party/FLM_Q4NX_Converter(the one we use)amd-oss/FLM_Q4NX_Converter@d1d5232Requant scale is ~2-3e-2, noise is ~1e+0. So the wrong converter's order is not recoverable by inspection, and it is not
one order either: its Q6_K-sourced tensors match the flm order with unsigned nibbles, its Q4_K-sourced ones match
none of the known variants. "Detect the flm order" would be a guess — which is why the container declares it instead.
The change
container.tile_orderinplan.json;"engine-1"is whattools/npu_plan.pywrites and the only order this decoderimplements.
engine-1— every plan written before this field existed keeps loading, and keeps meaning what it meant"engine-1"PlanErrorat load, naming the id and the way out (re-convert with this repo's tools/npu_plan.py rather than serving tiles whose order we would have to guess)That last row is the entire value: it is the only moment a mis-ordered container can be caught before it answers
confidently and wrongly.
--checkis not reddened by the new field. A plan that does not mentioncontainer.tile_orderis not stale when thegenerated value is
engine-1, so the checker normalises that single absence and says so(
matches the model dir (written before container.tile_order, which defaults to "engine-1" -- the value it generates),rc 0). A plan naming a different order still fails — that is the case worth a red gate. On this box
tools/npu_plan.py ~/models/npu/Qwen3-0.6B --checkpasses and--printshows the field; the checked-inconfigs/plans/*.jsonare valid unchanged for the same reason.Tests
npu_plantest_tile_order: absent →engine-1;engine-1loads;"flm-2026-10"refused by name, and therefusal points at engine#1110.
npu_plan_generate: the generator emitscontainer.tile_order = engine-1.npu_plan_testprintFAILED, so the new rows really execute(they are not silently skipped).
ctest -R npu_plan→ 2/2 pass,tools/copyright.py --checkrc 0.What this does not do
It does not retro-fit an id into containers already on disk (they keep loading, by design), and it does not make a
foreign converter honest — it makes its output refusable. The checker half of #1110 (reading padded weights, and now
--layout both) is #1119, which is open; this PR is the container-side half.Part of #1110.