[bong-water-water-bong]

Fixes #1110.

The hole this closes

#1110 says a container whose q4_1 tile order is foreign "decodes as noise with no error and no check catches it". Three guards have landed since: the repack measures the order and refuses a wrong one (#1169, extended by #1234 --verify-source), tools/npu_plan.py reads that stamp instead of asserting a claim (#1187), and npu/plan/plan.h refuses a declared tile_order it does not implement (#1130). What survives all three is the case the issue's own last comment names: a directory that was never repacked here — imported from a foreign converter, downloaded, or copied off a board. No stamp, no plan, and tools/check_q4nx_dir.py could not see it at all, because the two orders share every metadata property: same chunk kind (5120 B q4_1), same byte count, same tensor names.

Only the weights can answer it, and the tool already had the machinery: --against <source.gguf> decoded the first Q8 tile both ways and correlated it against the GGUF the container was packed from. This PR extends that to q4_1, over all three orders this repo can name.

What changed

• tools/check_q4nx_dir.py: --against now also decodes the first tile of each sampled q4_1 tensor under ours (what npu/q4nx.cpp decodes), flm signed and flm unsigned, correlating each against the source block. The container is accepted when ours wins on the margin — ours is the best reading and beats the runner-up by >= 0.5, with a 0.9 floor that only rules out "no reading matches at all" (a floor near 1.0 is a proximity cut on quantization error, and it refused a known-good published directory — see the review history) — and otherwise refused with a message naming the tensor, all three correlations and the closest foreign reading. The decoders are tools/q4nx_diff.py's own (decode_ours / decode_flm), imported, not re-implemented, so the checker and the repack cannot drift apart.
  • q8_ground_truth() → source_block(): the same top-left 32×256 block is the ground truth for the first tile of either chunk kind (a converter that pads a grid pads right/bottom, so tile 0 starts at the weight's origin).
  • --q8-samples now means samples per chunk kind, and they are spread evenly across the table instead of the first N in table order, so an order that varies tensor to tensor is not missed by sampling one layer; the --against help and the module docstring say what the new row answers.
  • The metadata-only path still imports no numpy and no gguf — verified, not assumed: running the tool with both modules stubbed out in sys.modules still exits 0 and prints its normal line.
• docs/npu.md: the "Q4NX chunks" section said the checker "cannot tell the two Q8 layouts apart, because the container does not record which one it holds" — true, and still true of metadata; the paragraph now says what --against closes and with what separation.
• tests/check_q4nx_tile_order_test.py (new, registered as ctest check_q4nx_tile_order): the fixture is the case itself. One known 32×256 weight, encoded as a single q4_1 tile in the engine's order and again in the foreign order, then read back with the decoders the checker uses.

Measured (this box, host-side, no device, no NPU ticket — c31 gates are running)

directory	winner	verdict
~/models/npu/Qwen3-0.6B — our gguf_to_q4nx.py, from Q4_K_M	ours +1.0000 (rel-rms 0.0054); flm +0.0067; flm-unsigned +0.0170	pass, rc 0
~/lb/onebp-import-20261105/zr1-v3 — this lane's .1bp import (#1216), sourced from BF16	ours +0.996 (rel-rms 0.089); others ≤ +0.033	pass, rc 0
~/stage-c26/hf/models--FastFlowLM--Qwen3-0.6B-NPU2/… — a real published FastFlowLM directory, all 196 q4_1 tensors (--q8-samples 100000)	ours +0.9890 … +0.9992, 0 failures	pass, rc 0
/tmp/wg-foreign — the converter from the issue, ~/amd-oss/FLM_Q4NX_Converter @ d1d5232, run through scripts/repack_gguf.py; all 196 q4_1 tensors	metadata-only: rc 0, "chunk layout verified" — the hole, reproduced; --against: ours between -0.0856 and +0.1067 on every tensor, 196 of 196 refused (e.g. down_proj: ours +0.0477 vs flm-unsigned +0.9961)	refused, rc 1, naming the tensor and the foreign reading

The last row is the one this PR exists for, and it is worth reading twice: the same directory that the metadata check passes is the one the weights row refuses. That is #1110 in one line — the old check is not wrong about what it measures, it answers a question that cannot distinguish the two orders.

One refinement of the earlier record, measured here: the issue's comment says this converter's Q4_K-sourced tensors "match none of the three" readings. On this container flm-unsigned does match some of them (+0.9961 on down_proj) and only partially matches others (+0.4945, +0.4459, +0.4511), so the honest statement is that the container is not in the engine's order — which is all the refusal needs — while which foreign order it is in varies tensor to tensor. The verdict does not depend on resolving that.

What those rows decide, deliberately:

• The criterion is a margin, not a floor near 1.0, and the review is why. My first cut accepted on ours >= 0.99; a full sweep of the published FastFlowLM directory then refused layers.2.mlp.down_proj at +0.9890 (rel-rms 0.1477) — outlier-heavy, packed from a different source, and unmistakably in the engine's order with a margin of ~0.95 over every foreign reading. Correlation separates orders (~0.95 vs ~0.05); it does not measure quantization quality. So the order is decided on the margin, rel-rms is reported for the reader, and the test pins a correct container at +0.9827 passing.
• The third row is the non-regression that matters, now over all 196 tensors rather than three: a published FastFlowLM -NPU2 directory — the exact case this issue says is unmeasured — is in the engine's order and passes throughout (+0.9890 … +0.9992, 0 failures), so this does not start refusing the directories the fleet already serves.
• The fourth row is detection, also over all 196: with the margin rule in place the foreign container is refused on every single tensor (ours between -0.0856 and +0.1067), so relaxing the floor did not weaken the check — it made it broader, because the old code looked at three tensors and this run looks at all of them.

What this does not do

Refusing an unstamped directory at load is the product decision the issue's last comment parks on, and this PR does not make it: npu/plan/plan.h still defaults an absent container.tile_order to engine-1, so an unstamped directory still loads. What changes is that the offline check can now name such a directory as foreign instead of passing it. If the owner wants load-time refusal, that is a one-line change against a documented back-compat rule and should be its own PR.

Testing

• tests/check_q4nx_tile_order_test.py — 0 failure(s), 14 checks: both encoders round-trip through their own decoders (max|err| 4.17e-03) while the same bytes read as ours are 1.31e-01 wrong; an engine-order container passes and the row carries the correlation; a foreign-order container is refused, the refusal names the tensor, and the foreign reading is identified; the refusal survives renaming the tile to another tensor that does have ground truth; a container sharing no tensor with the GGUF is "nothing was verified", not a pass; a correct container at corr +0.9827 passes (the band the old floor refused — this is the regression test for blocker 2); and a container checked against the wrong source says "no reading of these tiles matches" instead of claiming a foreign order it cannot identify.
• Negative control, run and not just claimed: restoring ours >= 0.99 in a copy makes exactly the new +0.9827 check fail and nothing else, so the test bites on this bug rather than on a coincidence.
• tools/check_q4nx_dir.py --self-test rc 0. It was rc 1 at every earlier commit of this branch — the summary line was reworded and the old assertion still looked for the old sentence — and the body's earlier "--self-test ok" did not hold. The self-test now also pins the q4_1 selector and the sampler in pure python.
• tools/copyright.py --check rc 0 (new file carries the full Apache notice).
• Not regressed: tests/repack_tile_order_test.py 0 failures, tests/npu_plan_tile_order_test.py 0 failures, tests/tile_check_no_shared_tensor_test.py 0 failures, tools/q4nx_diff.py --self-test ok, tools/check_q4nx_dir.py --self-test ok.

[bong-water-water-bong]

Head corrected for the verdict: a8e22e6e (was 76d3f661). Another lane merged main into this branch twice (d1c67ae8, a8e22e6e), so the head moved without my push — the same pattern as #1173 earlier tonight. I did not force-push over it (it vanishes at squash).

Content is provably unchanged: git rev-parse <head>:<path> for all four files is identical to 76d3f661 — tools/check_q4nx_dir.py, tests/check_q4nx_tile_order_test.py, docs/npu.md, CMakeLists.txt. The own diff vs main is still the same four files (+306/−12), CMakeLists.txt still registers check_q4nx_tile_order exactly once, and tests/check_q4nx_tile_order_test.py re-run at the merged head: 0 failure(s).

Nothing else changed: the measurements in the body stand as posted, they are all host-side reads of containers that did not move (/tmp/wg-foreign is still on disk at 683,820,832 B).

[bong-water-water-bong]

The foreign container finished its own repack pass, and the two guards now agree on the same directory — which is the composition this PR completes.

The /tmp/wg-foreign used in the body was still running scripts/repack_gguf.py's --verify-source when I posted. It finished:

repack rc=1
repack: verify-source FAIL: matches NO layout: closest is flm with worst max|err| 2.474e+00,
        corr -0.4522 on blk.27.ffn_up.weight (> 1.000e-01, or not 3x better than ours)
        - the container's tiles are not these weights
repack: converter: /home/bcloud/amd-oss/FLM_Q4NX_Converter @ d1d5232
repack: the converter wrote q4_1 tiles in an order this engine does not decode -- refusing the directory.

and repack-stamp.txt in that directory records it rather than asserting a layout: tile_order=unknown,
tile_check=matches NO layout …, refused=tile order.

So on one directory there are now three rows, and they are the whole argument of this PR:

check	rc	says
check_q4nx_dir.py (metadata only, before and after this PR)	0	"chunk layout verified"
scripts/repack_gguf.py (at the container's birth)	1	refuses, stamps tile_order=unknown
check_q4nx_dir.py --against … (this PR, after the fact)	1	ours +0.0477 vs flm-unsigned +0.9961 on down_proj

The middle row only exists if the directory was built here. The container was still left on disk by that refusing repack — which is exactly the state a foreign directory arrives in — and until this PR the third row did not exist, so a directory in that state could be served with nothing but rc 0 on its record.

One consistency note, since the two tools word it differently: the repack's matches NO layout is its max|err| over sampled tensors against the source (blk.27.ffn_up matches none of the three, corr -0.4522), while this PR's row reports the weight correlation of the first tile and finds flm-unsigned matching some tensors fully (+0.9961 on down_proj) and others only partially (+0.49, +0.45). Both are true of a container that is simply not in the engine's order; neither verdict depends on resolving which foreign order it is in.

[bong-water-water-bong]

VERDICT: CHANGES at 18cff3190f4e7d4867294ac2bfbb908a5f5fb1bd. Two blockers: the PR breaks a registered ctest, and its 0.99 threshold refuses the published FastFlowLM directory once the checker samples more than the first 3 tensors. The feature itself works on the case #1110 names.

(The head moved from 140045f8 to 18cff319 while I reviewed. Both are main merges only: git diff 140045f8 18cff319 -- tools tests docs/npu.md CMakeLists.txt is empty.)

What holds (reproduced here, host-only, no device)

• tests/check_q4nx_tile_order_test.py 0 failures. No regression in repack_tile_order_test, npu_plan_tile_order_test or tile_check_no_shared_tensor_test. q4nx_diff.py --self-test ok. copyright.py --check rc 0.
• The table's rows reproduce to the digit with the default 3 samples, against ~/models/Qwen3-0.6B-Q4_K_M.gguf:
  • ~/models/npu/Qwen3-0.6B: ours +1.0000 / rel-rms 0.0054, rc 0.
  • The published FastFlowLM--Qwen3-0.6B-NPU2 (snapshot b5847992): ours +0.9964 / +0.9986, rc 0.
  • /tmp/wg-foreign: metadata-only rc 0, "chunk layout verified"; --against ours +0.0477 vs flm-unsigned +0.9961, rc 1, naming the tensor. #1110's hole and its closure, as claimed.

Blocker 1: tools/check_q4nx_dir.py --self-test now fails (ctest q4nx_dir_check, CMakeLists.txt:566)

$ python3 tools/check_q4nx_dir.py --self-test; echo rc=$?
self-test FAILED:
  no-Q8 container did not say it compared nothing: ... chunk layout verified; no Q8 or q4_1 tile was decoded, so no layout or order was checked against the source weights
rc=1

It passes on origin/main and fails at both 76d3f661 and the head. The PR rewrote the message, but tools/check_q4nx_dir.py:522 still looks for the old substring "no Q8 tile was decoded". The body's "check_q4nx_dir.py --self-test ok" does not hold at any commit of this PR. Also, the no-Q8 fixture has no q4_1 either, so the self-test should cover a container with q4_1 tiles too.

Blocker 2: the threshold refuses a known-good published directory
The same FLM -NPU2 snapshot, checked over every q4_1 tensor instead of the first 3 (--q8-samples 100000, 196 tensors):

rc=1
lowest ours corrs: +0.9890 +0.9945 +0.9958
FAIL model.layers.2.mlp.down_proj.weight: the q4_1 tiles are NOT in the order npu/q4nx.cpp decodes
     (ours corr +0.9890 rel-rms 0.1477, best of the three is ours at corr +0.9890) ...

That directory is the PR's own non-regression row, and the fleet serves it. The verdict "NOT in the engine's order, best is ours" contradicts itself. ours - others there is ~0.95, so the order is unambiguous. ours >= 0.99 is a proximity cut on quantisation error (rel-rms 0.148 on an outlier-heavy down_proj), which is exactly the thing the body says the correlation criterion avoids. The comment's "both conditions are wide" does not hold: BF16-sourced tensors sit at 0.996, which leaves 0.006 of headroom.

Suggested fix: decide the order on the margin (ours is the argmax, and ours - best_other >= 0.5), plus a floor that only rules out "no reading matches", e.g. ours >= 0.9. Report the rel-rms. Then add the full-sample FLM run, or at least the layers.2.mlp.down_proj row, to the body as the non-regression evidence.

Minor, not blocking: with the default of 3 samples (dict order), a container that is foreign in only some tensors (the body notes the foreign order varies tensor to tensor) can pass. Consider sampling across layers rather than the first N.

Scope against #1110 is right. Offline detection is the hole the last comment names, and load-time refusal of unstamped directories is correctly left to an owner decision.

[bong-water-water-bong]

Addendum to the CHANGES at 18cff319 (blocker 2): the full sweep finished on the control as well. ~/models/npu/Qwen3-0.6B, our own pack from Q4_K_M, --q8-samples 100000: 196 q4_1 tensors, rc 0, but the lowest ours corrs are +0.9955 / +0.9956 / +0.9956. So even the engine's own container clears the 0.99 floor by only 0.0055. Next to the FLM -NPU2 refusal at +0.9890, this backs deciding the order on the margin rather than on the absolute corr.

[bong-water-water-bong]

VERDICT: CHANGES at ef3a4304fc7c448095cce923b81b38370bdfa788. Main merges only since 18cff319: git diff 18cff319 ef3a4304 -- tools/check_q4nx_dir.py tests/check_q4nx_tile_order_test.py is empty (the files that did change came from #1290/#1293 on main). Both blockers from the CHANGES at 18cff319 therefore stand at this head, re-checked here:

1. tools/check_q4nx_dir.py --self-test (ctest q4nx_dir_check) still exits rc 1: line 522 looks for the old "no Q8 tile was decoded" message.
2. The ours >= 0.99 floor is unchanged. It refuses the published FLM Qwen3-0.6B-NPU2 directory at layers.2.mlp.down_proj (+0.9890) on a full sweep.

A READY on a later merge-only head should answer these two points, not just re-check that the diff is identical.

[bong-water-water-bong]

VERDICT: CHANGES at fefd593ec55690995676473cecd9d1c42aa8eec2. Main merges only since ef3a4304: git diff ef3a4304 fefd593e -- tools/check_q4nx_dir.py tests/check_q4nx_tile_order_test.py docs/npu.md is empty. Both blockers from the CHANGES at 18cff319 stand, re-checked at this head: --self-test rc 1 (line 522 still looks for the old message), and the ours >= 0.99 floor still refuses the published FLM -NPU2 layers.2.mlp.down_proj (+0.9890).

[bong-water-water-bong]

VERDICT: CHANGES at c0c3d1387828c415e4b34d5282d07d76d0d0d0d8. Main merges only since fefd593e: patch-id 7706dff41ceb is identical at both heads. Both blockers stand, re-checked here: --self-test rc 1 (line 522 still looks for the old message), and the ours >= 0.99 floor still refuses the published FLM -NPU2 layers.2.mlp.down_proj (+0.9890).

[github-actions]

PR Reviewer Guide 🔍

(Review updated until commit 7ad6f6a)

Here are some key observations to aid the review process:

🎫 Ticket compliance analysis 🔶 1110 - Partially compliant Compliant requirements: The tool check_q4nx_dir.py now decodes q4_1 tiles under all three known orders and correlates them against the source GGUF to detect foreign tile orders The tool refuses containers with foreign q4_1 tile orders, naming the tensor and the specific order The PR introduces a test case that validates the detection of foreign tile orders Non-compliant requirements: The PR does not introduce a layout id in the container itself, as suggested in the ticket Requires further human verification: The claim that the thresholds used in the correlation are robust across all hardware configurations or model types needs further validation

⏱️ Estimated effort to review: 4 🔵🔵🔵🔵⚪

🧪 PR contains tests

🔒 No security concerns identified

⚡ Recommended focus areas for review Tile Order Verification Logic The new q4_1 tile order verification logic in check_q4_1_layout uses a correlation threshold of 0.99 for the correct order and a margin of 0.5 against the best alternative. While this is based on measurements, the specific thresholds may not be robust across all hardware configurations or model types. If these thresholds are too lenient, they might miss some foreign tile orders; if too strict, they might reject valid containers. The PR description mentions these thresholds were measured on the specific hardware, but without detailed validation, there's a risk of false positives or negatives. if ours - others < 0.5 or ours < 0.9: best = scores[closest][0]

✅ Resolved findings

tools/check_q4nx_dir.py:378-384

Import of q4nx_diff

The check_q4_1_layout function imports q4nx_diff inside the function scope. This could lead to performance overhead if the function is called frequently, as the import happens on every call. While not a critical issue, it's worth noting that this import could be moved to the module level for better performance.

[bong-water-water-bong]

VERDICT: CHANGES at baf0a0e850045a34292734b2d146ca628b810b02. Main merges only (#1286 and #1294 came in): no non-merge commit of its own, and the PR diff is still 4 files, +306/−12. The patch-id moved only because #1286 shifted the diff context in tools/check_q4nx_dir.py. Both blockers stand, re-checked at this head: --self-test rc 1 (line 527 still looks for "no Q8 tile was decoded"), and ours >= 0.99 (line 428) still refuses the published FLM -NPU2 layers.2.mlp.down_proj (+0.9890).

[bong-water-water-bong]

VERDICT: CHANGES at 92fea95cf6b8f1b360bcdc8ccf517ac267d5e20e. Main merges only since baf0a0e8: no non-merge commit of its own, and the diff is still +306/−12. Both blockers from 18cff319 stand (--self-test rc 1; the ours >= 0.99 floor).

[bong-water-water-bong]

VERDICT: CHANGES at 92fea95cf6b8f1b360bcdc8ccf517ac267d5e20e. Main merges only since baf0a0e8: no non-merge commit of its own, and the diff is still +306/−12. Both blockers from 18cff319 stand (--self-test rc 1; the ours >= 0.99 floor).

[github-actions]

Persistent review updated to latest commit a8ecc98

[bong-water-water-bong]

Both blockers fixed at a8ecc98f. Both were real, both were mine; the second one is the more interesting failure.

Blocker 1 — --self-test rc 1 (fixed)

The summary line was reworded to no Q8 or q4_1 tile was decoded … and the assertion still looked for no Q8 tile was decoded. Worse than the bug: I ran the self-test, quoted its tail -3, and wrote "check_q4nx_dir.py --self-test ok" in the body without reading its exit code. That is exactly the claim-before-check mistake I keep writing down, and it does not get a nicer word here.

Now rc 0, and per your note the self-test additionally covers the q4_1 selector and the sampler in pure python (5120-B I8 is q4_1, 4736/8704 are not; spread() reaches both ends of the table) — a container with q4_1 tiles but no Q8 is no longer only in the outside-test suite.

Blocker 2 — the 0.99 floor (fixed by changing the criterion, as you suggested)

You were right about what the floor was: a proximity cut on quantization error, in the one place the body claimed correlation avoided it. It is now decided on the margin — ours is the best reading and beats the runner-up by >= 0.5 — with a 0.9 floor that only rules out "no reading matches at all".

Full sweeps at a8ecc98f, all 196 q4_1 tensors each (--q8-samples 100000):

directory	ours corr, min … max	verdict
published FastFlowLM Qwen3-0.6B-NPU2	+0.9890 … +0.9992	rc 0, 0 FAIL — including layers.2.mlp.down_proj at +0.9890 (rel-rms 0.1477), the tensor you flagged
/tmp/wg-foreign (FLM_Q4NX_Converter @ d1d5232)	-0.0856 … +0.1067	rc 1, 196 of 196 refused

So relaxing the floor did not weaken the check — it broadened it: the old code sampled three tensors, this refuses all 196 of the foreign container.

The refusal no longer says one thing for three different situations. It now reports which happened:

• no reading matches (best below 0.9) — wrong GGUF, or an order this repo cannot name. Before, this case printed "NOT in the order, best is ours", the same self-contradiction you flagged, with all three correlations ≈ 0;
• a foreign order it can name — ours +0.0477 vs flm-unsigned +0.9961;
• an ambiguous margin — ours leads but by < 0.5 (a partly foreign container), refused and said so.

Test (14 checks, 0 failures) and a negative control

New checks: a correct container at corr +0.9827 must pass — inside the band the old floor refused — and a container checked against the wrong source must say "no reading of these tiles matches". Negative control, run rather than claimed: restoring ours >= 0.99 in a copy makes exactly that +0.9827 check fail and nothing else.

Also taken (your non-blocking minor): samples are spread evenly across the table instead of the first N in table order; the default three now land on lm_head (skipped — absent from the GGUF), layers.20 and layers.9.

[bong-water-water-bong]

VERDICT: READY at a8ecc98fe3e2a66583da0c827b4f4161e3d803fd. Both blockers from my CHANGES (first at 18cff319, re-posted through 92fea95c) are fixed by a8ecc98f, measured on all 196 q4_1 tensors of every directory involved. One wording nit, not blocking.

Blocker 1, the self-test: tools/check_q4nx_dir.py --self-test → self-test passed, rc 0. The check now looks for the current "no Q8 or q4_1 tile was decoded" message, and the self-test gains a q4_1-selector case and a spread() case. check_q4nx_tile_order_test.py, repack_tile_order_test.py and tile_check_no_shared_tensor_test.py all report 0 failures. copyright.py --check rc 0.

Blocker 2, the 0.99 floor: the order is now decided on the margin (ours − best_other ≥ 0.5), with a 0.9 floor that only means "no reading matches". Full sweeps, --against ~/models/Qwen3-0.6B-Q4_K_M.gguf --q8-samples 100000:

directory	q4_1 tensors	rc	refused	min ours
published FLM Qwen3-0.6B-NPU2 (snapshot b5847992)	196	0	0 (was 1: layers.2.mlp.down_proj +0.9890)	+0.9890
~/models/npu/Qwen3-0.6B (ours, from Q4_K_M)	196	0	0	+0.9955
/tmp/wg-foreign (the #1110 converter)	196	1	196	−0.0856
/tmp/wg-foreign, default 3 samples (now spread() across the table)	2 compared	1	2

So the published directory the fleet serves is no longer falsely refused, and the foreign one is still refused on every tensor. That is the #1110 hole closed without a false positive. spread() also answers my minor point about sampling the first N in table order.

Nit, not blocking: on the foreign container, 168 of the 196 refusals use the new branch no reading of these tiles matches the source weights … Qwen3-0.6B-Q4_K_M.gguf is probably not the GGUF this container was packed from, and 28 say NOT in the order npu/q4nx.cpp decodes. In #1110's own case the GGUF is the source, and the tiles are in an order that varies tensor to tensor (the body says so). An operator reading the first message would go looking for a different GGUF. If the "or the tiles are in an order this repo cannot name" clause led, or the run-level summary said "not the engine's order" whenever any tensor proved foreign, the common reading would be the right one.

[github-actions]

Persistent review updated to latest commit 7ad6f6a