Skip to content

npu-pack: the tied-head check tests for a scaled head (scale, rms, rel vs the 4-bit noise floor), not rel <= 0.02 (Refs #1503) #2194

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
bong-water-water-bong merged 2 commits into from
Nov 11, 2026
Merged

npu-pack: the tied-head check tests for a scaled head (scale, rms, rel vs the 4-bit noise floor), not rel <= 0.02 (Refs #1503) #2194

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
98531e9
npu-pack: the tied-head check tests for a scaled head (scale, rms, re…
Nov 11, 2026
1123461
npu_pack_import_test: fit_head's floor follows the head's block (bloc…
Nov 11, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 22 additions & 0 deletions docs/npu.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1467,6 +1467,28 @@ the decode, the tiling against `tests/golden/oga_pack_tiles.bin` — written by
the attention tile against `tests/golden/oga_attn_tile.bin`, and, against a real pack, the refusal
list, the tile round trip, the self-check numbers and that the container opens with `npu::Model`.

### The tied-head self-check

On a tied pack, the pack's own embedding is an independent reference for the decoded 4-bit head. The import compares the two and refuses a head that is scaled or only partly decoded. A pack in that state imports cleanly and answers garbage on the device, so `--allow-incomplete` does not override this refusal. The check has four parts. A head passes when all of them hold:

- **correlation** above 0.99;
- **best-fit scale** within 5 % of 1. This is the least-squares `a` in `embedding ≈ a · head`, so a scaled copy shows up here directly;
- **RMS ratio** `rms(head) / rms(embedding)` within 5 % of 1;
- **relative error** no more than 1.5× the embedding's own **4-bit noise floor**, or 0.02 if that is larger. The floor is what quantizing the embedding at the head's block size (asymmetric 4-bit, min/max per block) costs.

The relative error on its own cannot separate a scaled head from a correct 4-bit head whose embedding is kept at higher precision: both score about 0.09. Measured (engine#1503):

| pack | rel | best-fit scale | RMS ratio | noise floor | result |
|---|---|---|---|---|---|
| `Qwen2-1.5B_rai_1.9.0_npu` (block-32 tail head) | 0.094316 | 0.994173 | 1.001808 | 0.083865 | imports; on the NPU (`kernels-1536-2048-b8`, greedy) it answers "The capital of France is Paris." and "12 times 12 is 144." |
| `SmolLM2-135M` rai_1.7.1 hybrid | 0.103587 | 0.992172 | 1.003540 | 0.097113 | imports |
| `Qwen2.5-0.5B` rai_1.9.0 | 0.000700 | 0.999993 | 1.000006 | 0.006941 | imports |
| `gemma-2-2b-awq-…-hybrid` | 0.094874 | 0.992235 | 1.003849 | 0.088985 | head passes, but refused by `model_type` (below) |

A relative error near 0 means the embedding *is* the dequantized head (as on `Qwen-2.5_1.5B_Instruct` 1.9). It does not mean that only such packs are correct. The earlier rule (relative error ≤ 0.02, with nothing else checked) refused the first two packs above.

**`model_type: gemma2` is refused by name.** The two gemma-2-2b hybrid packs answer garbage on this route (`myſelf…`), but their head is a correct 4-bit head, so the cause is somewhere else and has not been found. The import refuses them with that reason, rather than letting them import cleanly.

### Cross-check against the official converter

ROCm's **FLM_Q4NX_Converter** (`convert.py`) is the official gguf → Q4NX converter, and since
Expand Down
115 changes: 75 additions & 40 deletions npu/pack_import.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4520,6 +4520,7 @@ Report import_pack(const std::string& pack_dir, const std::string& out_dir, cons
"tile decoder does not apply (engine#1503)");
Report r;
r.out_dir = out_dir;
r.model_type = p.dims.model_type;
r.refused = p.refused();
if (!r.refused.empty() && !opt.allow_incomplete) return r; // refuse by name, write nothing

Expand Down Expand Up @@ -5030,24 +5031,15 @@ Report import_pack(const std::string& pack_dir, const std::string& out_dir, cons
mm.weight.c_str(), q.N, q.K, mm.block, int(V), int(H), (int)head_geometry);
if (tied || head_geometry) r.selfcheck_reference = true;
if (tied || head_geometry) {
double sum_abs = 0, sum_ref = 0;
for (int n = 0; n < q.N; ++n)
for (int k = 0; k < q.K; ++k) {
const size_t g = size_t(n) * (q.K / 32) + size_t(k) / 32;
const double x = double(q.codes[size_t(n) * q.K + k]) * q.scale[g] + q.zero[g];
double e;
e = embed_at(embed_tbl, int(n), k);
sum_abs += std::fabs(x - e);
sum_ref += std::fabs(e);
}
r.lmhead_rel = sum_abs / sum_ref;
r.head_fit = fit_head(q, embed_tbl, q.N, mm.block);
r.lmhead_rel = r.head_fit.rel;
r.lmhead_corr = corr(q, embed_tbl, q.N);
// The diagnostic follows the NUMBER, not the pack's declaration: a genre-A pack ships no config.json, so
// `tied` is false there while the head plainly reproduces its embedding, and the readout used to call those
// packs untied while the container it wrote said tied. Same threshold the container's own decision uses.
r.selfcheck_tied = r.lmhead_corr > kTieCorr;

r.selfcheck_scaled = r.selfcheck_tied && !selfcheck_head_ok(r.lmhead_corr, r.lmhead_rel);
r.selfcheck_scaled = r.selfcheck_tied && !selfcheck_head_ok(r.lmhead_corr, r.head_fit);
const Q4 swap = decode_matmulnbits(bytes_of(wt, qw_need), scales.data(), zp, int(mm.N), int(mm.K),
mm.block, /*swap_nibbles*/ true, false);
const Q4 nozp = decode_matmulnbits(bytes_of(wt, qw_need), scales.data(), nullptr, int(mm.N), int(mm.K),
Expand Down Expand Up @@ -5121,22 +5113,12 @@ Report import_pack(const std::string& pack_dir, const std::string& out_dir, cons
: " (untied: layout checked, values not verified)");
if (embed_tbl.ok()) {
r.lmhead_corr = c32;
{
double sum_abs = 0, sum_ref = 0;
for (int n = 0; n < int(V); ++n)
for (int k = 0; k < q.K; ++k) {
const size_t gg = size_t(n) * (q.K / 32) + size_t(k) / 32;
const double x = double(q.codes[size_t(n) * q.K + k]) * q.scale[gg] + q.zero[gg];
const double e = embed_at(embed_tbl, n, k);
sum_abs += std::fabs(x - e);
sum_ref += std::fabs(e);
}
if (sum_ref > 0) r.lmhead_rel = sum_abs / sum_ref;
}
r.head_fit = fit_head(q, embed_tbl, int(V), 32);
r.lmhead_rel = r.head_fit.rel;
r.selfcheck_reference = true;
r.selfcheck_tied = c32 > kTieCorr;

r.selfcheck_scaled = r.selfcheck_tied && !selfcheck_head_ok(r.lmhead_corr, r.lmhead_rel);
r.selfcheck_scaled = r.selfcheck_tied && !selfcheck_head_ok(r.lmhead_corr, r.head_fit);
double sxy = 0, sx = 0, sy = 0, sxx = 0, syy = 0;
const int nsh = std::min(2048, int(V) - 1);
for (int rr = 0; rr < nsh; ++rr)
Expand Down Expand Up @@ -5450,22 +5432,13 @@ Report import_pack(const std::string& pack_dir, const std::string& out_dir, cons
// q.N may exceed V by the pack's tile padding (Hunyuan-0.5B: 120,832 rows for 120,818),
// so the reference is compared over the vocabulary's own rows.
if (q.N >= int(V) && q.K == int(H)) {
double sum_abs = 0, sum_ref = 0;
for (int n = 0; n < int(V); ++n)
for (int k = 0; k < q.K; ++k) {
const size_t g = size_t(n) * (q.K / 32) + size_t(k) / 32;
const double x = double(q.codes[size_t(n) * q.K + k]) * q.scale[g] + q.zero[g];
const double e =
embed_at(embed_tbl, int(n), k);
sum_abs += std::fabs(x - e);
sum_ref += std::fabs(e);
}
r.lmhead_rel = sum_abs / sum_ref;
r.head_fit = fit_head(q, embed_tbl, int(V), q.block > 0 ? q.block : 128);
r.lmhead_rel = r.head_fit.rel;
r.lmhead_corr = corr(q, embed_tbl, q.N);
r.selfcheck_reference = true;
r.selfcheck_tied = r.lmhead_corr > kTieCorr;

r.selfcheck_scaled = r.selfcheck_tied && !selfcheck_head_ok(r.lmhead_corr, r.lmhead_rel);
r.selfcheck_scaled = r.selfcheck_tied && !selfcheck_head_ok(r.lmhead_corr, r.head_fit);
// The control this path can actually run: the same tiles and the same decoder, with each
// head row matched against the NEXT embedding row. If a misaligned head still correlated,
// the number above would be measuring the vocabulary's mean brightness rather than the
Expand Down Expand Up @@ -6067,14 +6040,26 @@ int run(int argc, char** argv) {
// packs measured: rel_err 0.095 and 0.136 against 0.000000 on the pack that serves). `--allow-incomplete` is not an
// override here: that flag means "write the tensors I could not represent", and this is "the tensors I wrote do not
// decode", which no amount of writing fixes.
// engine#1503: an architecture this route is measured to answer garbage on, whatever its head scores. The head fit
// above stopped refusing gemma-2-2b AWQ (its head is a correct 4-bit head), so the refusal that #2101 made for the
// wrong reason is kept here for the right one, by name, until the cause is found.
if (const char* why = known_bad_architecture(r.model_type)) {
std::fprintf(stderr, "1bit npu-pack: refusing to import %s: model_type %s -- %s. The partial output this run "
"already wrote under %s is not servable -- delete it.\n",
dir.c_str(), r.model_type.c_str(), why, out.c_str());
return 3;
}
if (r.selfcheck_scaled) {
std::fprintf(stderr,
"1bit npu-pack: refusing to import %s: the head correlates with the pack's own embedding (%.6f) but "
"its relative error is %.6f, which is a scaled or partially decoded head rather than a rounding "
"difference (the threshold is %.3f). A pack in this state imports cleanly and answers garbage on "
"it is a scaled or partially decoded head rather than a correct 4-bit one: relative error %.6f "
"(allowed %.6f = max(%.3f, %.1f x the embedding's own 4-bit noise floor %.6f)), best-fit scale %.6f "
"and rms ratio %.6f (allowed 1 +- %.2f). A pack in this state imports cleanly and answers garbage on "
"the device; --allow-incomplete does not override this. The partial output this run already\n"
"wrote under %s is not servable -- delete it before retrying.\n",
dir.c_str(), r.lmhead_corr, r.lmhead_rel, kSelfcheckRelMax, out.c_str());
dir.c_str(), r.lmhead_corr, r.lmhead_rel,
std::max(kSelfcheckRelMax, kHeadFloorFactor * r.head_fit.floor), kSelfcheckRelMax, kHeadFloorFactor,
r.head_fit.floor, r.head_fit.scale, r.head_fit.rms_ratio, kHeadScaleTol, out.c_str());
return 3;
}
if (!r.refused.empty() && !allow) {
Expand All @@ -6091,6 +6076,10 @@ int run(int argc, char** argv) {
std::printf(" self-check: the head reproduces the pack's own embedding (as imported, bf16) corr %.6f rel_err %.6f\n"
" (tied: this correlation is the check, and it passed)\n",
r.lmhead_corr, r.lmhead_rel);
if (r.head_fit.floor > 0)
std::printf(" head fit: best-fit scale %.6f, rms ratio %.6f, embedding's own 4-bit noise floor %.6f "
"(rel_err is %.2fx of it): a 4-bit head, not a scaled one\n",
r.head_fit.scale, r.head_fit.rms_ratio, r.head_fit.floor, r.lmhead_rel / r.head_fit.floor);
if (r.ctrl_corr_rowshift >= 0)
std::printf(" controls: row-shifted corr %.4f (must fail)\n", r.ctrl_corr_rowshift);
else
Expand Down Expand Up @@ -6335,4 +6324,50 @@ void sweep_fconst_scratch_dirs(const std::vector<std::string>& dirs, const std::
// box measures are rel_err 0.000000 (serves) and 0.095/0.136 (imports cleanly, answers garbage).
bool selfcheck_head_ok(double corr, double rel) { return corr > kTieCorr && rel <= kSelfcheckRelMax; }

const char* known_bad_architecture(const std::string& model_type) {
if (model_type == "gemma2")
return "gemma-2-2b hybrid packs import cleanly and answer garbage on this route (`myselfGetBytes...`, two packs, "
"measured for #2101), and their head is a correct 4-bit head (best-fit scale 0.992, rel 1.07x its 4-bit "
"noise floor), so the fault is elsewhere and not yet found (engine#1503)";
return nullptr;
}

bool selfcheck_head_ok(double corr, const HeadFit& f) {
return corr > kTieCorr && std::fabs(f.scale - 1.0) <= kHeadScaleTol && std::fabs(f.rms_ratio - 1.0) <= kHeadScaleTol &&
f.rel <= std::max(kSelfcheckRelMax, kHeadFloorFactor * f.floor);
}

HeadFit fit_head(const Q4& q, const EmbedTable& embed, int rows, int block) {
HeadFit f;
rows = std::min(rows, q.N);
if (block <= 0 || q.K % block) block = 32;
double sum_abs = 0, sum_ref = 0, sxe = 0, sxx = 0, see = 0, floor_abs = 0;
std::vector<double> e(size_t(q.K));
for (int n = 0; n < rows; ++n) {
for (int k = 0; k < q.K; ++k) e[size_t(k)] = embed_at(embed, n, k);
for (int k = 0; k < q.K; ++k) {
const size_t g = size_t(n) * (q.K / 32) + size_t(k) / 32;
const double x = double(q.codes[size_t(n) * q.K + k]) * q.scale[g] + q.zero[g];
sum_abs += std::fabs(x - e[size_t(k)]);
sum_ref += std::fabs(e[size_t(k)]);
sxe += x * e[size_t(k)];
sxx += x * x;
see += e[size_t(k)] * e[size_t(k)];
}
for (int g0 = 0; g0 < q.K; g0 += block) {
double lo = e[size_t(g0)], hi = lo;
for (int k = g0; k < g0 + block; ++k) { lo = std::min(lo, e[size_t(k)]); hi = std::max(hi, e[size_t(k)]); }
const double sc = (hi - lo) / 15.0;
for (int k = g0; k < g0 + block; ++k) {
const double c = sc > 0 ? std::min(15.0, std::max(0.0, std::round((e[size_t(k)] - lo) / sc))) : 0.0;
floor_abs += std::fabs(c * sc + lo - e[size_t(k)]);
}
}
}
if (sum_ref > 0) { f.rel = sum_abs / sum_ref; f.floor = floor_abs / sum_ref; }
if (sxx > 0) f.scale = sxe / sxx;
if (see > 0) f.rms_ratio = std::sqrt(sxx / see);
return f;
}

} // namespace onebit::npu::pack_import
26 changes: 25 additions & 1 deletion npu/pack_import.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -434,6 +434,10 @@ std::string weight_layout_refusal(const std::string& want);
// broken container, so it is refused here, by name, before anything is written.
std::string kernel_set_refusal(const std::vector<std::string>& file_names, const std::string& layout);

// The tied-head fit (engine#1503); what each field means is stated at fit_head() below.
struct HeadFit {
double rel = 0, scale = 1, rms_ratio = 1, floor = 0;
};
struct Report {
std::string out_dir;
int tensors = 0;
Expand All @@ -456,6 +460,8 @@ struct Report {
// cleanly and answers nonsense, so it is refused rather than reported.
bool selfcheck_scaled = false;
double lmhead_corr = 0, lmhead_rel = 0, ctrl_corr_nibble = 0, ctrl_corr_nozp = 0;
HeadFit head_fit; // engine#1503: scale, rms ratio and the 4-bit noise floor beside lmhead_rel
std::string model_type; // the pack's, so the CLI can refuse a known-bad architecture by name
// The tail (.fconst) path has no MatMulNBits bytes to mis-decode, so its control is an ALIGNMENT one: the same
// tiles, the same decoder, rows matched against the neighbouring embedding row. A correct decode correlates ~1
// with the embedding; if a row-misaligned head still did, the check would be measuring nothing. -1 = not run.
Expand Down Expand Up @@ -794,7 +800,15 @@ const Fconst19Rule* select_fconst19_rule_from(const Fconst19Rule* rules, int cou
// embedding (corr) AND not be a scaled copy of it (rel_err). The constant is the measured gap between the two states --
// 0.000000 for the pack that serves, 0.095/0.136 for the two that do not -- so 0.02 sits far from both.
inline constexpr double kSelfcheckRelMax = 0.02;
bool selfcheck_head_ok(double corr, double rel);
bool selfcheck_head_ok(double corr, double rel); // the same rule with no fit: scale 1, no noise floor
// engine#1503: the rule on what it claims. A head passes when it correlates (kTieCorr), is NOT scaled (scale and
// rms_ratio within kHeadScaleTol of 1), and its rel error is no worse than kHeadFloorFactor x the embedding's own 4-bit
// noise floor at the head's block (and never refused below kSelfcheckRelMax, so a head that IS its embedding passes).
inline constexpr double kHeadScaleTol = 0.05;
inline constexpr double kHeadFloorFactor = 1.5;
bool selfcheck_head_ok(double corr, const HeadFit& fit);
// A model_type this route is MEASURED to serve garbage on, with the reason; nullptr for every other (engine#1503).
const char* known_bad_architecture(const std::string& model_type);

int64_t attn_tile_bytes(int64_t k);
// Pearson correlation between a quantised head and a bf16 embedding table (engine#1690; the self-check's own statistic,
Expand All @@ -815,6 +829,16 @@ struct EmbedTable {
// The only decode: bf16 or f16 per t.fmt. out-of-range rows still read (callers clamp), as before.
double embed_at(const EmbedTable& t, int row, int k);

// engine#1503: what the tied-head self-check measures, beyond the correlation. `rel` alone cannot tell a scaled head
// from a correct 4-bit head whose embedding is kept at higher precision: both score ~0.09 (measured: Qwen2-1.5B 1.9,
// tail head block 32, rel 0.094316; gemma-2-2b AWQ hybrid, rel 0.094874), while the pack 0.02 was calibrated on
// scores 0 only because its embedding IS the dequantised head. So the fit also carries:
// scale -- the least-squares a in e ~= a * head. A scaled copy shows it directly; the two above are 0.994 / 0.992.
// rms_ratio -- rms(head) / rms(embedding), the same question without the fit (1.0018 / 1.0038).
// floor -- the rel error of the EMBEDDING quantised the way the head is (asymmetric 4-bit, min/max per `block`
// columns): what a correct head costs. The two above are 1.12x / 1.07x of theirs (0.083865 / 0.088985).
HeadFit fit_head(const Q4& q, const EmbedTable& embed, int rows, int block);

double corr_bf16(const Q4& q, const uint8_t* embed_bf16, int H, int rows);
// Where a layer block's attention payload starts: the block's 64-byte header, then its
// input_layernorm weight as bf16 [hidden]. So it is `64 + 2 x hidden`, not a constant. It reads 1,856 on
Expand Down
Loading
Loading