diff --git a/docs/npu.md b/docs/npu.md index baedf4d8e..823fe351e 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -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 diff --git a/npu/pack_import.cpp b/npu/pack_import.cpp index 0f8071d5f..99cd91bd1 100644 --- a/npu/pack_import.cpp +++ b/npu/pack_import.cpp @@ -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 @@ -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), @@ -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) @@ -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 @@ -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) { @@ -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 @@ -6335,4 +6324,50 @@ void sweep_fconst_scratch_dirs(const std::vector& 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 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 diff --git a/npu/pack_import.h b/npu/pack_import.h index 16c1af0b4..1f2fc16f2 100644 --- a/npu/pack_import.h +++ b/npu/pack_import.h @@ -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& 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; @@ -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. @@ -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, @@ -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 diff --git a/tests/npu_pack_import_test.cpp b/tests/npu_pack_import_test.cpp index ae29dc890..fbdd56e55 100644 --- a/tests/npu_pack_import_test.cpp +++ b/tests/npu_pack_import_test.cpp @@ -5744,6 +5744,102 @@ static void selfcheck_head_rule_arms() { check(!pi::selfcheck_head_ok(0.50, 0.000000), "head rule: an uncorrelated head is still refused on correlation"); check(pi::selfcheck_head_ok(0.996053, pi::kSelfcheckRelMax), "head rule: exactly at the threshold passes (inclusive bound)"); std::printf(" head rule: 6 arms (two measured bad packs, the serving one, the unscaled control, no-corr, boundary)\n"); + + // engine#1503: the rule on what it claims (scale, rms ratio, rel against the embedding's own 4-bit noise floor). + // Measured fits, from the import's own numbers: + const pi::HeadFit qwen2_15b{0.094316, 0.994173, 1.001808, 0.083865}; // Qwen2-1.5B 1.9, block-32 tail head + const pi::HeadFit gemma_awq{0.094874, 0.992235, 1.003849, 0.088985}; // gemma-2-2b AWQ hybrid, graph head + check(!pi::selfcheck_head_ok(0.995969, qwen2_15b.rel), + "fit rule: the rel-only rule refuses Qwen2-1.5B's head (the reason this rule exists)"); + check(pi::selfcheck_head_ok(0.995969, qwen2_15b), + "fit rule: Qwen2-1.5B's head passes: scale .994, rms 1.002, rel 1.12x its 4-bit floor"); + check(pi::selfcheck_head_ok(0.996053, gemma_awq), + "fit rule: gemma-2-2b AWQ's head passes too -- by these numbers its head is not what served garbage"); + check(pi::selfcheck_head_ok(0.996901, pi::HeadFit{0.0, 1.0, 1.0, 0.08}), + "fit rule: a head that IS its embedding (rel 0) passes"); + check(!pi::selfcheck_head_ok(0.996, pi::HeadFit{0.09, 0.5, 2.0, 0.08}), "fit rule: a head scaled by 2 is refused"); + check(!pi::selfcheck_head_ok(0.996, pi::HeadFit{0.09, 0.94, 1.0, 0.08}), "fit rule: scale 0.94 (outside 5 %) is refused"); + check(!pi::selfcheck_head_ok(0.996, pi::HeadFit{0.09, 1.0, 1.06, 0.08}), "fit rule: rms ratio 1.06 is refused"); + check(!pi::selfcheck_head_ok(0.996, pi::HeadFit{0.13, 1.0, 1.0, 0.08}), + "fit rule: rel 1.6x the floor is refused (a partial decode is noisier than quantisation)"); + check(pi::selfcheck_head_ok(0.996, pi::HeadFit{0.12, 1.0, 1.0, 0.08}), "fit rule: exactly 1.5x the floor passes"); + check(!pi::selfcheck_head_ok(0.50, qwen2_15b), "fit rule: an uncorrelated head is still refused on correlation"); + + // fit_head itself, on a synthetic table: the head is the embedding quantised asymmetric 4-bit at block 32, the way + // a correct pack's head is, so the fit must read scale ~1 and rel == floor; the same head scaled by 0.5 must not. + const int N = 64, K = 128; + std::vector emb(size_t(N) * K * 2); + std::vector ef(size_t(N) * K); + uint32_t st = 12345; + for (size_t i = 0; i < ef.size(); ++i) { + st = st * 1664525u + 1013904223u; + float v = (float(st >> 8) / float(1u << 24) - 0.5f) * 0.2f; + uint32_t b; + std::memcpy(&b, &v, 4); + const uint16_t h = uint16_t(b >> 16); // bf16: the top half, so ef[] holds exactly what the table says + b = uint32_t(h) << 16; + std::memcpy(&v, &b, 4); + ef[i] = v; + emb[2 * i] = uint8_t(h & 0xff); + emb[2 * i + 1] = uint8_t(h >> 8); + } + pi::EmbedTable et; + et.p = emb.data(); + et.H = K; + pi::Q4 q; + q.N = N; q.K = K; q.block = 32; + q.codes.assign(size_t(N) * K, 0); + q.scale.assign(size_t(N) * (K / 32), 0.f); + q.zero.assign(size_t(N) * (K / 32), 0.f); + for (int n = 0; n < N; ++n) + for (int g = 0; g < K / 32; ++g) { + float lo = 1e30f, hi = -1e30f; + for (int k = 0; k < 32; ++k) { lo = std::min(lo, ef[size_t(n) * K + g * 32 + k]); hi = std::max(hi, ef[size_t(n) * K + g * 32 + k]); } + const float sc = (hi - lo) / 15.f; + q.scale[size_t(n) * (K / 32) + g] = sc; + q.zero[size_t(n) * (K / 32) + g] = lo; + for (int k = 0; k < 32; ++k) + q.codes[size_t(n) * K + g * 32 + k] = + uint8_t(std::min(15.0, std::max(0.0, std::round((double(ef[size_t(n) * K + g * 32 + k]) - lo) / sc)))); + } + const pi::HeadFit good = pi::fit_head(q, et, N, 32); + check(std::fabs(good.scale - 1) < 0.01 && std::fabs(good.rms_ratio - 1) < 0.01 && good.rel > 0 && + std::fabs(good.rel - good.floor) < 1e-6 && pi::selfcheck_head_ok(0.999, good), + "fit_head: a correct 4-bit head reads scale ~1, rel == its floor, and passes (scale " + + std::to_string(good.scale) + " rel " + std::to_string(good.rel) + " floor " + std::to_string(good.floor) + ")"); + // The floor follows the head's block (the block-128 tail head passes 128): the same table quantised at block 128 + // must read rel == its block-128 floor, which is LARGER than the block-32 floor of the same table. + pi::Q4 q128 = q; + for (int n = 0; n < N; ++n) + for (int g = 0; g < K / 128; ++g) { + float lo = 1e30f, hi = -1e30f; + for (int k = 0; k < 128; ++k) { lo = std::min(lo, ef[size_t(n) * K + g * 128 + k]); hi = std::max(hi, ef[size_t(n) * K + g * 128 + k]); } + const float sc = (hi - lo) / 15.f; + for (int j = 0; j < 4; ++j) { // a 128-block's scale and zero apply to its four 32-column groups + q128.scale[size_t(n) * (K / 32) + g * 4 + j] = sc; + q128.zero[size_t(n) * (K / 32) + g * 4 + j] = lo; + } + for (int k = 0; k < 128; ++k) + q128.codes[size_t(n) * K + g * 128 + k] = + uint8_t(std::min(15.0, std::max(0.0, std::round((double(ef[size_t(n) * K + g * 128 + k]) - lo) / sc)))); + } + const pi::HeadFit f128 = pi::fit_head(q128, et, N, 128); + check(std::fabs(f128.rel - f128.floor) < 1e-6 && f128.floor > good.floor && pi::selfcheck_head_ok(0.999, f128), + "fit_head: at block 128 the floor is the block-128 one (rel " + std::to_string(f128.rel) + " == floor " + + std::to_string(f128.floor) + " > block-32 floor " + std::to_string(good.floor) + ")"); + pi::Q4 half = q; + for (auto& v : half.scale) v *= 0.5f; + for (auto& v : half.zero) v *= 0.5f; + const pi::HeadFit bad = pi::fit_head(half, et, N, 32); + check(std::fabs(bad.scale - 2) < 0.05 && !pi::selfcheck_head_ok(0.999, bad), + "fit_head: the same head at half scale reads scale ~2 and is refused (scale " + std::to_string(bad.scale) + ")"); + check(pi::known_bad_architecture("gemma2") != nullptr && + std::string(pi::known_bad_architecture("gemma2")).find("not yet found") != std::string::npos, + "known-bad: gemma2 stays refused, by name and with the reason (the head fit no longer refuses it)"); + check(pi::known_bad_architecture("qwen2") == nullptr && pi::known_bad_architecture("llama") == nullptr && + pi::known_bad_architecture("") == nullptr, + "known-bad: nothing else is on the list"); + std::printf(" head fit rule: 15 arms (two measured heads, scaled/rms/noisy/no-corr controls, fit_head at blocks 32 and 128, gemma2 by name)\n"); } int main(int argc, char** argv) {