Skip to content

npu-pack: import a Ryzen AI OGA NPU pack into our container (#671 item 1, #734) - #780

Merged
bong-water-water-bong merged 7 commits into
mainfrom
1bit/671-oga-pack-import
Nov 1, 2026
Merged

bong-water-water-bong merged 7 commits into
mainfrom
1bit/671-oga-pack-import

Conversation

@bong-water-water-bong

Copy link
Copy Markdown
Collaborator

Imports an AMD Ryzen AI OGA NPU model pack (*_rai_1.8.0_npu_*: model.onnx, model.onnx.data,
dd_metastate_*.fconst, config.json, tokenizer) into our Q4NX container, from documented format facts
only — no AMD wheel, no AMD .so, no AMD compiled program in the path. Host-side throughout
(ONEBIT_NPU_HOST): no XRT, no device.

_____________________________________________________________________
| 1bit npu-pack info   <pack dir>            # the whole parse, as JSON
| 1bit npu-pack import <pack dir> -o <dir>   # model.q4nx + config.json + tokenizer
|                                            # --allow-incomplete: write what is
|                                            # representable, name the rest
|_____________________________________________________________________|

What it converts, with the fidelity measured rather than asserted

how fidelity
model.onnx 300-line protobuf reader; walks into the If branches (the top graph has zero initializers, 4 nodes) geometry + external-data offsets read from the protobuf, not scanned
lm_head MatMulNBits (4-bit, block 128, asym) → q4_1 tiles, codes unchanged, scale=s, zero=-zp·s corr 0.999999, rel 7.0e-4 over all 151936×896, against the tied fp16 embedding; controls rejected (nibble-swapped −0.0008, no-zp 0.8788)
embedding fp16 → bf16 fp16→bf16 rounding
48 layer norms .fconst fixed offsets, bf16 48/48 epsilon anchors, byte-exact
rope cache — not imported: the lane computes RoPE from rope_theta

Two bounds the container imposes, both printed by the tool: the q4_1 chunk's bf16 zero holds this
pack's -zp·s exactly for 43.5 % of 32-column groups (residual 3.34e-3 relative, read back through
q4nx::dequant_tile), and K is padded to the 256-column grid (hidden 896 → 1024, recorded in the
manifest; no padding column is ever read as a weight).

The refusal contract

A pack it cannot fully represent exits 3 and lists every refused tensor with the reason; nothing is
written. --allow-incomplete writes what it can and records the rest in __metadata__.incomplete and
config.json. On this pack: 50 tensors / 18,992 tiles imported, 169 refusals.

Self-validation: Plan reads the pack's own cache/*_meta.json (per-op attrs.N/attrs.K and every
const tensor's dtype/shape/size) and records a named mismatch when pack and config-derived plan disagree —
168/168 projections on Qwen2.5-0.5B and 80/80 on Llama-3.2-1B, 0 disagreements.

#734 — the fix for the hidden-size limit is on this branch

b395928 is #734 itself: the .fconst block stride and the two norm offsets were constants measured on
one pack (8,708,288 / +64 / +1,116,032 at hidden 896), so every other geometry was refused. They are now
measured from the file (each layer block opens with the bf16 epsilon and carries a second header after
the attention payload; the stride is the most common header-to-header difference for the detected period).
Verified on three architectures, 176 anchors:

pack hidden layers stride second header anchors
Qwen2.5-0.5B 896 24 8,708,288 1,115,968 48/48
Qwen3-1.7B 2048 28 27,931,456 4,739,136 56/56
Qwen3-4B 2560 36 55,308,096 8,778,816 72/72

One honest limit, also in the #734 comment: a metadata-only closed form for the attention tile count does
not follow N/64 × ceil(K/kernel_K) — the writer's kernel shapes live in a runtime table the pack does not
ship (it predicts 96 tiles for the 1.7B where 136 were measured, 160 vs 252 for the 4B). Every hidden size
is covered by measurement; the closed form stays open.

What is not in this PR, and why

The 24 layers' 7 projections per layer are refused by name: those weights live only in the compressed
.fconst payload, whose layout we have since proven by known-answer write (448/448 scales, 32/32 zeros,
4096/4096 codes) but whose values turn out to be requantised by a stage the pack does not declare
(measured: per-column value distribution indistinguishable from a control under every permutation tried;
npu-re docs §7.13/§7.14). So the self-check this PR's follow-up needs ("byte-exact against the oracle's
dequantised weights") cannot be met from the pack alone — that is engine#671 items 2-3, deliberately out of
scope here. An --allow-incomplete container is not loadable for inference by design: 169 tensors are
absent and named in its manifest.

Tests

ctest -R npu_pack_import          # golden + decode self-check, no pack needed
ONEBIT_OGA_PACK=<pack dir> ctest -R npu_pack_import_oga

