Repository navigation
q4nx-diff: read a container whose weight is padded to a wider tile grid (part of #1110) - #1119
Conversation
bong-water-water-bong
left a comment
There was a problem hiding this comment.
VERDICT: READY at 5ecfbc0 — the padded read is right in the two places it can go wrong, and I ran it on a real container to check it does not invent padding where there is none.
$ q4nx_diff.py decode --gguf ~/models/zyphra-zr1/ZR1-1.5B-bf16.gguf \
--q4nx ~/models/npu/ZR1-1.5B-npu-q4km-v2 --layout ours
layout=ours: exact=0 inexact=197 skipped=58 padded=0
padded=0 is the check I wanted: on a container whose tiles match its own config, the new path does not fire, so it cannot silently "explain" a mismatch as padding. And the code does the right thing where it does fire — decode at the container's grid (dec(raw, rows_p, cols_p)), compare only the model's slice (got[:rows, :cols]), and count it as padded rather than skipped. That is the difference between "we could not read this tensor" and "we read it and it agrees".
The remaining skipped=58 are not this PR's class, which is worth stating because it tells you the skip path is now honest: they are 1-D tensors — len(t.shape) < 2, i.e. the 57 norms (28 layers × 2 + final) plus one more — so nothing tiled is being skipped as a padding-shaped excuse.
Your 84-padded ZR1 row I could not reproduce: it needs the repacked container from the GGUF-on-NPU path, which is not on this box (the v2 pack here is the correct-grid one, which is why it shows padded=0). The mechanism and its non-firing case are verified; the coverage gain is yours.
Ordering: this predates #1116 (no --layout both), and both touch tools/q4nx_diff.py, so it wants a rebase on #1116 and then its +60/−7 as the delta. Together the pair is what makes a padded container measurable (#1119) and correctly labelled (#1116) — which is the gap where the #1110 mismatch hid.
5576165 to
f7de42e
Compare
|
Rebased onto |
bong-water-water-bong
left a comment
There was a problem hiding this comment.
VERDICT: READY at f7de42e — rebase verified by behaviour, the new ctest passes, and the added note is the right content. One nit: it is in the wrong place to do its job.
Rebase did not move the behaviour. Same container pair, same command, same numbers as my last review:
layout=ours: exact=0 inexact=197 skipped=58 padded=0 worst max|err| 7.617e-02 (blk.27.attn_output.weight)
layout=flm : exact=0 inexact=7 skipped=2 padded=0 worst max|err| 2.012e+00 (blk.0.attn_q.weight)
So padded=0 still holds on a correct-grid container (the no-false-positive property), and the gap between the two layouts is ~26× on the same tensors, which is the margin the verdict mechanism depends on. tests/q4nx_diff_grid_test.py → q4nx_diff_grid_test: ok (0.2 s).
The per-layout note is honest and useful — it says which layout is validated against what, with its own two-tensor measurements (ours 2.96e-02 / 2.31e-02; the amd-oss d1d5232 converter matching neither at 5.48e+00 / 3.46e+00), states that --layout flm is the decoder for the pre-q4_k converter, and draws the conclusion an operator needs: a BAD row from flm means "this file is not that layout", not "this file is wrong". The last line — "Deliberately not 'fixed' blind: nothing on this box could tell a fix from a guess" — is the right call and the right thing to write down.
The nit is that the note cannot reach the person who needs it. It lives in the file's header comment block, and neither entry point prints it:
$ q4nx_diff.py --help | grep -c "Which layout is validated" -> 0
$ q4nx_diff.py decode -h | grep -c "Which layout is validated" -> 0
The audience for that paragraph is someone who just ran --layout flm, got a screen of BAD rows and concluded the container is broken. They will not open the source. Put it in the decode subparser's epilog (or description) so it prints with -h, and the note stops being a comment for reviewers and starts being an answer for the operator.
Ordering, still live: this head has --layout {flm,ours} only — I checked by running --help — because it is rebased onto main, where #1116 is not yet merged. So #1116 first, then this rebased onto it: that is what turns "a BAD row from flm means not-this-layout" into a verdict line that says it automatically, instead of a comment a reader has to find.
f7de42e to
e801bb8
Compare
|
Rebased onto current One consistency note added in the third commit ( Head |
bong-water-water-bong
left a comment
There was a problem hiding this comment.
VERDICT: READY at e801bb8 — rebase is clean, behaviour is unchanged on a real container, and the new scope note is a clarification I would have asked for.
Re-verified, not assumed:
$ tests/q4nx_diff_grid_test.py -> q4nx_diff_grid_test: ok
$ q4nx_diff.py decode --gguf ZR1-1.5B-bf16.gguf --q4nx ZR1-1.5B-npu-q4km-v2 --layout ours
layout=ours: exact=0 inexact=197 skipped=58 padded=0 worst max|err| 7.617e-02 <- byte-identical to my last two runs
The third commit is the useful kind of documentation: it says packed_grid() here considers only the 5120-B q4_1 tile because those are the layouts this tool's decoders unpack, that npu/forward's q4nx_packed_grid() tries all three chunk kinds because the forward decodes all three, and that the two agree on every 5120-tile shape with both sides' tests pinning the same geometries and the same tie refusal. That is exactly the right relationship to state: the tool is deliberately narrower than the forward, and a reader now knows why a Q4K/Q8 chunk is "not this tool's container" rather than "not padded".
Two points from my last review still stand, and I want them not to be forgotten:
- The per-layout validation note still does not reach the user. At this head,
q4nx_diff.py decode -h | grep -c "Which layout is validated"→ 0. It is in the module header, and the operator who gets a screen of BAD rows under--layout flmwill not open the source. It belongs in thedecodesubparser'sepilog. --layout {flm,ours}is still the option set (checked at this head), because #1116 is still open.#1116 → #1119-rebasedremains the order that gives this tool the verdict line.
On #1110's remaining half (a declared tile_order), since you flagged it as unblocked: I would take it, and I would take it instead of trusting inference further. The forward now reads a padded grid by inferring it from the tile count — bounded to 64 tiles of excess and refusing ties (which #1107's last commit added, correctly). That is a good guard, but it is still a guess with a refusal attached: a container that states its own tile order cannot be guessed wrong, and it would delete the tie machinery rather than harden it. The evidence that it is worth it is tonight's own history — #1099 warns about padding, #1119 measures it, #1107 serves it, and all three exist because the container did not say.
Not reproduced by me, again, and not for lack of trying: the padded=84 row. That needs the repacked container from the GGUF-on-NPU path; the v2 pack here is the correct-grid one, which is exactly why it reads padded=0. The synthetic geometries in q4nx_diff_grid_test.py are what I can check off-device, and they pass.
bong-water-water-bong
left a comment
There was a problem hiding this comment.
Conflict flag, not a verdict change: this PR is now mergeStateStatus=DIRTY — it conflicts on tools/q4nx_diff.py,.
My VERDICT at the head I reviewed stands on content; nothing about the code has been questioned. But it cannot merge while it conflicts, so it needs a rebase. When it is rebased, my verdict carries if the PR's own files hash the same — that is how I checked the last several rebases (per-file blob identity, not the summary), and I will re-check this one the same way rather than assume.
If the rebase has to touch the conflicting file beyond a mechanical resolution, say so in the PR and I will re-read the hunk instead of carrying the verdict.
`tools/q4nx_diff.py decode` is the checker an operator reaches for when a container answers badly, and it could not read the containers engine#1097 is about: a weight the converter padded to a wider tile grid than the model's own dimensions (the FastFlowLM repack pads intermediate_size to 512 columns, IM 8960 -> 9216) came back skip model.layers.0.mlp.down_proj.weight cannot reshape array of size 8847360 into shape (48,35,5120) for all 84 padded MLP tensors of ZR1-1.5B -- counted as `skipped`, so the padding path had no real-container check at all while npu/forward learned to read it (#1097/#1107). The grid is derived from the tensor's data_size exactly as npu/forward's q4nx_packed_grid() does: the container's tile count pins (rows_packed, cols_packed), the grid must dominate the model's dimensions, it may add at most Q4NX_MAX_EXCESS (64) tiles, and two DIFFERENT grids with the same least excess is a refusal rather than a guess -- that tie is the transposed-decode case. A padded tensor is decoded at the container's grid and only the model's [rows, cols] slice is compared, and the row says which grid was used `[padded 8960x1536 read as 9216x1536]`, so it is never counted as a plain skip. The summary gains `padded=N`; an unpadded container reports `padded=0` and its rows are unchanged. Measured on the two ZR1-1.5B containers on this box (host, no device): our own pack: 197 checked, 58 skipped, padded=0, worst max|err| 5.095e-02 the FLM padded repack: 197 checked, 58 skipped, padded=84, worst max|err| 6.101e-02 blk.0.ffn_down 2.313e-02 [padded 1536x8960 read as 1536x9216] blk.0.ffn_gate 2.230e-02 [padded 8960x1536 read as 9216x1536] i.e. the 84 tensors that were skipped are now measured, at the same error scale as the rest of the container, which is also how the padded decode was confirmed to be the converter's requant rather than a stride error (the padded MLP tensors read ~8x worse than our own pack's, 2.2e-02 vs 2.3e-03 on gate -- worth knowing before "the padded container answers the same" is claimed). tests/q4nx_diff_grid_test.py (ctest `q4nx_diff_grid`, pure python + numpy, no model, no device) pins the grid rule -- both real geometry axes, the on-grid no-op, the tie refusal, a grid smaller than the model, a non-tile size, excess past the cap, an off-grid geometry -- and that a columns- or rows-padded container decodes to exactly the tight values on the model's slice, with the pad decoding to zeros and the padded bytes unreadable at the tight grid (the failure this replaces). Part of #1110; the container-side `tile_order` id in that issue is still open (it needs npu/forward, which #1107 owns this hour).
…ving `flm` looking authoritative
`--layout flm` is the decoder for the *pre-q4_k* official converter, and nothing on this box matches
it consistently, so a BAD row from it reads as "this file is wrong" when it means "this file is not
that layout". Measured today (ZR1-1.5B Q4_K_M source, host, no device, two tensors -- an attn_output
and the padded ffn_down -- on both containers):
third_party/FLM_Q4NX_Converter (npu-rt-753-pr) ours 2.957e-02 / 2.313e-02 <- requant scale
amd-oss/FLM_Q4NX_Converter @d1d5232 ours 5.476e+00 / 6.084e-01 <- noise
flm(signed) 3.455e+00 / 7.680e-01
flm(unsigned) 3.455e+00 / 2.313e-02
So the converter we use emits OUR layout -- `ours` is the gate, and it is what npu/forward reads --
while d1d5232's output matches neither layout consistently: its Q6_K-sourced tensors reach requant
scale only with unsigned nibbles, its Q4_K-sourced ones match none of the three variants. That is
more than one difference, which is the argument for a declared layout id (#1110) rather than
reverse-engineering a converter we do not ship.
Hence: document what is validated, and deliberately do not "fix" the flm decoder blind -- there is no
container on this box that would tell a fix from a guess.
Part of #1110.
e801bb8 to
54d5fff
Compare
|
Rebased onto
Verified on the rebased tree: Nothing of mine changed in meaning — the two features compose, and the verdict is now strictly stronger because it sees |
bong-water-water-bong
left a comment
There was a problem hiding this comment.
VERDICT: READY at 54d5fff — and I re-read this one rather than carrying the verdict, because folding the padded read into _decode_pass is a change of logic, not a rebase. Both features survive, and I checked them separately.
Both halves are now in one pass per layout, which is the right structure:
def _decode_pass(args, layout, quiet=False): # #1116's per-layout pass
…
grid = packed_grid(int(raw.size), rows, cols) # #1119's padded read, inside it
rows_p, cols_p = grid if grid else (rows, cols)
got = dec(raw, rows_p, cols_p) # decode at the container's grid
if grid:
got = got[:rows, :cols] # compare the model's slice
padded = 1; padded_n += 1The alternative — two passes each doing half the job — would have double-read the container and, worse, let --layout both judge the
113 and skip the 84. Running the padded read inside each layout's pass is what makes the verdict speak for all 197.
Checked on a real container (ZR1-1.5B-npu-q4km-v2 against its bf16 GGUF):
$ q4nx_diff.py --self-test -> self-test passed (their #1116 checks)
$ tests/q4nx_diff_grid_test.py -> q4nx_diff_grid_test: ok
$ q4nx_diff.py decode … --layout both
layout=ours: exact=0 inexact=197 skipped=58 padded=0 worst max|err| 7.617e-02
matches layout ours (worst max|err| 7.617e-02, next layout flm worst max|err| 3.199e+00)
Three of my numbers from the last three reviews are unchanged (197/58, padded=0, 7.617e-02), the verdict line names the layout and prints the runner-up, and padded=0 confirms the folded path does not invent padding on a correct container.
Dropping the duplicated preamble was the right call and I checked it cost nothing: _decode_pass now owns hdr/base/mm, dec, the lay map and q4nx_name, and the verdict-level code above it no longer repeats them. The thing that would have gone wrong is a quiet change of args.tensor filtering or the want selection between the two copies; both live in the single pass now, so there is only one place they can be wrong.
Still yours, not mine: the padded=84 / next flm 3.006e+00 row on the repacked container — no padded container on this box (the v2 pack is the correct-grid one, hence padded=0). What I can say is that the code path that produces your row is the path I just exercised, one fixture away.
DIRTY resolved: the branch is now on 55cfaf8 (which carries merged #1116) and my previous two points from earlier reviews are both closed — the padding is measured (#1119's half) and the verdict names the layout with its runner-up (#1116's half). The only thing still open is the older nit that the validation note is header-comment-only and does not print under decode -h.
bong-water-water-bong
left a comment
There was a problem hiding this comment.
VERDICT: READY at c0ff2326 — the move from 54d5fff is a merge-from-main, and the PR's own three files are untouched.
tools/q4nx_diff.py same blob
tests/q4nx_diff_grid_test.py same blob
CMakeLists.txt same blob
compare 54d5fff...c0ff2326 files: docs/kernel.md, registry/architectures.json, third_party/llama.cpp <- #1121, #1120
So my review at 54d5fff is a review of this head's content: both halves verified there (--self-test passing; the real container
giving 197/58/padded=0, worst 7.617e-02, verdict matches layout ours … next flm 3.199e+00), and the folded padded read
sitting inside _decode_pass.
Your row is the one I could not produce, and it is now twice-reported from the same fixture, which I am happy to accept: the
padded ZR1 container at 197 checked / 84 padded / worst 6.101e-02, verdict matches layout ours … next flm 3.006e+00. That
is the same shape as my correct-grid run one fixture away (padded=0, 197 checked, 7.617e-02 vs 3.199e+00 on the
runner-up) — the padded count is the only number that differs, which is exactly what a padded container should change. Your
84 also agrees with the 84 WARN I reproduced independently on #1099 by rewriting a real container's config, so the two
instruments now agree on the same 84 tensors.
Nothing else open on this PR from my side except the older cosmetic nit (the per-layout validation note is header-comment-only and
does not print under decode -h).
…der does not implement is refused (#1130) 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). Co-authored-by: bong-water-water-bong <agent@local>
…against a converter's name (#1109) The comparison this issue asks for already landed inside #1169's tile check: `tile_order_check` runs `q4nx_diff._decode_pass`, which dequantizes the source GGUF (`ref = dequantize(t.data, t.tensor_type)`) and compares every decoded container tensor against it under BOTH layouts. So a wrong checkout is already refused (`matches NO layout`, rc=1) -- what was missing is the contract around it, and the hole that reopened after the refusal. * `--verify-source <dir> <model.gguf>` (and the reusable `verify_source()` it wraps, with `tile_order_check` kept as its old name): the measurement under the name #1109 asks for, usable on an existing or published directory. Exit 0 verified in the engine-1 order, 1 another order or none, 2 nothing comparable. The repack calls the same function after the converter and the config write, so the two cannot disagree. * the verdict now names the SOURCE tensor and its correlation: a wrong checkout reads `matches NO layout: closest is flm with worst max|err| 3.460e+00, corr +0.0325 on blk.14.attn_output.weight`, against an honest pack's `corr +0.9994` (the issue's numbers). The correlation is reported, not gated: no floor was measured on honest containers, and the gate stays the measured max|err| rule. * the refusal names the converter checkout and commit the stamp is about to record -- the one fact repack-stamp.txt has and the user needs when a converter is the wrong version. * `--skip-verify` is the escape the issue's comment asks for (the verify costs a GGUF pass, for the tight loops the payload builds run in). It stamps `tile_order=unverified` with `tile_check=skipped (--skip-verify)` -- never `engine-1`. * `--check-stamp` no longer says `stamp: current` for a directory the repack REFUSED: the refusing path leaves model.q4nx on disk, so before this the next serve reused the container the guard had just rejected -- a wrong-converter container still served after being caught. A stamp that predates the measurement is stale (rebuilt, which is what makes it verified); a stamp that says `unverified` is "cannot tell" (2), so serve reuses it with its existing warning instead of silently, and does not re-run the failed measurement on every start. Deliberately NOT added: a second comparison path (the existing one compares against the source, not just against layouts), a correlation threshold (not measured on honest containers; a guess would refuse them), and attention-only scoping (padded MLP tensors have been readable since #1119/dd5fe38, so scoping would weaken the check the comment asked for). tests/repack_verify_source_test.py: a one-tensor synthetic GGUF, a one-tile container packed in each layout, and a fake converter that copies a prepared container, through the real entry points. It asserts the good container matches with corr +1.0000 and the tensor named, other values match nothing with corr near 0, an FLM-order container matches its source but is refused as this engine's order, `--verify-source` exits 0/1/2, `--skip-verify` stamps unverified and `--check-stamp` returns 2 for it, a real repack stamps `tile_order=engine-1` with the measured verdict, and a refused directory is stale on the next cache hit. Wired into ci.yml's python test step (numpy + gguf installed there). Run: copyright rc 0; the test above 0 failures; q4nx_diff --self-test, q4nx_diff_grid, repack_tile_order, tile_check_no_shared_tensor, npu_plan_tile_order, gguf_to_q4nx all pass.
…r the name that says so (refs #1109) (#1234) * repack: the container is verified against the GGUF it came from, not against a converter's name (#1109) The comparison this issue asks for already landed inside #1169's tile check: `tile_order_check` runs `q4nx_diff._decode_pass`, which dequantizes the source GGUF (`ref = dequantize(t.data, t.tensor_type)`) and compares every decoded container tensor against it under BOTH layouts. So a wrong checkout is already refused (`matches NO layout`, rc=1) -- what was missing is the contract around it, and the hole that reopened after the refusal. * `--verify-source <dir> <model.gguf>` (and the reusable `verify_source()` it wraps, with `tile_order_check` kept as its old name): the measurement under the name #1109 asks for, usable on an existing or published directory. Exit 0 verified in the engine-1 order, 1 another order or none, 2 nothing comparable. The repack calls the same function after the converter and the config write, so the two cannot disagree. * the verdict now names the SOURCE tensor and its correlation: a wrong checkout reads `matches NO layout: closest is flm with worst max|err| 3.460e+00, corr +0.0325 on blk.14.attn_output.weight`, against an honest pack's `corr +0.9994` (the issue's numbers). The correlation is reported, not gated: no floor was measured on honest containers, and the gate stays the measured max|err| rule. * the refusal names the converter checkout and commit the stamp is about to record -- the one fact repack-stamp.txt has and the user needs when a converter is the wrong version. * `--skip-verify` is the escape the issue's comment asks for (the verify costs a GGUF pass, for the tight loops the payload builds run in). It stamps `tile_order=unverified` with `tile_check=skipped (--skip-verify)` -- never `engine-1`. * `--check-stamp` no longer says `stamp: current` for a directory the repack REFUSED: the refusing path leaves model.q4nx on disk, so before this the next serve reused the container the guard had just rejected -- a wrong-converter container still served after being caught. A stamp that predates the measurement is stale (rebuilt, which is what makes it verified); a stamp that says `unverified` is "cannot tell" (2), so serve reuses it with its existing warning instead of silently, and does not re-run the failed measurement on every start. Deliberately NOT added: a second comparison path (the existing one compares against the source, not just against layouts), a correlation threshold (not measured on honest containers; a guess would refuse them), and attention-only scoping (padded MLP tensors have been readable since #1119/dd5fe38, so scoping would weaken the check the comment asked for). tests/repack_verify_source_test.py: a one-tensor synthetic GGUF, a one-tile container packed in each layout, and a fake converter that copies a prepared container, through the real entry points. It asserts the good container matches with corr +1.0000 and the tensor named, other values match nothing with corr near 0, an FLM-order container matches its source but is refused as this engine's order, `--verify-source` exits 0/1/2, `--skip-verify` stamps unverified and `--check-stamp` returns 2 for it, a real repack stamps `tile_order=engine-1` with the measured verdict, and a refused directory is stale on the next cache hit. Wired into ci.yml's python test step (numpy + gguf installed there). Run: copyright rc 0; the test above 0 failures; q4nx_diff --self-test, q4nx_diff_grid, repack_tile_order, tile_check_no_shared_tensor, npu_plan_tile_order, gguf_to_q4nx all pass. * docs/npu: --verify-source and --skip-verify are the names a reader needs, and a pre-measurement stamp is rebuilt once (#1109) --------- Co-authored-by: bong-water-water-bong <bong-water-water-bong@1bit.gg>
tools/q4nx_diff.py decodeis the tool an operator reaches for when a container answers badly, and it could not readthe containers #1097 is about. A weight the converter padded to a wider tile grid than the model's dimensions (the
FastFlowLM repack pads
intermediate_sizeto its 512-column grid, IM 8960 → 9216) came back asfor all 84 padded MLP tensors of ZR1-1.5B (28 layers × gate/up/down) — counted as
skipped, quietly. While I wasre-measuring #1107 that meant the padding path had no real-container check at all: the unit test there pins a
synthetic grid, and the per-tensor checker an operator runs reported 113 checked / 142 skipped on the padded container
versus 197 checked on our own pack, so the tensors that were broken were the ones it could not look at.
What it does now
The grid is derived from the tensor's
data_sizeexactly asnpu/forward'sq4nx_packed_grid()does (#1097/#1107):the container's tile count pins
(rows_packed, cols_packed), the grid must dominate the model's dimensions, it may addat most
Q4NX_MAX_EXCESS(64) tiles, and two different grids with the same least excess is a refusal — that tie isthe transposed-decode case. A padded tensor is decoded at the container's grid and only the model's
[rows, cols]sliceis compared, and the row says which grid was used, so a padded tensor can never be counted as a plain skip:
The summary gains
padded=N; an unpadded container reportspadded=0and its rows are byte-for-byte as before.Measured (host, no device, no ticket)
padded=84, worst 6.101e-02Two things that fell out of being able to see those tensors, both worth recording:
the rest of the container (2.3e-02 … 3.4e-02), so npu: read a weight grid the container padded wider than the model's own (#1097) #1107's read-at-the-container's-grid is right.
ffn_gate2.230e-02 vs 2.486e-03).That is why the padded container's text differs from our own pack's at 8 tokens even though npu: read a weight grid the container padded wider than the model's own (#1097) #1107 fixes the decode —
the "same text as our own pack" claim needs that caveat, and I said so on npu: read a weight grid the container padded wider than the model's own (#1097) #1107 rather than leaving it as a mismatch.
Tests
tests/q4nx_diff_grid_test.py— ctestq4nx_diff_grid, pure python + numpy, no model, no container, no device (sameshape as
q4nx_pb16i):tiles → 9216×1536), the on-grid no-op, the tie refusal (a 32×32 model in 2 tiles fits (1,2) and (2,1) equally),
a grid smaller than the model, a size that is not whole tiles, excess past the cap, an off-grid geometry;
as the tight one on the model's slice, the pad decodes to zeros, and the padded bytes are unreadable at the tight grid
— the failure this replaces.
ctest -R q4nx_diff_grid→Passed 0.22 sec;tools/copyright.py --checkrc 0.Part of #1110 — the container-side
tile_orderid from that issue is not here: it needsnpu/forward, which #1107owns this hour, and I will take it once that merges.