[bong-water-water-bong]

Host-side pieces the MTP window needs, stacked on #797 (step_batch_all) and #798 (deferred commit + tau). No device
in any of it, and the head is checked against the reference before any window opens.

• npu/mtp/mtp_head.{h,cpp} — the head on the host in the pb16i-exact arithmetic the container holds: per 32-K
  group a signed int4 code and an int16 magnitude, one exponent per column, and the source's additive min as a host
  table (y = Σ_g (Σ_k a_k·code)·ms·2^eref + Σ_g A_g·add_g). Same decomposition the device does, so a draft measured
  here is a draft the NPU head can reproduce. Its vocabulary projection goes through the base model's lm head — the
  drafter has none of its own.
• npu/mtp/forward_base.h — BaseModel over Q4nxNpuForward: step, step_batch_all, last_hidden /
  last_residual, and commit_verify() as snapshot-then-replay while the engine has no GDN deferral (correct, not
  free; the deferred design is npu: the draft loop commits a verify prefix instead of rolling back, and gates the chain on tau #798's target).
• engine accessors (read-only, three small ones): lm_head_logits(), embed_row(), last_residual().
• tools/mtp_acceptance.cpp — the measurement itself: N prompts × M tokens, draft k, verify through
  step_batch_all, commit the accepted prefix, report acceptance length / hit rate / passes per token, both
  hidden-state variants (--residual), plus --selftest-head which needs no device.
• tools/mtp_head_gemv_harness.py — freezes the per-matrix fixture for the six head matrices that map onto sealed
  M=16 ELFs (int8 A, the exact int32 C the device must produce, and the float oracle for the same rows), so a device
  ticket can report cost and byte-exactness rather than just "it ran".

Checked host-side: the C++ head against the Python reference on a frozen input — rel-L1 1.838e-04, max|d|
1.567e-02, cosine 0.999999985, i.e. the int16 mantissa, the same limit the image itself has. mtp_draft ctest still
green (69 assertions) with pos threaded into DraftHead::draft (rope cannot guess the position of the first draft).

Not merged here.

[bong-water-water-bong]

VERDICT: CHANGES

Two items, both small; the pipeline itself builds and the draft loop is still green.

What I ran (-DONEBIT_NPU=ON -DONEBIT_NPU_RT_DIR=<private runtime>, so the NPU lane is actually compiled — CI does not):

$ ninja mtp_draft_test mtp_batch_all_test mtp_acceptance        rc 0
$ ./mtp_draft_test …                                            all checks passed   (69 assertions, incl. pos threading:
                                                                "commits=12 last_commit=1 accepted=12 drafted=24")

Item 1 — the head-vs-reference number cannot be reproduced from the tree. tools/mtp_acceptance.cpp:162 reads
hidden.bin, emb.bin and out_hidden.bin from --gold DIR, the tree ships none of them, and nothing in the tree writes
them — so --selftest-head, which is the only check of the C++ head's arithmetic, fails with cannot read the golden files under for anyone who is not the author:

$ ./mtp_acceptance --selftest-head --head-exact …/pb16i-exact --head-q4k …/q4_k
cannot read the golden files under

The claim in the body (rel-L1 1.838e-04, max|d| 1.567e-02, cosine 0.999999985) is self-consistent — it lands where the image's
own int16-mantissa limit does (#791 measured 1.56e-04 on the image itself), which is exactly why I want it runnable rather than
taken on trust. Ask: ship the three files (3 × 5,120 floats ≈ 60 KB, and the repo already ships binary goldens under
tests/golden/), or a ~15-line generator built on tools/mtp_head_ref.py head_step so the fixture is derivable from the
reference that defines it.

Item 2 — the branch is behind #807/#809/#810, and the rebase touches the same file as the plan layer. Its own change is
additive: three read-only accessors (lm_head_logits(), embed_row(), last_residual()), one last_residual_.assign(...) in
the step path, and the new npu/mtp/* + tools. What looks like deletion in npu/forward/include/q4nx_forward.h against current
main is #807's plan_/plan_path_/missing_kernel(), which this branch predates. So: rebase needed, and re-verdict after
it, because the accessors will land next to the plan members in one file.

Nothing else is in question: the arithmetic decomposition is the same one #791's artifact and #809's chunk use (signed int4
codes × int16 magnitudes, one exponent per column, the additive min as a host table), forward_base.h's
commit_verify() as snapshot-then-replay is correctly described as "correct, not free" until #798's deferred design is wired,
and the vocabulary projection going through the base model's own lm head is the right call (the drafter has no head of its own —
lm_head_logits() exists for exactly that).

[bong-water-water-bong]

VERDICT: READY

Correction to my own review: item 1 is closed, and I closed it myself rather than waiting. The head-vs-reference number is
not only reproducible, it reproduces — I generated the missing fixture from the Python reference (the same convention
check_head() already uses) and ran the tool's selftest against it:

# tools/ → the three files mtp_acceptance --selftest-head reads
W, _ = mtp_head_ref.load_q4k_head('…/bonsai2-27b-mtp-head/q4_k')
h, e = np.random.default_rng(0).standard_normal(H), np.random.default_rng(0).standard_normal(H)   # unit-RMS, as check_head does
_, out, _ = mtp_head_ref.head_step(h, e, W, pos=3)          # pos=3: the position the C++ selftest uses
h.astype('<f4').tofile('hidden.bin'); e.astype('<f4').tofile('emb.bin'); out.astype('<f4').tofile('out_hidden.bin')

$ ./mtp_acceptance --selftest-head --gold <that dir> --head-exact …/pb16i-exact --head-q4k …/q4_k
head selftest vs Python reference: rel-L1 1.933e-04, max|d| 1.758e-03, cosine 0.999999981      rc 0

Same order as the body's 1.838e-04 / 0.999999985, well inside the tool's own thresholds (rel-L1 < 5e-3, cosine > 0.9999), and
max|d| differs only because my fixture is a different draw than the author's. So the C++ head really is the reference's
arithmetic at the int16-mantissa limit — the claim stands, checked by me on this box.

That turns my item 1 from an objection into a reproducibility nit, which I am not holding the PR for: the recipe above (or
the three 20 KB files) belongs in the tree, because the next person who wants to check the head should not have to rediscover
the convention. One sentence or one script.

Item 2 stands as a process note, not a defect: the branch predates #807/#809/#810, its own change is additive, and the
apparent deletions in q4nx_forward.h are #807's plan members. When pi-bonsai rebases it I will re-post at the new head (one
line, own-path diff) — the label on this head means "content READY", not "merge without the rebase".

Everything else from my first review holds: NPU lane compiles (rc 0, private runtime, which CI does not do) and mtp_draft_test
is green at 69 assertions with pos threaded through DraftHead::draft.

[github-actions]

PR Reviewer Guide 🔍

(Review updated until commit e32370d)

Here are some key observations to aid the review process:

🎫 Ticket compliance analysis 🔶 797 - Partially compliant Compliant requirements: Implemented step_batch_all() function that returns per-row logits from one batched pass Function runs the same batched body once, norms every row, and runs LM head once for the batch Uses pb16i head path with per-row activation scales for M rows Host-side test (tests/mtp_batch_all_test.cpp) included and passes Non-compliant requirements: The test evidence for pb16i per-row head path was validated on device (mentioned as pre-existing issue) Requires further human verification: The segfault in host-GEMM for pb16i container (mentioned as pre-existing) needs further verification 798 - Partially compliant Compliant requirements: Implemented deferred commit instead of rollback using recurrence inputs storage Implemented tau gate on the chain with optional confidence parameter Validated with toy base (69 assertions, all green) Non-compliant requirements: The deferred design is targeted for npu: the draft loop commits a verify prefix instead of rolling back, and gates the chain on tau #798, but the current implementation uses snapshot-then-replay approach Requires further human verification: The deferred design implementation is not yet complete (deferred design is npu: the draft loop commits a verify prefix instead of rolling back, and gates the chain on tau #798's target)

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

🧪 PR contains tests

🔒 No security concerns identified

⚡ Recommended focus areas for review Potential Memory Leak in GDN Recurrence Storage The gdn_apply_rows function in Q4nxNpuForward allocates memory for row_in_ in begin_verify but does not ensure proper cleanup in all code paths. If gdn_apply_rows is called with invalid parameters or fails, the allocated memory might not be properly freed, leading to a potential memory leak. The row_in_ vector is cleared in commit_verify_rows but not in error cases within gdn_apply_rows. bool Q4nxNpuForward::gdn_apply_rows(int from, int to) { const int KD = lin_kd_, CD = 2 * lin_kd_ + lin_vd_; const int NKH = lin_nkh_, NVH = lin_nvh_, KHD = lin_khd_, VHD = lin_vhd_; const int CK = lin_ck_; const int rep = (NKH > 0) ? std::max(1, NVH / NKH) : 1; std::vector<float> co((size_t)CD), ql((size_t)NVH * KHD), kl((size_t)NVH * KHD), vv((size_t)NVH * VHD); std::vector<float> kvm((size_t)VHD); for (int l = 0; l < (int)row_in_.size(); l++) { if (row_in_[(size_t)l].empty() || rec_state_[(size_t)l].empty()) continue; const int lt = (l < sem_.n_layer_types) ? sem_.layer_types[l] : Q4NX_LAYER_FULL; if (!(lin_kd_ > 0 && lt == Q4NX_LAYER_LINEAR)) continue; // only GDN layers carry this state for (int m = from; m < to; m++) { const GdnRowIn& rec = row_in_[(size_t)l][(size_t)m]; if (rec.qkv.empty()) continue; float* st = conv_state_[(size_t)l].data(); for (int ch = 0; ch < CD; ch++) { const float* cs = st + (size_t)ch * (CK - 1); float ssum = 0; for (int j = 0; j < CK - 1; j++) ssum += cs[j] * gdn_small_[(size_t)l].conv[(size_t)j * CD + ch]; ssum += rec.qkv[(size_t)ch] * gdn_small_[(size_t)l].conv[(size_t)(CK - 1) * CD + ch]; co[(size_t)ch] = ssum / (1.0f + expf(-ssum)); } for (int ch = 0; ch < CD; ch++) { float* cs = st + (size_t)ch * (CK - 1); for (int j = 0; j + 1 < CK - 1; j++) cs[j] = cs[j + 1]; if (CK > 1) cs[CK - 2] = rec.qkv[(size_t)ch]; } for (int hh = 0; hh < NVH; hh++) { const int src = hh / rep; for (int d = 0; d < KHD; d++) ql[(size_t)hh * KHD + d] = co[(size_t)src * KHD + d]; for (int d = 0; d < KHD; d++) kl[(size_t)hh * KHD + d] = co[KD + (size_t)src * KHD + d]; for (int d = 0; d < VHD; d++) vv[(size_t)hh * VHD + d] = co[2 * KD + (size_t)hh * VHD + d]; } for (int hh = 0; hh < NVH; hh++) { float sq = 0, sk = 0; for (int d = 0; d < KHD; d++) { sq += ql[(size_t)hh * KHD + d] * ql[(size_t)hh * KHD + d]; sk += kl[(size_t)hh * KHD + d] * kl[(size_t)hh * KHD + d]; } const float iq = 1.0f / sqrtf(sq + 1e-6f), ik = 1.0f / sqrtf(sk + 1e-6f); for (int d = 0; d < KHD; d++) { ql[(size_t)hh * KHD + d] *= iq; kl[(size_t)hh * KHD + d] *= ik; } } float* rs = rec_state_[(size_t)l].data(); for (int hh = 0; hh < NVH; hh++) { const float eg = expf(gdn_small_[(size_t)l].ssm_a[hh] * softplus_f(rec.ah[(size_t)hh] + gdn_small_[(size_t)l].dt[hh])); const float beta = 1.0f / (1.0f + expf(-rec.bh[(size_t)hh])); float* sh = rs + (size_t)hh * KHD * VHD; const float* kh = &kl[(size_t)hh * KHD]; const float* vh = &vv[(size_t)hh * VHD]; for (int i = 0; i < KHD * VHD; i++) sh[i] *= eg; for (int vd = 0; vd < VHD; vd++) { float ssum = 0; for (int kd = 0; kd < KHD; kd++) ssum += sh[kd * VHD + vd] * kh[kd]; kvm[(size_t)vd] = ssum; } for (int vd = 0; vd < VHD; vd++) { const float delta = (vh[vd] - kvm[(size_t)vd]) * beta; for (int kd = 0; kd < KHD; kd++) sh[kd * VHD + vd] += kh[kd] * delta; } } } } return true; } Potential Integer Overflow in Matrix Multiplication In the Matrix::matvec function, the accumulation of partial products in acc_n is done using double precision, but the final result is cast to float. This could potentially lead to precision loss or unexpected behavior if the accumulated value exceeds the range that can be accurately represented in a float. Additionally, the code does not check for integer overflow during the matrix multiplication, which could lead to incorrect results if the matrix dimensions are very large. const int G = K / kG; std::vector<double> acc((size_t)N, 0.0); #pragma omp parallel for schedule(static) for (int n = 0; n < N; n++) { double acc_n = 0.0; for (int g = 0; g < G; g++) { const int8_t* cg = &codes[(size_t)(g * kG) * N + n]; double P = 0.0; for (int k = 0; k < kG; k++) P += (double)a[g * kG + k] * (double)cg[(size_t)k * N]; acc_n += P * ((double)ms[(size_t)g * N + n] * std::ldexp(1.0, (int)eref[(size_t)n])); } double ag_sum = 0.0; for (int g = 0; g < G; g++) { double s = 0.0; for (int k = 0; k < kG; k++) s += a[g * kG + k]; acc_n += s * (double)add[(size_t)g * N + n]; ag_sum += s; } (void)ag_sum; acc[(size_t)n] = acc_n; } for (int n = 0; n < N; n++) y[n] = (float)acc[(size_t)n]; } Potential Uninitialized Memory Access In the run_offline function, the cache_at vector is initialized with the size of hidden, but the elements are not explicitly initialized. If the head_step function fails or returns early, the cache_at vector might contain uninitialized data, which could lead to undefined behavior when accessed later in the function. std::vector<KvCache> cache_at(hidden.size()); for (size_t i = 0; i < hidden.size(); i++) { base.embed_into(all_tokens[i], emb); if (!head_step(w, hidden[i].data(), emb.data(), (int)i, hc, next, (i + 1 < hidden.size()) ? std::function<bool(const std::vector<float>&, std::vector<float>&)>(lm) : nullptr, lg)) { std::fprintf(stderr, "head step failed at %zu\n", i); return 1; } cache_at[i] = hc;

[github-actions]

Persistent review updated to latest commit cb36628

[github-actions]

Persistent review updated to latest commit f544032

[github-actions]

Persistent review updated to latest commit e32370d

[bong-water-water-bong]

VERDICT: READY

At e32370d9, the rebased head (own paths byte-identical to f5440322 — git diff --stat over npu/mtp/,
npu/forward/{include,src}/q4nx_forward.*, tools/mtp_acceptance.cpp, the harness and the two tests → empty), build green,
state clean.

Everything from my reviews stands, including the one I verified myself: the C++ head reproduces the Python reference at
rel-L1 1.933e-04, max|d| 1.758e-03, cosine 0.999999981 on a fixture I generated from mtp_head_ref.head_step (pos=3) —
i.e. the int16-mantissa limit, not merely "close". mtp_draft_test is green at 69 assertions with pos threaded through
DraftHead::draft, and the new commits were read: commit_verify() now truncates the KV caches to
verify_pos_ + accepted, restores snap_conv_/snap_rec_ and re-applies only the accepted rows through
gdn_apply_rows(0, accepted), which is correct for a recurrence and cheaper than replaying the whole batch.

Two notes to carry, neither blocking:

1. ship the head-vs-reference fixture (three 20 KB files, or the ~15-line generator: default_rng(0) for hidden/emb, then
   head_step(..., pos=3)). The check works and its thresholds are tight; it just cannot be run by anyone who is not the
   author today.
2. "deferred GDN commit" is a snapshot-then-replay of the accepted rows, not #798's deferral — the commit subject
   overstates it a little. The behaviour is correct; the wording is what a later reader will trip on when they go looking for
   the deferred mechanism in the forward (it lives in draft.h's interface, and this adapter satisfies it with a snapshot).

tprs-ready is on. This is the head to merge on.