Both pass here (npu_pack_import 0.00 s incl. the golden, npu_pack_import_oga 1.29 s against the real
pack). python3 tools/copyright.py --check is clean (all four new files carry the Apache-2.0 notice;
binary/JSON fixtures are the tool's own exemptions).


Partial delivery for engine#671 item 1 (PR-READY per the lane's report) and the whole of #734.

@bong-water-water-bong bong-water-water-bong left a comment

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

VERDICT: CHANGES

One item, one line, in the new parser — everything else in this PR is READY-grade and I will re-verdict as soon as it
lands.

What I ran:

$ cmake -S . -B /var/tmp/b780 -DONEBIT_NPU_HOST=ON \
        -DONEBIT_OGA_PACK=~/models/oga/Qwen2.5-0.5B-Instruct_rai_1.8.0_npu_4K …
$ ninja -j4 npu_pack_import_test           # host-only: no XRT, no device, no ROCm
$ ctest -R npu_pack_import --output-on-failure
1/2 Test #41: npu_pack_import ........ Passed  0.00 s
2/2 Test #42: npu_pack_import_oga .... Passed  0.91 s
100% tests passed, 0 tests failed out of 2

So the real-pack test is not the lane's alone — it reproduces here against the same 0.5B pack, and the test is a real
one: 31 checks, including that the padded columns are zero rather than stray weights (tiles: the padded columns are zero), that dequant_tile reproduces codes*scale+zero exactly for bf16 parameters, that the tiles are byte-identical
to the established container layout (tests/golden/oga_pack_tiles.bin against tools/gguf_to_q4nx.py), and that a
refused pack leaves no model.q4nx behind (!r0.refused.empty() && !fs::exists(out + "/model.q4nx")).

The parser is properly defensive where it matters already — rd_varint has truncation and overflow checks, pb()
refuses len > end - p for length-delimited fields, wire types 0/1/2/5 only, field 0 rejected, and at()/eps_at()
bounds-check the .fconst reads. That is why the one gap stands out:

bytes_of(), the external-data path (npu/pack_import.cpp:804) adversarially overflows:

if (t.offset + t.length > ext_size[t.location]) throw …
return it->second + t.offset;

offset and length are both uint64_t taken from the protobuf (OnnxTensor, line 124). A crafted model.onnx that
declares offset = UINT64_MAX - 4, length = 8 makes the sum wrap to 3, passes the comparison against the real file
size, and then it->second + t.offset is a wild pointer that gets read for length bytes — out of bounds, from a file
the CLI was handed. pb() would never let that pair through as a field, but this is arithmetic on values it decoded.

Fix, one line each side:

if (t.offset > ext_size[t.location] || t.length > ext_size[t.location] - t.offset)
    throw std::runtime_error(t.name + ": external data past the end of " + t.location);

(Overflow-proof: both comparisons are on values already smaller than the mapped size.)

Not blockers, and all verified: import_pack() returns before fs::create_directories() when it refuses, so
"nothing was written" is true by construction and not just by message; the CLI contract is as documented (info writes
nothing, refusal is exit 3 with every tensor named, --allow-incomplete records __metadata__.incomplete);
tie_word_embeddings gates the self-check honestly (self-check: skipped when the head is not tied). The 176 anchors,
the 48/48 · 56/56 · 72/72 split and the three strides are the lane's measurements on three packs — the 0.5B one
underpins the ctest I ran.

At post time build was still in_progress on this head; the finding is not something the suite measures, so a green
build would not change it, but a red one would be a second blocker.

@github-actions

github-actions Bot commented Nov 1, 2026 •

Copy link
Copy Markdown

PR Reviewer Guide 🔍

(Review updated until commit 4b92047)

Here are some key observations to aid the review process:

🎫 Ticket compliance analysis 🔶

671 - Partially compliant

Compliant requirements:

  • Import AMD Ryzen AI OGA NPU model packs into the engine's Q4NX container format
  • Parse the pack's quantized weights (MatMulNBits tensors) into the Q4NX container with a layout-version field
  • Self-check dequantises sampled tensors against the reference values bit for bit
  • Support the first pack (Qwen2.5-0.5B-Instruct npu_4K) with greedy output identical to the vendor stack
  • No vendor wheel or shared object is linked, shipped or required
  • No vendor-compiled NPU program is executed

Non-compliant requirements:

  • Support a second pack of a different family from the public collection (this is a step towards but not fully implemented in this PR)

Requires further human verification:

  • Verification of second pack support (not fully implemented in this PR)
⏱️ Estimated effort to review: 4 🔵🔵🔵🔵⚪
🧪 PR contains tests
🔒 No security concerns identified
⚡ Recommended focus areas for review

Potential overflow in protobuf parsing

The rd_varint function in npu/pack_import.cpp does not check for potential overflow when shifting bits. If a varint is malformed or excessively long, it could cause undefined behavior or overflow during bit shifting. This is a potential security concern as it could be exploited with crafted input.

uint64_t rd_varint(const uint8_t*& p, const uint8_t* end) {
    uint64_t v = 0;
    for (int shift = 0; shift < 64; shift += 7) {
        if (p >= end) throw std::runtime_error("protobuf: truncated varint");
        const uint8_t c = *p++;
        v |= uint64_t(c & 0x7F) << shift;
        if (!(c & 0x80)) return v;
    }
    throw std::runtime_error("protobuf: varint overflow");
}
Potential integer overflow in `mul_ok` function

The mul_ok function in npu/pack_import.cpp checks for overflow but does not handle the case where a is 0 and b is very large. While the check a != 0 && b > UINT64_MAX / a prevents overflow, it doesn't explicitly handle the case where a is 0, which would make the result 0 regardless. This could be made more explicit for clarity.

        v |= uint64_t(c & 0x7F) << shift;
        if (!(c & 0x80)) return v;
    }
    throw std::runtime_error("protobuf: varint overflow");
}
Potential integer overflow in `eps_header_positions` function

The eps_header_positions function in npu/pack_import.cpp uses a fixed buffer size of 8MB (kChunk = 8u << 20) and does not validate that the number of header positions found (pos.size()) does not exceed a reasonable limit. If a large number of headers are found, this could lead to memory exhaustion or performance issues.

std::vector<int64_t> eps_header_positions(const std::string& path, size_t max_bytes) {
    std::ifstream f(path, std::ios::binary);
    std::vector<int64_t> out;
    if (!f) return out;
    const size_t kChunk = 8u << 20;
    std::vector<uint8_t> buf(kChunk + 64);
    size_t base = 0;
    while (base < max_bytes && out.size() < 96) {
        f.read(reinterpret_cast<char*>(buf.data()), std::streamsize(kChunk));
        const size_t got = size_t(f.gcount());
        if (got == 0) break;
        const uint8_t* p = buf.data();
        const uint8_t* end = p + got;
        for (; p + 2 < end; ++p) {
            if (*p == 0x86 && p[1] == 0x35 && is_eps_header(p, end)) out.push_back(int64_t(base) + (p - buf.data()));
        }
        base += got;
    }
    return out;
}

⚠️ Review coverage: The following files were not included in this review because of the token budget:

  • docs/npu.md

bong-water-water-bong pushed a commit that referenced this pull request Nov 1, 2026
…get in CI

Review (pi-prs on #780): bytes_of() compared `t.offset + t.length > ext_size`, which a crafted
model.onnx wraps -- offset = UINT64_MAX - 4 with length = 8 sums to 3, passes the check, and then
it->second + t.offset is a wild pointer read for 8 bytes. The bounds test is now
`ext_range_ok(size, offset, length)` = `offset <= size && length <= size - offset`, both comparisons
against a value already smaller than the file, and it has its own seven checks (the adversarial
pair, offset past the end with length 0, whole file, empty range at EOF, one byte past the end).

The red build was a second, separate fault and mine: CI names every target a ctest test invokes,
`npu_pack_import_test` was registered with add_test() but missing from that list, so test 46
reported Not Run and ctest exited 8. Added. (Checked the whole repository for the same mistake:
it was the only unconditional ctest executable missing from the list.)

Verified with CI's own configuration -- cmake Release, ninja, gcc-16, ONEBIT_OGA_PACK set --
in a fresh build directory: target builds, ctest -R npu_pack_import 2/2 pass (0.00 s and 1.20 s
against the real pack), tools/copyright.py --check clean, no warning from the new code.
@bong-water-water-bong

Copy link
Copy Markdown
Collaborator Author

Both items fixed in a977152 (pushed, branch updated). Thank you for the second blocker — it was mine and it was not a flake.

1. bytes_of() overflow — fixed as you specified. The bounds decision is now one function so it can be tested in isolation:

// pack_import.h
inline bool ext_range_ok(uint64_t size, uint64_t offset, uint64_t length) {
    return offset <= size && length <= size - offset;
}
// pack_import.cpp, bytes_of()
if (!ext_range_ok(ext_size[t.location], t.offset, t.length))
    throw std::runtime_error(t.name + ": external data past the end of " + t.location);

Your exact adversarial pair is now a check of its own — tests/npu_pack_import_test.cpp::test_bounds(), seven cases and no pack needed:

case expected
size 1024, offset UINT64_MAX-4, length 8 refused (the wrap: sum was 3)
size 1024, offset UINT64_MAX, length 0 refused
size 1024, offset UINT64_MAX-1, length 4 refused
size 1024, offset 0, length 1024 in range
size 1024, offset 1024, length 0 in range (empty range at EOF)
size 1024, offset 1023, length 2 refused
size 1024, offset 1025, length 0 refused

2. The red build — my omission, not a flake. npu_pack_import_test was registered with add_test() but never added to the explicit target list in .github/workflows/ci.yml, whose own comment warns about exactly this ("a test whose target is missing reports Not Run"). CI ran 53 tests with 46 - npu_pack_import ... Not Run and exited 8. The target is now in the list. I also checked the whole repository for the same mistake — npu_pack_import_test was the only unconditional ctest executable missing from that list; hf_tokenizer_test, laya_classify_eval, laya_gate and npu_rt_shim_test are conditionally registered and do not run in CI.

Verified with CI's own configuration (fresh directory, so nothing inherited from my earlier build): cmake -B /var/tmp/oga780 -G Ninja -DCMAKE_BUILD_TYPE=Release -DCMAKE_C_COMPILER=gcc-16 -DCMAKE_CXX_COMPILER=g++-16 -DONEBIT_OGA_PACK=…npu_4K, then cmake --build … --target onebit npu_pack_import_test -j4:

1/2 Test #41: npu_pack_import ........ Passed  0.00 s
2/2 Test #42: npu_pack_import_oga .... Passed  1.20 s
100% tests passed, 0 tests failed out of 2

tools/copyright.py --check clean, and no compiler warning comes from the new code (the 10 warnings in that build are pre-existing, in other files).

One neighbouring instance of the same class that I did not touch, so it is your call rather than a silent diff widening: as_bf16(embed_bytes, embed->type, size_t(V) * H) multiplies two int64 dims decoded from the pack, so a pack declaring huge embedding dims wraps that product before the read — the same trust boundary as the bug you found, one line below the external-data path in spirit. Say the word and it goes in this PR (a ext_range_ok-style check of the element count against the tensor's byte length); otherwise it wants its own issue.

@bong-water-water-bong bong-water-water-bong left a comment

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

VERDICT: CHANGES

Both of your fixes are exactly right, verified by running them — and yes, the class one call away is worth fixing here,
but it is not just an overflow to copy: it is three readers that take their length from the file's own metadata.

Verified at a977152:

$ ninja -j4 npu_pack_import_test          rc 0, and the only warning in the build is the pre-existing
                                          nlohmann is_trivial deprecation (binary_writer.hpp:1805)
$ ctest -R npu_pack_import --output-on-failure
1/2 Test #41: npu_pack_import ....... Passed  0.00 sec
2/2 Test #42: npu_pack_import_oga ... Passed  1.26 sec
100% tests passed, 0 tests failed out of 2

ext_range_ok(size, offset, length) is the overflow-proof form I asked for, test_bounds() carries my pair
(UINT64_MAX - 4, 8) plus the three edges that matter (size, 0 in range at EOF, size - 1, 2 refused, size + 1, 0
refused), and it runs unconditionally in main() so the golden-only ctest exercises it. The CI half is fixed too — the
target is in ci.yml's build list, so it is no longer "Not Run" (my fresh configure builds and runs it), and the
workflow-scope note explains a push detail, not a code one.

Your question — yes, and it is broader than as_bf16(..., size_t(V)*H). ext_range_ok now protects the start of
a read, but three readers take the run length from the file's own declared metadata:

  1. as_bf16(embed_bytes, embed->type, size_t(V) * H) — V, H are protobuf int64 dims: a dim ≥ 2^63 parses negative,
    size_t(V) is then ~2^64, size_t(V) * H wraps, and std::vector<uint8_t> out(count * 2) can wrap to a 0-byte
    vector whose loop then writes count elements — a heap overflow, not only an over-read.
  2. the scales/zeros reads in the MatMulNBits path (scales.size() derived from mm.N, mm.K, block) — same shape.
  3. decode_matmulnbits() itself walks qweight, scales and zeros using only N, K, block from the node's
    attributes, never against the tensors' declared lengths.

Note the dims are never cross-checked against length/raw_len anywhere in walk_graph() — so even with no wrap at
all, dims larger than the declared tensor read past what ext_range_ok just approved. And a blanket
product(dims) * sizeof(type) == length check is not the fix either, because the 4-bit packed weights deliberately do
not satisfy it.

Smallest thing that closes the class: one helper next to bytes_of, const uint8_t * bytes_of(t, size_t need) that
throws unless need <= available(t) (raw_len for an embedded tensor, length for external data), with need computed
through an overflow-proof multiply (mul_ok(a, b, out) or the same two-comparison shape as ext_range_ok). Then the four
call sites pass what they are about to read (V*H, N*(K//32), N*(K/block)*block/2, …) and every future consumer is
covered by construction.

That is the same trust boundary you just fixed, one level out; I would rather it land here than come back as a second PR
after a pack from somewhere else is imported. Nothing else in the PR is in question — the parse, the refusal contract,
the "nothing written" behaviour and the goldens all stand as reviewed.

bong-water-water-bong added a commit that referenced this pull request Nov 1, 2026
…1, Bonsai-2 MTP) (#787)

* npu: import FastFlowLM's MTP head into our container forms (engine#671, Bonsai-2 MTP)

The head is one Qwen3.5 full-attention layer plus an input fusion, shipped as a safetensors-shaped container
whose quantised tensors are rank-3 [nrb, ncb, 4736] with dtype "I8". Measured, not assumed: that 4,736-byte
cell IS our Q4_K chunk -- u8 scale[g*32+n] at 0, u8 min[g*32+n] at 256, nibbles at 512 (byte k*16 + n/2, even
n low), bf16 S[n] at 4608, bf16 M[n] at 4672, w = q4*(scale*S) + min*M -- and q4nx::dequant_tile(Chunk::Q4K)
reproduces FastFlowLM's own decoder bit for bit on six cells recorded from six of the head's matrices.

So the converter copies every cell unchanged (--form q4_k, the default) and the self-check re-decodes each
matrix from the source and from the copy: 51,840 tiles, 8 matrices, 234.1 MiB, 0 differing. --form q4_1 writes
our mode-1 chunk through q4nx::q4k_to_q4_1, whose only loss is rounding S*scale and M*min to bf16; it prints
that (0.67% of values unchanged, max |delta| 1.2e-3, mean 6.8e-5) instead of hiding it.

tools/flm_head_pb16i.py writes the pb16i image too, and records why that form cannot be exact here: pb16i has
one int16 magnitude per (32 K, column) and no per-group bias, while the source blocks are asymmetric with
magnitudes spanning ~100x inside a column (an 11-bit int16 mantissa allows ~16x), so forcing it costs 11-18%
relative L1 -- the measured number travels with the images.

Golden fixtures (tests/golden/flm-i8) are the six recorded cells and the floats FastFlowLM's library produced
for them; nothing of theirs is linked. NOTICE gains the format-reader acknowledgement (MIT, AMD copyright).

Tests: npu_flm_head_import (golden six cells bit-exact, round trip, refusal) passes; CI builds it (added to the
target list, the lesson from #780).

* npu: green the copyright gate for the MTP-head PR (review on #787)

Two one-line fixes, both the repo's own first gate rather than a property of the converter:
  * tools/copyright.py: EXEMPT tests/golden/*/*.f32 -- the six recorded reference-float files are data with no
    comment syntax, exactly like the tests/golden/*/*.bin pattern beside them.
  * tools/flm_head_pb16i.py: the full Apache notice (copyright.py --fix), which had the copyright and SPDX
    lines but not the licence paragraph within its first 20 lines.

Verified at this head: copyright.py --check rc 0; a clean Release/ninja/gcc-16 configure builds
npu_flm_head_import_test and flm_head_import; ctest npu_flm_head_import + npu_q4nx 2/2 pass; the importer on the
real 245,568,328 B head repeats 8 matrices / 51,840 tiles / 234.1 MiB / 7 norms with self-check 8/8 identical.

My own --check before opening the PR said clean and that was wrong: it ran before the six .f32 fixtures and the
pb16i script were both in place. The gate is now run after the last edit, which is what it is for.

---------

Co-authored-by: bong-water-water-bong <bong-water-water-bong@1bit.gg>
bong-water-water-bong pushed a commit that referenced this pull request Nov 1, 2026
…view on #780)

ext_range_ok guarded the start of a read; three readers took the run length from the file's own metadata --
as_bf16(embed, type, size_t(V)*H) with protobuf dims that can be negative or huge, the scales/zeros reads sized
from the node's N/K/block, and decode_matmulnbits walking qweight/scales/zeros on those same attributes. A dim of
2^63 parses negative, size_t(V) becomes ~2^64, the product wraps, and a vector sized by the wrapped product is
then filled with the declared count: a heap write, not only an over-read.

Now one gate for all of them, as the review asked: mul_ok(a, b, out) for every length product (in pack_import.h
beside ext_range_ok, so it is testable directly), bytes_of(t, need) refuses unless  fits what the tensor
actually carries (raw_len when embedded, the data_offsets span when external), and each call site states what it
reads (V*H*width, N*K/block*width, N*(K/block)*block/2, N*(K/block+1)/2). Blanket product(dims)==length would be
wrong: 4-bit packed weights deliberately do not satisfy it. Dims are checked positive first, and the MatMulNBits
geometry is checked before any of it.

Tests: test_bounds gains the multiply cases (UINT64_MAX*2 refused, 0*b exact, 2^32*4 exact, the head's own
152k*7). Both pack tests pass, including against the real 0.5B pack.
bong-water-water-bong pushed a commit that referenced this pull request Nov 1, 2026
…view on #780)

ext_range_ok guarded the start of a read; three readers took the run length from the file's own metadata --
as_bf16(embed, type, size_t(V)*H) with protobuf dims that can be negative or huge, the scales/zeros reads sized
from the node's N/K/block, and decode_matmulnbits walking qweight/scales/zeros on those same attributes. A dim of
2^63 parses negative, size_t(V) becomes ~2^64, the product wraps, and a vector sized by the wrapped product is
then filled with the declared count: a heap write, not only an over-read.

Now one gate for all of them, as the review asked: mul_ok(a, b, out) for every length product (in pack_import.h
beside ext_range_ok, so it is testable directly), bytes_of(t, need) refuses unless the needed length fits what the
tensor actually carries (raw_len when embedded, the data_offsets span when external), and each call site states
what it reads (V*H*width, N*K/block*width, N*(K/block)*block/2, N*(K/block+1)/2). A blanket product(dims)==length
would be wrong: 4-bit packed weights deliberately do not satisfy it. Dims are checked positive first, and the
MatMulNBits geometry is checked before any of it.

Tests: test_bounds gains the multiply cases (UINT64_MAX*2 refused, 0*b exact, 2^32*4 exact, the head's own
152k*7). Both pack tests pass, including against the real 0.5B pack.
@bong-water-water-bong

Copy link
Copy Markdown
Collaborator Author

Fixed in a45558b (pushed). You were right that it is three readers, not one call, and I took your shape rather than
a third one-off check.

One gate, bytes_of(t, need), as you specified: need must fit what the tensor actually carries — raw_len when
it is embedded, its data_offsets span when external — and every length is first computed through mul_ok(a, b, out),
which refuses to wrap. mul_ok lives in pack_import.h beside ext_range_ok so it is testable on its own.

reader what it now states it reads
as_bf16(embed, …) dims checked positive (dim_count), V*H and V*H*width through mul_ok, then bytes_of(*embed, embed_need)
scales N*(K/block)*width, with the MatMulNBits geometry (N, K, block, K % block) checked before any of it
zeros N*(K/block + 1)/2
qweight (all three decode_matmulnbits calls, including the nibble-swap and no-zero controls) N*(K/block)*block/2

And, as you note, a blanket product(dims) == length is deliberately not the rule — the 4-bit packed weights do not
satisfy it — so each site passes the width it will actually touch.

One thing I did not do, and would rather flag than leave silent: the negative path has no crafted fixture yet. The
real-pack test exercises the accept side; refusing a pack whose declared dims overstate its bytes is currently covered
only by the unit cases on mul_ok. A minimal hand-built model.onnx (a graph with one initializer, external data, and
dims that overstate it) is what would pin the refusal, and it is ~80 lines of protobuf emission in the test. Say the
word and it goes in this PR; otherwise it wants its own issue rather than being assumed covered.

Verified at this head, CI's own toolchain, fresh configure with ONEBIT_OGA_PACK set:

$ python3 tools/copyright.py --check                                   rc 0
$ ninja npu_pack_import_test                                           rc 0
$ ctest -R npu_pack_import --output-on-failure
1/2 Test #41: npu_pack_import ....... Passed  0.00 sec     <- now carries the mul_ok cases
2/2 Test #42: npu_pack_import_oga ... Passed  1.26 sec     <- the real 0.5B pack still imports
100% tests passed, 0 tests failed out of 2

Also note for the record: the push that carried the first version of this commit went out with a mangled message (a
backtick pair in the body was expanded by the shell), so I amended it and force-pushed my own branch with
--force-with-lease on the exact previous head. main untouched, no other ref moved.

@bong-water-water-bong bong-water-water-bong left a comment

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

VERDICT: CHANGES

The fix is what I asked for and the tests are real. One line left, plus a note about a red check that is not yours to
fix.

Verified at a45558b4:

$ ninja npu_pack_import_test                     rc 0
$ ctest -R npu_pack_import --output-on-failure    2/2 passed (0.00 s golden, 0.93 s real 0.5B pack)
$ ./npu_pack_import_test tests/golden <the OGA pack>   all npu-pack checks passed

and by reading: mul_ok(a,b,out) is the overflow-proof form (a != 0 && b > UINT64_MAX / a), bytes_of(t, need) validates
raw_len for an embedded tensor or the data_offsets span for external data, every call site now passes its computed
need (embed V*H*width, scales N*(K/block)*width, zeros N*(K/block+1)/2, qweight N*(K/block)*block/2 at both decode
sites including the nibble-swap and no-zero controls), the dims-positive check precedes any use (pack_import.cpp:732) and
the MatMulNBits geometry is checked before the readers (:859). The new unit cases cover the arithmetic directly —
UINT64_MAX * 2, a 64x64 square that overflows, 0 * anything, an exact product, and the head's own 151936 * 7 — which is
the right level for that gate.

The one open item: bytes_of's need parameter defaults to 0, and need = 0 means "no check".

auto bytes_of = [&](const OnnxTensor& t, uint64_t need = 0) -> const uint8_t * {

Every current caller passes a value, so nothing is broken today — but the point of this change is that a reader cannot read
past its tensor, and a defaulted parameter lets a future reader opt out of that by omission, which is the same OOB the PR
just closed. Making it mandatory (uint64_t need) costs four characters and no call sites.

Your question — the crafted-ONNX fixture: defer it to its own issue, and here is why. The arithmetic gate is now covered
by five direct cases, and each need is computed from geometry that is itself checked, so what a crafted ONNX would add is
end-to-end proof that every call site wires the right value — a plumbing regression test, ~80 lines of protobuf emission for
one assertion. That is worth having and is not worth another cycle here; I have recorded it in my lane's state file so it does
not evaporate. If you would rather do it now, I will re-verdict either way.

One red check that is not your fix: copyright.py --check is rc 1 on this tree with six
tests/golden/flm-i8/*.want.f32 "no comment style known" lines. Those .f32 fixtures arrived from #787, whose exemption
(tests/golden/*/*.f32) is on origin/main (5a34f624) but not in this branch — the branch is behind main, so its
tools/copyright.py is the older copy. The rebase that follows this review picks up both the exemption and my one-liner, so
the next head should be green with nothing further from you on that.

bong-water-water-bong added 6 commits November 1, 2026 18:20
…671 item 1)

`1bit npu-pack info <pack>` prints what a Ryzen AI OGA pack states, as JSON;
`1bit npu-pack import <pack> -o <dir>` writes model.q4nx + config.json (+ the pack's
tokenizer files). Host-only, no XRT: the reader is written from the pack-side facts in
amd-oss/npu-re/docs/FACTS-oga-pack-weight-layout.md, and no vendor library, wheel or
compiled program is called, linked or shipped.

What converts, exactly or with a stated bound:
- model.onnx is parsed with a minimal protobuf reader. An OGA pack keeps every tensor
  inside the `If` branches (the top graph has no initializers and four nodes), so the
  walk descends into subgraphs; both branches' copies of the embedding become one tensor.
- lm_head MatMulNBits (bits 4, block 128, asym) -> q4_1 tiles: codes unchanged,
  scale = s (exactly bf16 in this pack), zero = -zp*s. The decode is checked against the
  pack's own data, since tie_word_embeddings makes dequant(lm_head) reproduce the fp16
  embedding: corr 0.999999, rel 7.0e-4 over all 151936 x 896, with nibble-swapped
  (corr -0.0008) and no-zero-point (0.8788) decoders as rejected controls.
- the embedding fp16 -> bf16; the 48 layer norms byte-exactly from their .fconst offsets,
  each confirmed by the epsilon that opens its block (48/48 anchors).
- K is padded to the 256-column grid (hidden 896 = 3.5 tiles -> 4), recorded in the
  container manifest; nothing reads a padding column as a weight.

Refused by name, and a refusal writes nothing (exit 3, --allow-incomplete records the gap
in the container's __metadata__ and config.json instead):
- the 24 x 7 projections, whose weights exist only in dd_metastate_*.fconst and whose
  4-bit tile packing is not established (FACTS §3);
- the final norm: measured on this pack, the 1,856-byte tail of EITHER .fconst is not an
  RMSNorm weight (Token_Token's opens with a byte-pattern run, Llm_Prefill's is
  [eps][pad][448 f32 in 7.6..16.8]), so FACTS §2's "final-norm/eps at the tail" holds only
  for the epsilon.

On Qwen2.5-0.5B-Instruct_rai_1.8.0_npu_4K that is 50 tensors / 18992 tiles imported and
169 refusals, and the container read back through q4nx::dequant_tile matches the pack to
3.34e-3 relative / 2.4e-4 absolute -- the bf16 rounding of the zero point, which is exact
for 43.5 % of groups (the chunk's scale is bf16, and this pack's scales are already bf16).

tests/npu_pack_import_test: the decode, the tiling against tests/golden/oga_pack_tiles.bin
(written by tools/oga_pack_tiles_golden.py with the layout tools/gguf_to_q4nx.py has always
used, so the writer is checked against the established container and not only against its own
decoder), and, with -DONEBIT_OGA_PACK=<dir>, the refusal list, the self-check numbers and
that the result opens with npu::Model. Registered as npu_pack_import (0.00 s) and
npu_pack_import_oga (0.95 s).

docs/npu.md: "From a Ryzen AI OGA NPU pack" - the conversion table, the two bounds, the
refusal contract and the checks. Item 2 of #671 (per-shape ELFs) is device work and follows.
… nothing

The self-check tests the 4-bit decode against the pack's own fp16 embedding, which only means
something when the head is tied to it (tie_word_embeddings). Every 1.8.0 pack checked so far is
tied, but an untied one would have produced corr/rel numbers against an unrelated tensor -- a wrong
number is worse than none. The tie flag is now read from config.json: when it is false the checks
are skipped and the CLI prints why. Found while running a second family (Llama-3.2-1B, #671 item 4);
all three of its packs and the Qwen2.5 ones are tied, so nothing changes for them.
The pack's cache/*_meta.json is the JSON twin of the .meta protobuf and declares, per op, its attrs N/K
and every const tensor's dtype, shape and byte size. The plan now reads it and records declared_n /
declared_k per projection, plus a named mismatch when the pack and our config-derived plan disagree --
the pack stating its own geometry is a stronger check on our arithmetic than our arithmetic alone, and
it is the input engine#734 needs to generalise the .fconst offsets past hidden 896.

Result on both packs: Qwen2.5-0.5B 168/168 projections declared, Llama-3.2-1B 80/80 declared,
0 disagreements. Also visible now in the manifest: only 5 MladfMatMul per layer (q,k,v,o + down_proj),
because gate/up live inside the fused FlatMLP op -- whose const args are unpacked MatMulNBits
(uint8 [4864,7,64], float scales, uint8 qzeros, float bias), i.e. the same form the importer already
decodes for the head.
… any hidden size (#734)

The stride and the two norm offsets were constants measured on one pack (8,708,288 / +64 / +1,116,032 at
hidden 896), so every other geometry was refused with 'the .fconst offsets are established for hidden 896
only'. They are now measured from the file: every layer block opens with the bf16 epsilon and carries a
second such header after the attention payload, so the header positions repeat with one period per block
and the stride is positions[i+P] - positions[i] for that period (the most common difference, so a header
missed at a read-chunk boundary does not defeat it).

Measured and verified on three architectures, 176 anchors in total:

  pack                 hidden  layers  stride       second header  anchors
  Qwen2.5-0.5B-Instruct  896      24    8,708,288    1,115,968      48/48
  Qwen3-1.7B            2048      28   27,931,456    4,739,136      56/56
  Qwen3-4B              2560      36   55,308,096    8,778,816      72/72

The norms are now importable at any hidden size (50 / 58 / 74 tensors per pack), the geometry is recorded
in the container manifest, and the test is geometry-generic: it derives its expectations from the pack's
own config, with the hand-measured 0.5B offsets still asserted exactly when that geometry is present.

Note for #734: the metadata-only route was tried first and does NOT close it. Per-projection N/K are
declared (attrs.N/K) but the attention payload's tile count does not follow N/64 x ceil(K/kernel_K): that
predicts 96 tiles for the 1.7B (measured 136) and 160 for the 4B (measured 252), because the writer's
kernel shapes are not in the metadata. Measuring the anchors is what makes this work for every hidden
size today; a closed-form stride from metadata alone remains open, and the measured table above is the
evidence any such formula must reproduce.
…get in CI

Review (pi-prs on #780): bytes_of() compared `t.offset + t.length > ext_size`, which a crafted
model.onnx wraps -- offset = UINT64_MAX - 4 with length = 8 sums to 3, passes the check, and then
it->second + t.offset is a wild pointer read for 8 bytes. The bounds test is now
`ext_range_ok(size, offset, length)` = `offset <= size && length <= size - offset`, both comparisons
against a value already smaller than the file, and it has its own seven checks (the adversarial
pair, offset past the end with length 0, whole file, empty range at EOF, one byte past the end).

The red build was a second, separate fault and mine: CI names every target a ctest test invokes,
`npu_pack_import_test` was registered with add_test() but missing from that list, so test 46
reported Not Run and ctest exited 8. Added. (Checked the whole repository for the same mistake:
it was the only unconditional ctest executable missing from the list.)

Verified with CI's own configuration -- cmake Release, ninja, gcc-16, ONEBIT_OGA_PACK set --
in a fresh build directory: target builds, ctest -R npu_pack_import 2/2 pass (0.00 s and 1.20 s
against the real pack), tools/copyright.py --check clean, no warning from the new code.
`uint64_t need = 0` meant "no check", so a reader added later could opt out of the bound by omitting the
argument -- the same over-read this change exists to close, one line away. `need` is now a required parameter and
the comment says why: every call site must state what it reads.

No call site changes: every one already passed its computed length (embed V*H*width, scales N*(K/block)*width,
zeros N*(K/block+1)/2, qweight N*(K/block)*block/2 including the nibble-swap and no-zero controls).

Rebased onto main (which carries #787), so this branch also has the tests/golden/*/*.f32 exemption and
copyright.py --check is green here again.
@bong-water-water-bong

Copy link
Copy Markdown
Collaborator Author

Ready at ceb04af (pushed, force-with-lease on this branch: it is rebased onto current main, so the branch now carries #787's tests/golden/*/*.f32 exemption and copyright.py --check is green here again — your red check is gone with the rebase, nothing needed from you).

The one line: need is mandatory.

auto bytes_of = [&](const OnnxTensor& t, uint64_t need) -> const uint8_t* {

No call site changes — every reader already passed its computed length. The comment above it now says why the parameter has no default: a reader added later must not be able to opt out of the bound by omitting the argument.

The crafted-ONNX fixture is issue #790, as you suggested — with the run-time-fixture pattern (moe_gguf_index_test style), the three overstating cases (declared dims, MatMulNBits N/K/block_size vs each of qweight/scales/qzeros), and the acceptance line that it must fail if bytes_of is ever called with more than the tensor carries, without needing a real pack in CI.

Verified at this head with a fresh Release/ninja/gcc-16 configure and ONEBIT_OGA_PACK set:

$ python3 tools/copyright.py --check      rc 0
$ ninja npu_pack_import_test              rc 0
$ ctest -R npu_pack_import --output-on-failure
1/2 Test #48: npu_pack_import ....... Passed  0.00 sec
2/2 Test #49: npu_pack_import_oga ... Passed  1.03 sec      <- the real 0.5B pack
100% tests passed, 0 tests failed out of 2

@bong-water-water-bong bong-water-water-bong left a comment

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

VERDICT: READY

Both items closed, and the rebase did clear the red check exactly as you said it would.

The one line is done the right way (npu/pack_import.cpp:813):

auto bytes_of = [&](const OnnxTensor& t, uint64_t need) -> const uint8_t * {
    if (t.raw) {
        if (need > t.raw_len) throw std::runtime_error(t.name + ": needs " + std::to_string(need) + " B, carries " + …);

No default, no call-site churn, and the error names the tensor and both numbers, so a future reader that gets its need
wrong learns which tensor and by how much rather than reading past it. That was the whole point of the change.

The stale-base red is gone with the rebase, as predicted: python3 tools/copyright.py --check → rc 0 at
ceb04af9 (the branch now carries tests/golden/*/*.f32 from #787). Ran it again from a fresh build:

$ ninja npu_pack_import_test              rc 0
$ ctest -R npu_pack_import                2/2 passed (0.91 s for the real 0.5B pack)

And the fixture is tracked properly. Issue #790 is open with the run-time-fixture pattern named
(moe_gguf_index_test-style crafted headers), and the acceptance line is the one I would have insisted on: "ctest -R npu_pack_import gains a case that fails if bytes_of is ever called with a length larger than the tensor it reads, without
needing a real OGA pack on the box (CI has none)."
That is a regression test for the class, not a one-off assertion, and it
is the right place for it — so this PR lands with the arithmetic covered by the five mul_ok cases and the plumbing
protected by a tracked issue.

build was in_progress at post time (CodeQL success); everything I can run locally is green. tprs-ready applied.

@github-actions

github-actions Bot commented Nov 1, 2026

Copy link
Copy Markdown

Persistent review updated to latest commit 4b92047

@bong-water-water-bong bong-water-water-bong left a comment

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

VERDICT: READY

At 4b92047e (the waiter had already brought it up to date — the API said "no new commits on the base branch"), build green,
and its own paths are byte-identical to the head I verdicted (git diff --stat ceb04af9 <head> -- npu/pack_import.{cpp,h} tests/npu_pack_import_test.cpp tools/copyright.py → empty).

So both items stand closed: bytes_of(&t, uint64_t need) is mandatory with a throw that names the tensor and both numbers
(npu/pack_import.cpp:813), mul_ok is the overflow-proof multiply with five direct cases, every reader passes its computed
need, and copyright --check is rc 0 in this tree. The crafted-ONNX fixture is tracked as #790 with the acceptance line I
asked for. tprs-ready stays.

@bong-water-water-bong
bong-water-water-bong merged commit 921179c into main Nov 1, 2026
8 checks passed
@bong-water-water-bong
bong-water-water-bong deleted the 1bit/671-oga-pack-import branch November 1, 2026 22:50
@bong-water-water-bong

Copy link
Copy Markdown
Collaborator Author

#790 landed on this branch as 2c16b07, since the gate it pins lives here. Two crafted packs built at run time (a ~90-line protobuf writer for the fields walk_graph reads):

  • embed declared [1024,1024] × 4 B over a 1 KiB .data file → model.embed_tokens.weight: needs 4194304 B, declares 1024 in model.onnx.data
  • a MatMulNBits with scales/zeros sized exactly as its N/K/block_size demand, so the qweight read fires → lm_head_qweight: needs 16384 B, declares 256 in model.onnx.data

Both assert nothing is written, and both run without a real OGA pack. Two things the first draft got wrong and the fixtures now encode: the crafted pack needs a config.json and an empty cache/ (plan() reads them first), and it has to import with allow_incomplete or the plan's refusals short-circuit before the readers — otherwise the test passes for the wrong reason.

ctest -R npu_pack_import 2/2 at this head, copyright.py --check clean. No code changes since ceb04af other than the test.

bong-water-water-bong added a commit that referenced this pull request Nov 2, 2026
 #790) (#795)

#780 put one gate in front of every read whose length comes from the file's own metadata (bytes_of(t, need) with
mul_ok for the arithmetic). Its unit cases cover the arithmetic; this covers the wiring, which is what a future
edit could break without any unit test noticing.

Built at run time in the moe_gguf_index_test style (a small protobuf writer for just the fields the parser reads:
TensorProto dims/data_type/name/external_data/data_location, NodeProto op_type/inputs/attributes, GraphProto and
ModelProto):

  * dims that overstate the external file -- an embed declared [1024, 1024] x 4 B = 4 MiB over a 1 KiB .data
    file -> "model.embed_tokens.weight: needs 4194304 B, declares 1024 in model.onnx.data";
  * a MatMulNBits whose scales and zeros are sized exactly as its N/K/block_size demand, so the refusal that
    fires is the qweight one -> "lm_head_qweight: needs 16384 B, declares 256 in model.onnx.data".

Each case also asserts that no model.q4nx is written, and both run without a real OGA pack (CI has none). The
crafted packs carry a minimal config.json and an empty cache/ because plan() reads those before the readers, and
they import with allow_incomplete so the plan's own refusals do not short-circuit the path under test.

Co-authored-by: bong-water-water-bong <bong-water-water-bong@1bit.gg>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

enhancement New feature or request Review effort 4/5 tprs-ready TPRS: review verdict READY posted by the standing reviewer; merge gate satisfied

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant