From 7eb633ca602a7b98f8c1cd7fcef7b26021bc666c Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 05:53:03 -0400 Subject: [PATCH 1/7] npu: LFM2's gated short convolution as host arithmetic, with known answers (Refs #1786) npu/forward/include/lfm2_shortconv.h: the block between an LFM2 conv layer's two projections (B*X, a causal depthwise convolution over the last L-1 values, gated by C), with its per-layer state, defined once for decode, batched prefill and state snapshots. Not wired into the forward yet; this is the reference the wiring will be checked against. Co-Authored-By: Claude Fable 5.2 --- CMakeLists.txt | 6 ++ npu/forward/include/lfm2_shortconv.h | 94 +++++++++++++++++ tests/lfm2_shortconv_test.cpp | 152 +++++++++++++++++++++++++++ 3 files changed, 252 insertions(+) create mode 100644 npu/forward/include/lfm2_shortconv.h create mode 100644 tests/lfm2_shortconv_test.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 2746671bd..2ee02342b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1374,6 +1374,12 @@ add_executable(ffn_cache_policy_test tests/ffn_cache_policy_test.cpp) target_include_directories(ffn_cache_policy_test PRIVATE npu/forward/include) add_test(NAME ffn_cache_policy COMMAND ffn_cache_policy_test) add_dependencies(onebit ffn_cache_policy_test) +# LFM2's gated short convolution as host arithmetic (npu/forward/include/lfm2_shortconv.h, engine#1786): the batched +# form must equal the token-by-token form bit for bit. Header-only, no model, no device. +add_executable(lfm2_shortconv_test tests/lfm2_shortconv_test.cpp) +target_include_directories(lfm2_shortconv_test PRIVATE npu/forward/include) +add_test(NAME lfm2_shortconv COMMAND lfm2_shortconv_test) +add_dependencies(onebit lfm2_shortconv_test) # The device weight layout contract (q4nx_weight_layout): pure index math, no device # and no runtime. CI builds an explicit target list, so the check compiles itself with diff --git a/npu/forward/include/lfm2_shortconv.h b/npu/forward/include/lfm2_shortconv.h new file mode 100644 index 000000000..88d2466b2 --- /dev/null +++ b/npu/forward/include/lfm2_shortconv.h @@ -0,0 +1,94 @@ +// Copyright 2026 bong-water-water-bong +// SPDX-License-Identifier: Apache-2.0 +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +// lfm2_shortconv.h -- the gated short convolution of an LFM2 layer, as host arithmetic (engine#1786). +// +// LFM2 is a hybrid: ten of LFM2-1.2B's sixteen layers have no attention at all. In place of it they run +// +// bcx = in_proj(norm(x)) H -> 3H, the three thirds are B, C and X +// bx = B * X elementwise +// conv[h] = sum_j taps[h][j] * seq[t + j][h] depthwise, causal, L taps: seq is the last L-1 values of bx +// followed by this token's, so taps[h][L-1] meets the present +// y = out_proj(C * conv) H -> H +// +// and then the same residual and feed-forward as an attention layer. The only state a layer carries from one token +// to the next is the last L-1 values of bx, per channel: (L-1) * H floats, where an attention layer carries a KV row. +// +// The two projections are ordinary GEMMs and belong to the device. What is here is everything between them, written +// once so that the three places that will need it (token-by-token decode, a batched prefill, and a snapshot of the +// state for the prefix cache) share one definition, and tested without a model: the batched form must equal the +// token-by-token form bit for bit, because a prefill that disagrees with decode is wrong at every later position. +// +// The definition follows the graph our own llama.cpp fork builds for this family (src/models/lfm2.cpp, +// build_shortconv_block), not any vendor's code. Layout of `taps`: [H][L] row-major, the container's +// `shortconv.conv.weight` after its bf16 decode. + +#pragma once + +#include +#include +#include + +namespace onebit::npu { + +// The conv state of one layer: the last L-1 values of bx, oldest first, [L-1][H] row-major. All zeros at position 0. +struct ShortConvState { + int hidden = 0; + int taps = 0; // L, the kernel length (config.json conv_L_cache) + std::vector history; // (L-1) * hidden + + void reset(int hidden_size, int kernel_len) { + hidden = hidden_size; + taps = kernel_len; + history.assign(size_t(kernel_len > 1 ? kernel_len - 1 : 0) * size_t(hidden_size), 0.0f); + } + bool ready() const { return hidden > 0 && taps > 0 && history.size() == size_t(taps - 1) * size_t(hidden); } +}; + +// One token. `bcx` is in_proj's output, 3*H values [B | C | X]; `taps` is [H][L]; `y` receives H values, the input of +// out_proj. The state is advanced by one position. Returns false, touching nothing, when the state is not sized. +inline bool shortconv_step(ShortConvState& st, const float* bcx, const float* taps, float* y) { + if (!st.ready() || bcx == nullptr || taps == nullptr || y == nullptr) return false; + const int H = st.hidden, L = st.taps; + const float* B = bcx; + const float* C = bcx + H; + const float* X = bcx + 2 * H; + for (int h = 0; h < H; ++h) { + const float bx = B[h] * X[h]; + const float* k = taps + size_t(h) * size_t(L); + float acc = 0.0f; + for (int j = 0; j < L - 1; ++j) acc += k[j] * st.history[size_t(j) * size_t(H) + size_t(h)]; + acc += k[L - 1] * bx; + y[h] = C[h] * acc; + } + // slide the window: drop the oldest row, append this token's bx + if (L > 1) { + if (L > 2) std::memmove(st.history.data(), st.history.data() + H, size_t(L - 2) * size_t(H) * sizeof(float)); + float* last = st.history.data() + size_t(L - 2) * size_t(H); + for (int h = 0; h < H; ++h) last[h] = B[h] * X[h]; + } + return true; +} + +// M tokens in order: `bcx` is [M][3H] row-major, `y` is [M][H]. Defined as M calls of shortconv_step(), in that +// order, so there is nothing for it to disagree with. +inline bool shortconv_rows(ShortConvState& st, const float* bcx, int rows, const float* taps, float* y) { + if (rows < 0) return false; + for (int r = 0; r < rows; ++r) + if (!shortconv_step(st, bcx + size_t(r) * 3 * size_t(st.hidden), taps, y + size_t(r) * size_t(st.hidden))) return false; + return true; +} + +} // namespace onebit::npu diff --git a/tests/lfm2_shortconv_test.cpp b/tests/lfm2_shortconv_test.cpp new file mode 100644 index 000000000..fd6e1c28f --- /dev/null +++ b/tests/lfm2_shortconv_test.cpp @@ -0,0 +1,152 @@ +// Copyright 2026 bong-water-water-bong +// SPDX-License-Identifier: Apache-2.0 +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +// lfm2_shortconv_test: the gated short convolution (npu/forward/include/lfm2_shortconv.h) against the definition +// written out directly, with no model: a causal depthwise convolution over the whole sequence, zero history before +// position 0. Token-by-token, in one batch, and in uneven batches must all give the same bits. + +#include "lfm2_shortconv.h" + +#include +#include +#include +#include +#include + +using onebit::npu::ShortConvState; +using onebit::npu::shortconv_rows; +using onebit::npu::shortconv_step; + +static int failures = 0; +static void check(bool ok, const std::string& what) { + if (!ok) { + std::printf("FAIL %s\n", what.c_str()); + ++failures; + } +} + +// Deterministic inputs with a little of everything: both signs, zeros, different scales per channel. +static float bcx_of(int t, int i, int H) { + const int third = i / H, h = i % H; + float v = std::sin(0.37f * float(t + 1) * float(h + 1) + 1.3f * float(third)); + if ((t + h + third) % 7 == 0) v = 0.0f; + return v * (1.0f + 0.25f * float(h % 4)); +} +static float tap_of(int h, int j) { return 0.5f * std::cos(0.9f * float(h + 1) + 1.7f * float(j)) + (j == 0 ? 0.125f : 0.0f); } + +// The definition, on the whole sequence at once: y[t][h] = C[t][h] * sum_j taps[h][j] * bx[t - (L-1) + j][h], bx = 0 before 0. +// The sum runs in the same order as the code under test, so equality can be exact. +static std::vector reference(int T, int H, int L, const std::vector& bcx, const std::vector& taps) { + std::vector y(size_t(T) * size_t(H)); + for (int t = 0; t < T; ++t) + for (int h = 0; h < H; ++h) { + float acc = 0.0f; + for (int j = 0; j < L; ++j) { + const int src = t - (L - 1) + j; + const float bx = src < 0 ? 0.0f : bcx[size_t(src) * 3 * H + h] * bcx[size_t(src) * 3 * H + 2 * H + h]; + acc += taps[size_t(h) * L + j] * bx; + } + y[size_t(t) * H + h] = bcx[size_t(t) * 3 * H + H + h] * acc; + } + return y; +} + +int main() { + for (int L : {1, 2, 3, 4}) { + const int H = 37, T = 23; + std::vector bcx(size_t(T) * 3 * H), taps(size_t(H) * L); + for (int t = 0; t < T; ++t) + for (int i = 0; i < 3 * H; ++i) bcx[size_t(t) * 3 * H + i] = bcx_of(t, i, H); + for (int h = 0; h < H; ++h) + for (int j = 0; j < L; ++j) taps[size_t(h) * L + j] = tap_of(h, j); + const std::vector want = reference(T, H, L, bcx, taps); + const std::string tag = " (L=" + std::to_string(L) + ")"; + + ShortConvState one; + one.reset(H, L); + check(one.ready() && one.history.size() == size_t(L - 1) * H, "the state is L-1 rows of H" + tag); + std::vector got(size_t(T) * H); + bool ok = true; + for (int t = 0; t < T; ++t) ok = ok && shortconv_step(one, &bcx[size_t(t) * 3 * H], taps.data(), &got[size_t(t) * H]); + check(ok && std::memcmp(got.data(), want.data(), want.size() * sizeof(float)) == 0, + "token by token it is the causal depthwise convolution of B*X, gated by C, bit for bit" + tag); + + ShortConvState all; + all.reset(H, L); + std::vector batched(size_t(T) * H); + check(shortconv_rows(all, bcx.data(), T, taps.data(), batched.data()) && + std::memcmp(batched.data(), want.data(), want.size() * sizeof(float)) == 0, + "one batch of every token gives the same bits" + tag); + check(all.history == one.history, "and leaves the same state behind" + tag); + + ShortConvState mixed; // a prefill in uneven chunks, then decode: 5 + 1 + 9 + the rest one at a time + mixed.reset(H, L); + std::vector m(size_t(T) * H); + int at = 0; + bool mok = true; + for (int chunk : {5, 1, 9}) { + mok = mok && shortconv_rows(mixed, &bcx[size_t(at) * 3 * H], chunk, taps.data(), &m[size_t(at) * H]); + at += chunk; + } + for (; at < T; ++at) mok = mok && shortconv_step(mixed, &bcx[size_t(at) * 3 * H], taps.data(), &m[size_t(at) * H]); + check(mok && std::memcmp(m.data(), want.data(), want.size() * sizeof(float)) == 0 && mixed.history == one.history, + "uneven prefill chunks followed by decode steps give the same bits and state" + tag); + + // a snapshot of the state is all that a resume needs: copy it at position 11, replay from there + ShortConvState first; + first.reset(H, L); + std::vector scratch(size_t(T) * H); + shortconv_rows(first, bcx.data(), 11, taps.data(), scratch.data()); + ShortConvState resumed = first; + std::vector tail(size_t(T - 11) * H); + check(shortconv_rows(resumed, &bcx[size_t(11) * 3 * H], T - 11, taps.data(), tail.data()) && + std::memcmp(tail.data(), &want[size_t(11) * H], tail.size() * sizeof(float)) == 0, + "a copy of the state at position 11 resumes to the same bits" + tag); + + // the history really is the last L-1 values of B*X, oldest first + bool hist = true; + for (int j = 0; j < L - 1 && hist; ++j) + for (int h = 0; h < H; ++h) { + const int src = T - (L - 1) + j; + const float bx = bcx[size_t(src) * 3 * H + h] * bcx[size_t(src) * 3 * H + 2 * H + h]; + if (one.history[size_t(j) * H + h] != bx) { hist = false; break; } + } + check(hist, "the state is the last L-1 values of B*X, oldest first" + tag); + } + { // the present token meets the LAST tap, and the gate C multiplies the result, not the input + ShortConvState st; + st.reset(1, 3); + const float taps[3] = {100.0f, 10.0f, 1.0f}; + float y = 0.0f; + const float t0[3] = {2.0f, 5.0f, 3.0f}; // B, C, X -> bx 6 + check(shortconv_step(st, t0, taps, &y) && y == 5.0f * (1.0f * 6.0f), "at position 0 only the last tap sees anything"); + const float t1[3] = {1.0f, 2.0f, 4.0f}; // bx 4 + check(shortconv_step(st, t1, taps, &y) && y == 2.0f * (10.0f * 6.0f + 1.0f * 4.0f), "the token before meets the middle tap"); + const float t2[3] = {1.0f, 1.0f, 1.0f}; // bx 1 + check(shortconv_step(st, t2, taps, &y) && y == 1.0f * (100.0f * 6.0f + 10.0f * 4.0f + 1.0f * 1.0f), + "and the one before that the first tap"); + } + { // refusals: an unsized state, null buffers, a negative row count + ShortConvState st; + float b[3] = {1, 1, 1}, k[1] = {1}, y = 7.0f; + check(!shortconv_step(st, b, k, &y) && y == 7.0f, "an unsized state is refused and nothing is written"); + st.reset(1, 1); + check(!shortconv_step(st, nullptr, k, &y) && !shortconv_step(st, b, nullptr, &y) && !shortconv_step(st, b, k, nullptr), + "null buffers are refused"); + check(!shortconv_rows(st, b, -1, k, &y) && shortconv_rows(st, b, 0, k, &y), "a negative row count is refused; zero rows is a no-op"); + } + if (failures == 0) std::printf("lfm2_shortconv: ok\n"); + return failures == 0 ? 0 : 1; +} From 39013209fea17677f9e7be9367d07ee9b5168406 Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 06:42:20 -0400 Subject: [PATCH 2/7] npu: run LFM2's gated short-convolution layers in the forward A layer that carries shortconv.* tensors is no longer refused: in_proj (H -> 3H) derives one fused-projection design, out_proj reuses the O design, and the gate + depthwise causal conv between them is lfm2_shortconv.h on the host, shared by decode, batched prefill and the verify rollback. The state lives in conv_state_, so snapshot/restore/reset cover it. norm_eps is read when a config has no rms_norm_eps. Co-Authored-By: Claude Fable 5.2 --- npu/forward/include/lfm2_shortconv.h | 13 ++ npu/forward/include/q4nx_forward.h | 15 ++ npu/forward/include/q4nx_pack.h | 7 + npu/forward/src/q4nx_forward.cpp | 202 +++++++++++++++++++++++++-- npu/forward/src/q4nx_pack.c | 7 +- npu/forward/src/q4nx_semantics.c | 2 + tests/lfm2_shortconv_test.cpp | 26 ++++ tests/npu_design_grid_test.c | 36 +++++ 8 files changed, 293 insertions(+), 15 deletions(-) diff --git a/npu/forward/include/lfm2_shortconv.h b/npu/forward/include/lfm2_shortconv.h index 88d2466b2..624f7d5c3 100644 --- a/npu/forward/include/lfm2_shortconv.h +++ b/npu/forward/include/lfm2_shortconv.h @@ -82,6 +82,19 @@ inline bool shortconv_step(ShortConvState& st, const float* bcx, const float* ta return true; } +// Advance the state by one position WITHOUT computing an output: `bx` is that position's B*X, H values. This is +// what a rolled-back batch re-applies for the rows it keeps (the outputs were already produced; only the window +// has to match the accepted prefix). Leaves exactly the state shortconv_step() leaves for the same token. +inline bool shortconv_push(ShortConvState& st, const float* bx) { + if (!st.ready() || bx == nullptr) return false; + const int H = st.hidden, L = st.taps; + if (L > 1) { + if (L > 2) std::memmove(st.history.data(), st.history.data() + H, size_t(L - 2) * size_t(H) * sizeof(float)); + std::memcpy(st.history.data() + size_t(L - 2) * size_t(H), bx, size_t(H) * sizeof(float)); + } + return true; +} + // M tokens in order: `bcx` is [M][3H] row-major, `y` is [M][H]. Defined as M calls of shortconv_step(), in that // order, so there is nothing for it to disagree with. inline bool shortconv_rows(ShortConvState& st, const float* bcx, int rows, const float* taps, float* y) { diff --git a/npu/forward/include/q4nx_forward.h b/npu/forward/include/q4nx_forward.h index 22901baaf..05504c4a2 100644 --- a/npu/forward/include/q4nx_forward.h +++ b/npu/forward/include/q4nx_forward.h @@ -657,6 +657,21 @@ class Q4nxNpuForward { bool gdn_layer_batch(int l, const std::vector& xn, int M, std::vector& out); bool load_gdn_layer(int l); + + // ── LFM2 gated short-convolution layers ────────────────────────────── + // A layer that carries shortconv.* tensors has no attention: in_proj (H -> 3H, the QKVLIN design, weight + // key 0) and out_proj (H -> H, the O design, key 1, the slot an attention layer's o_proj has) run on the + // device, and everything between them is lfm2_shortconv.h on the host. The only state is the last L-1 + // values of B*X per channel, kept in conv_state_[l] so snapshot/restore/reset cover it with no new code. + bool shortconv_ = false; // the model has at least one such layer + int sc_taps_ = 0; // L, from shortconv.conv.weight's own shape + std::vector sc_layer_; // [NL] 1 = conv layer + std::vector> sc_w_; // [NL][H * L] the depthwise taps, bf16-decoded once + std::vector> sc_rows_; // [NL][rows * H] B*X of a verify batch's rows (see begin_verify) + bool is_conv_layer(int l) const { return shortconv_ && l >= 0 && l < (int)sc_layer_.size() && sc_layer_[(size_t)l]; } + bool conv_layer(int l, const std::vector& xn, std::vector& out); + bool conv_layer_batch(int l, const std::vector& xn, int M, std::vector& out); + void conv_apply_rows(int rows); // re-apply the first `rows` recorded rows after a restore std::string err_; std::string backend_ = "uninitialised"; bool ready_ = false; diff --git a/npu/forward/include/q4nx_pack.h b/npu/forward/include/q4nx_pack.h index a634a64c7..6db9a5fd3 100644 --- a/npu/forward/include/q4nx_pack.h +++ b/npu/forward/include/q4nx_pack.h @@ -246,6 +246,13 @@ typedef struct { int qk_nope; /* per-head qk width excluding rope */ int dense_im; /* leading dense layer's intermediate width */ int mla_kva_pad; /* KVA design N pad: ceil((kv_lora+rope_dim)/128)*128 */ + /* LFM2's gated short-convolution layers: in_proj is one fused (K=H, N=3*H) + * projection [B | C | X] and out_proj is (K=H, N=H), which is the O design + * whenever NH*HD == H. shortconv_n is in_proj's N; 0 means the model has no + * such layer. It derives one QKVLIN design and nothing else (no Z), and is + * exclusive with lin_kd: a model has GDN layers or conv layers, not both. + * Appended last so positional initialisers keep their meaning. */ + int shortconv_n; } Q4nxModelDims; /* No model-specific geometry constant lives here by design: every caller diff --git a/npu/forward/src/q4nx_forward.cpp b/npu/forward/src/q4nx_forward.cpp index 4466017fa..bab266aac 100644 --- a/npu/forward/src/q4nx_forward.cpp +++ b/npu/forward/src/q4nx_forward.cpp @@ -27,6 +27,7 @@ #include "pad_head_rows.h" // engine#1990: the padded-head row rule, header-only for its arm #include "batch_width.h" #include "ffn_cache_policy.h" +#include "lfm2_shortconv.h" // LFM2 conv layers: the host part between the two projections #include "q4nx_state_planes.h" // the plane-copy half of a snapshot, device-free and unit-tested #include "file_create.h" // dumps are created 0600, never through a symlink #include "model/op_dispatch.h" // the optional-op fallback seam (engine#803 A3) @@ -936,20 +937,58 @@ bool Q4nxNpuForward::init_impl(xrt::device* dev, const char* model_path, std::to_string(mw_->config.num_layers) + " layers in model.q4nx"; return false; } - // A layer kind nothing here implements is refused BY NAME (the same rule as an unknown layer_types entry below). - // LFM2 declares its hybrid through `full_attn_idxs`, not `layer_types`, so that check cannot see it: its ten - // gated short-convolution layers carry no self_attn.* tensors, and run as attention they fail in the first - // dequant with a size error that names neither the layer nor the kind. + // LFM2 declares its hybrid through `full_attn_idxs`, not `layer_types`, so the layer_types check below cannot see + // it: its gated short-convolution layers carry shortconv.* tensors and no self_attn.* at all. The container is + // the authority on which layers they are (model_shortconv_layers), and on the kernel length (the taps' own + // shape). What is refused, by name, is what the conv path does not cover: a kernel length that differs between + // layers or does not match the tensor's bytes, and the image layouts (pb16i/t2s), whose weights are sidecar + // images this path has no tags for. { int first_conv = -1; const int n_conv = model_shortconv_layers(mw_, cfg.num_layers, &first_conv); - if (n_conv > 0) { - err_ = "model.q4nx has " + std::to_string(n_conv) + " gated short-convolution layer(s) of " + - std::to_string(cfg.num_layers) + " (the first is model.layers." + std::to_string(first_conv) + - ".shortconv.*, LFM2's hybrid), a layer kind this forward does not implement (it runs " - "full_attention, sliding-window and linear/GDN layers), so serving it would run those layers as " - "attention; file the kind you need instead of converting it"; - return false; + shortconv_ = n_conv > 0; + sc_layer_.assign((size_t)cfg.num_layers, 0); + sc_taps_ = 0; + if (shortconv_) { + if (pb16i_) { + err_ = "model.q4nx has gated short-convolution layers (LFM2's hybrid) and config.json declares an " + "image weight layout (pb16i/t2s); the conv layers' projections have no image form here"; + return false; + } + for (int l = 0; l < cfg.num_layers; l++) { + const LayerWeights* lw = &mw_->layers[l]; + if (lw->shortconv_in_proj_weight.shape[0] == 0) continue; + sc_layer_[(size_t)l] = 1; + const TensorDesc* t = &lw->shortconv_conv_weight; + const int L = (int)t->shape[1]; + const std::string where = "model.layers." + std::to_string(l) + ".shortconv"; + if (lw->shortconv_out_proj_weight.shape[0] == 0 || t->shape[0] == 0) { + err_ = where + " has in_proj but not conv.weight and out_proj.weight"; + return false; + } + if ((int)t->shape[0] != cfg.hidden_size || L < 1 || L > 64 || + (size_t)t->data_size != (size_t)cfg.hidden_size * (size_t)L * 2) { + err_ = where + ".conv.weight is not bf16 [hidden_size][L] (shape " + + std::to_string((long long)t->shape[0]) + " x " + std::to_string((long long)t->shape[1]) + + ", " + std::to_string((unsigned long long)t->data_size) + " bytes, hidden_size " + + std::to_string(cfg.hidden_size) + ")"; + return false; + } + if (sc_taps_ == 0) sc_taps_ = L; + else if (L != sc_taps_) { + err_ = where + ".conv.weight has kernel length " + std::to_string(L) + " where layer " + + std::to_string(first_conv) + " has " + std::to_string(sc_taps_) + + "; one kernel length per model is what this forward runs"; + return false; + } + } + if (cfg.num_attention_heads * cfg.head_dim != cfg.hidden_size) { + err_ = "gated short-convolution layers need out_proj (hidden_size -> hidden_size) to be the attention " + "output geometry, and num_attention_heads * head_dim is not hidden_size here"; + return false; + } + fprintf(stderr, " [forward] %d gated short-convolution layer(s) of %d, kernel length %d (LFM2 hybrid)\n", + n_conv, cfg.num_layers, sc_taps_); } } @@ -1236,6 +1275,7 @@ bool Q4nxNpuForward::init_impl(xrt::device* dev, const char* model_path, dims.qk_nope = qk_nope_; dims.dense_im = dense_im_; dims.H = c_.H; dims.NH = c_.NH; dims.NKV = c_.NKV; dims.HD = c_.HD; dims.IM = c_.IM; + dims.shortconv_n = shortconv_ ? 3 * c_.H : 0; // LFM2 conv in_proj [B | C | X] // Activation tile of the PREFILL family. It used to be the only family (a // bare 128 here decided the whole engine's M); decode now selects the M=16 // twin instead -- see the m16 block after q4nx_derive_designs() and @@ -1355,6 +1395,27 @@ bool Q4nxNpuForward::init_impl(xrt::device* dev, const char* model_path, rec_state_.assign(NL, std::vector((size_t)lin_nvh_ * (size_t)lin_khd_ * (size_t)lin_vhd_, 0.0f)); } + // LFM2 conv layers: the two projections' f32 scratch (read only by assemble_B until the packed weight is on + // the device), the taps decoded once, and the per-layer state -- (L-1)*H floats on a conv layer, nothing on an + // attention layer -- in conv_state_, where snapshot/restore/reset already look. + if (shortconv_) { + if (lin_kd_ > 0) { err_ = "a model with both linear_attention and short-convolution layers is not implemented"; return false; } + wqkvlin_.resize((size_t)3 * (size_t)H * (size_t)H); + woutlin_.resize((size_t)H * (size_t)q4nx_pad_k(H)); + sc_w_.assign((size_t)NL, std::vector()); + sc_rows_.assign((size_t)NL, std::vector()); + conv_state_.assign((size_t)NL, std::vector()); + for (int l = 0; l < NL; l++) { + if (!sc_layer_[(size_t)l]) continue; + const size_t n = (size_t)H * (size_t)sc_taps_; + const uint16_t* p = bf16_data(mw_, &mw_->layers[l].shortconv_conv_weight, n); + if (!p) { err_ = "model.layers." + std::to_string(l) + ".shortconv.conv.weight has no data"; return false; } + sc_w_[(size_t)l].resize(n); + for (size_t i = 0; i < n; i++) sc_w_[(size_t)l][i] = bf16_to_f32(p[i]); + conv_state_[(size_t)l].assign((size_t)(sc_taps_ - 1) * (size_t)H, 0.0f); + } + } + fprintf(stderr, " [forward] derived %d designs (fused_gu=%d, H=%d NH=%d NKV=%d HD=%d IM=%d)\n", ndesigns_, dims.fused_gu, dims.H, dims.NH, dims.NKV, dims.HD, dims.IM); @@ -2047,7 +2108,12 @@ bool Q4nxNpuForward::load_layer(int l) { // copies of the big projections are never needed: skipping them here is the // difference between a first token that dequantises the whole 17.9 GB model and // one that opens 13.4 GB of images. The norms and every small member still load. - if (!xd_->pb16i()) { + if (is_conv_layer(l)) { + // LFM2 conv layer: two projections where an attention layer has four. Same f32-until-packed rule. + LhScope _t_p(LH_LOAD_PROJ); + if (!deq(&lw->shortconv_in_proj_weight, wqkvlin_.data(), 3 * H, H)) return false; + if (!deq(&lw->shortconv_out_proj_weight, woutlin_.data(), H, q4nx_pad_k(H))) return false; + } else if (!xd_->pb16i()) { if (!deq(&lw->q_proj_weight, wq_.data(), NH * hd * (attn_gate_ ? 2 : 1), kpad_)) return false; // q|gate rows if (!kv_shared) { @@ -2443,6 +2509,7 @@ void Q4nxNpuForward::begin_verify(int rows, int start_pos) { for (size_t l = 0; l < rec_state_.size(); l++) snap_rec_[l] = rec_state_[l]; row_in_.resize(rec_state_.size()); for (auto& v : row_in_) v.assign((size_t)rows, GdnRowIn{}); + for (auto& v : sc_rows_) v.clear(); // conv layers record their rows' B*X as the batch runs } // Keep the first `accepted` rows of the verify batch and drop the rest: restore the snapshot, then re-apply the @@ -2481,6 +2548,7 @@ bool Q4nxNpuForward::commit_verify_rows(int accepted) { } const auto cv_t1 = std::chrono::steady_clock::now(); verify_open_ = false; // the replay below must not re-record + if (accepted > 0 && shortconv_) conv_apply_rows(accepted); if (accepted > 0 && !gdn_apply_rows(0, accepted)) return false; if (cv_prof) { const auto cv_t2 = std::chrono::steady_clock::now(); @@ -2512,6 +2580,7 @@ bool Q4nxNpuForward::commit_verify_rows(int accepted) { accepted, maxd, scale, ok_state ? "OK" : "MISMATCH"); } for (auto& v : row_in_) v.clear(); + for (auto& v : sc_rows_) v.clear(); return true; } @@ -2962,6 +3031,107 @@ bool Q4nxNpuForward::gdn_layer_batch(int l, const std::vector& xn, int M, return true; } +// -------------------------------------------------------------------------- +// LFM2 gated short-convolution layers +// -------------------------------------------------------------------------- +// One token: in_proj on the device, the gate + depthwise causal conv on the host (lfm2_shortconv.h, the one +// definition decode, prefill and the tests share), out_proj on the device. The state is borrowed from +// conv_state_[l] for the call and handed back, so there is one copy of it. +bool Q4nxNpuForward::conv_layer(int l, const std::vector& xn, std::vector& out) { + const int H = c_.H; + const int i_in = idx_of(Q4NX_DESIGN_QKVLIN); + const int i_o = idx_of(Q4NX_DESIGN_O); + if (i_in < 0 || i_o < 0) { err_ = "short-convolution designs missing (QKV 3H / O)"; return false; } + std::vector bcx; + { + const Q4nxDesignGeom* g = &designs_[i_in]; + Q4nxProjections p{}; + p.ple = wqkvlin_.data(); + if (!weight_cached(0) && q4nx_assemble_B(g, &p, B_.data()) != 0) { err_ = "assemble conv in_proj"; return false; } + npu_gemm(i_in, xn.data(), g->K, g->N, B_.data(), bcx, 0, "conv_in", g->real_k); + } + if (bcx.size() < (size_t)3 * (size_t)H) { err_ = "conv in_proj returned fewer than 3*hidden values"; return false; } + const Q4nxDesignGeom* go = &designs_[i_o]; + std::vector y((size_t)std::max(H, go->K), 0.0f); + onebit::npu::ShortConvState st; + st.hidden = H; st.taps = sc_taps_; + st.history.swap(conv_state_[(size_t)l]); + const bool ok = onebit::npu::shortconv_step(st, bcx.data(), sc_w_[(size_t)l].data(), y.data()); + st.history.swap(conv_state_[(size_t)l]); + if (!ok) { err_ = "short-convolution state of layer " + std::to_string(l) + " is not sized"; return false; } + if (onebit::npu::ops_dump_enabled()) onebit::npu::ops_dump("conv_core", l, 1, y.data(), H); + { + Q4nxProjections p{}; + p.o = woutlin_.data(); + if (!weight_cached(1) && q4nx_assemble_B(go, &p, B_.data()) != 0) { err_ = "assemble conv out_proj"; return false; } + npu_gemm(i_o, y.data(), go->K, go->N, B_.data(), out, 1, "conv_out", go->real_k); + } + return true; +} + +// M tokens: each projection is one M-row launch, and between them the rows go through shortconv_step() in position +// order -- the same calls conv_layer() makes -- so a batched prefill leaves exactly the state token-by-token decode +// would. A verify batch also keeps each row's B*X, the only thing needed to re-apply an accepted prefix. +bool Q4nxNpuForward::conv_layer_batch(int l, const std::vector& xn, int M, std::vector& out) { + const int H = c_.H; + const int i_in = idx_of(Q4NX_DESIGN_QKVLIN); + const int i_o = idx_of(Q4NX_DESIGN_O); + if (i_in < 0 || i_o < 0) { err_ = "short-convolution designs missing (QKV 3H / O)"; return false; } + std::vector BCX; + const Q4nxDesignGeom* g = &designs_[i_in]; + { + Q4nxProjections p{}; + p.ple = wqkvlin_.data(); + if (!weight_cached(0) && q4nx_assemble_B(g, &p, B_.data()) != 0) { err_ = "assemble conv in_proj"; return false; } + npu_gemm_rows(i_in, xn.data(), M, g->K, g->N, B_.data(), BCX, 0, "conv_in", g->real_k); + } + const size_t stride = (size_t)g->N; + if (stride < (size_t)3 * (size_t)H || BCX.size() < (size_t)M * stride) { + err_ = "conv in_proj returned fewer than rows * 3*hidden values"; + return false; + } + const Q4nxDesignGeom* go = &designs_[i_o]; + const size_t ystride = (size_t)std::max(H, go->K); + std::vector Y((size_t)M * ystride, 0.0f); + std::vector* rec = verify_open_ ? &sc_rows_[(size_t)l] : nullptr; + if (rec) rec->resize((size_t)M * (size_t)H); + onebit::npu::ShortConvState st; + st.hidden = H; st.taps = sc_taps_; + st.history.swap(conv_state_[(size_t)l]); + bool ok = true; + for (int m = 0; m < M && ok; m++) { + const float* row = BCX.data() + (size_t)m * stride; + if (rec) { + float* bx = rec->data() + (size_t)m * (size_t)H; + for (int h = 0; h < H; h++) bx[h] = row[h] * row[2 * H + h]; + } + ok = onebit::npu::shortconv_step(st, row, sc_w_[(size_t)l].data(), Y.data() + (size_t)m * ystride); + } + st.history.swap(conv_state_[(size_t)l]); + if (!ok) { err_ = "short-convolution state of layer " + std::to_string(l) + " is not sized"; return false; } + { + Q4nxProjections p{}; + p.o = woutlin_.data(); + if (!weight_cached(1) && q4nx_assemble_B(go, &p, B_.data()) != 0) { err_ = "assemble conv out_proj"; return false; } + npu_gemm_rows(i_o, Y.data(), M, go->K, go->N, B_.data(), out, 1, "conv_out", go->real_k); + } + return true; +} + +// After commit_verify_rows() has put the pre-batch state back: advance each conv layer over the accepted rows +// only. No projection is recomputed; the window is B*X, and the batch kept it. +void Q4nxNpuForward::conv_apply_rows(int rows) { + const int H = c_.H; + for (size_t l = 0; l < sc_rows_.size(); l++) { + if (!sc_layer_[l] || sc_rows_[l].size() < (size_t)rows * (size_t)H) continue; + onebit::npu::ShortConvState st; + st.hidden = H; st.taps = sc_taps_; + st.history.swap(conv_state_[l]); + for (int m = 0; m < rows; m++) onebit::npu::shortconv_push(st, sc_rows_[l].data() + (size_t)m * (size_t)H); + st.history.swap(conv_state_[l]); + } +} + // -------------------------------------------------------------------------- // Gemma4 per-layer embeddings (PLE) // -------------------------------------------------------------------------- @@ -3772,7 +3942,9 @@ bool Q4nxNpuForward::step_batch(const std::vector& tokens, int start_pos, rmsnorm(X.data() + (size_t)m * H, in_norm_[l].data(), H, XN.data() + (size_t)m * H); O.assign((size_t)M * H, 0.0f); - if (lin_kd_ > 0 && lt == Q4NX_LAYER_LINEAR) { + if (is_conv_layer(l)) { + if (!conv_layer_batch(l, XN, M, O)) { if (err_.empty()) err_ = "conv batch failed"; return false; } + } else if (lin_kd_ > 0 && lt == Q4NX_LAYER_LINEAR) { if (!gdn_layer_batch(l, XN, M, O)) { err_ = "gdn batch failed"; return false; } } else { const Q4nxDesignGeom* g = &designs_[i_qkv_l]; @@ -5027,7 +5199,9 @@ bool Q4nxNpuForward::step(int token, int pos, std::vector& logits, int ca // ---- attention ---- (or, for a GDN layer, gdn_layer below) rmsnorm(x.data(), in_norm_[l].data(), H, xn.data()); - if (lin_kd_ > 0 && lt == Q4NX_LAYER_LINEAR) { + if (is_conv_layer(l)) { + if (!conv_layer(l, xn, o)) return false; + } else if (lin_kd_ > 0 && lt == Q4NX_LAYER_LINEAR) { LhScope _tlh(LH_GDN); if (!gdn_layer(l, xn, o)) return false; } else { diff --git a/npu/forward/src/q4nx_pack.c b/npu/forward/src/q4nx_pack.c index 1c6006948..1a53600b1 100644 --- a/npu/forward/src/q4nx_pack.c +++ b/npu/forward/src/q4nx_pack.c @@ -40,7 +40,7 @@ int q4nx_num_designs(const Q4nxModelDims* d) { return 10; return (d->fused_gu ? 4 : 5) + (d->ple_dim > 0 ? 2 : 0) + (d->double_wide_mlp ? 2 : 0) + (has_second_attn(d) ? 2 : 0) + - (d->lin_kd > 0 ? 2 : 0); + (d->lin_kd > 0 ? 2 : 0) + (d->lin_kd <= 0 && d->shortconv_n > 0 ? 1 : 0); } /* MLA + MoE design set (see q4nx_num_designs). dims->HD must be the per-head @@ -463,6 +463,11 @@ int q4nx_derive_designs(const Q4nxModelDims* d, Q4nxDesignGeom* out) { i++; geom(&out[i], d->M, d->H, d->lin_vd, Q4NX_DESIGN_Z, "Z"); + } else if (d->shortconv_n > 0) { + /* LFM2 short-convolution in_proj: the same fused-projection role and + * kernel format as the GDN q|k|v, at N = 3*H, and no Z. */ + i++; + geom(&out[i], d->M, d->H, d->shortconv_n, Q4NX_DESIGN_QKVLIN, "QKV"); } return i + 1; diff --git a/npu/forward/src/q4nx_semantics.c b/npu/forward/src/q4nx_semantics.c index c623e660d..4aa491a5e 100644 --- a/npu/forward/src/q4nx_semantics.c +++ b/npu/forward/src/q4nx_semantics.c @@ -463,6 +463,8 @@ int q4nx_semantics_from_config(const char* json_path, const char* model_dir, /* epsilon */ if (get_num(json, "rms_norm_eps", &d) && d > 0) s->eps = (float)d; + /* LFM2 spells it norm_eps (and has no rms_norm_eps at all) */ + else if (get_num(json, "norm_eps", &d) && d > 0) s->eps = (float)d; /* ── Gemma4-style per-layer embeddings (PLE) + KV sharing ── */ if (get_num(json, "hidden_size_per_layer_input", &d) && d > 0) diff --git a/tests/lfm2_shortconv_test.cpp b/tests/lfm2_shortconv_test.cpp index fd6e1c28f..e67414c65 100644 --- a/tests/lfm2_shortconv_test.cpp +++ b/tests/lfm2_shortconv_test.cpp @@ -115,6 +115,32 @@ int main() { std::memcmp(tail.data(), &want[size_t(11) * H], tail.size() * sizeof(float)) == 0, "a copy of the state at position 11 resumes to the same bits" + tag); + // a rolled-back batch: run all T rows from a copy, put the copy back, re-apply only the first 7 rows' B*X. + // That must be the state 7 decode steps leave, and decode must carry on from it to the same bits. + { + ShortConvState base; + base.reset(H, L); + std::vector junk(size_t(T) * H); + shortconv_rows(base, bcx.data(), 4, taps.data(), junk.data()); // 4 tokens already committed + const ShortConvState snap = base; + shortconv_rows(base, &bcx[size_t(4) * 3 * H], T - 4, taps.data(), junk.data()); // the speculative batch + base = snap; // rejected past row 7 + bool pok = true; + std::vector bx(H); + for (int t = 4; t < 11; ++t) { + for (int h = 0; h < H; ++h) bx[h] = bcx[size_t(t) * 3 * H + h] * bcx[size_t(t) * 3 * H + 2 * H + h]; + pok = pok && onebit::npu::shortconv_push(base, bx.data()); + } + std::vector rest(size_t(T - 11) * H); + pok = pok && shortconv_rows(base, &bcx[size_t(11) * 3 * H], T - 11, taps.data(), rest.data()); + check(pok && base.history == one.history && + std::memcmp(rest.data(), &want[size_t(11) * H], rest.size() * sizeof(float)) == 0, + "restore + push of the accepted rows' B*X equals having decoded them" + tag); + ShortConvState unsized; + check(!onebit::npu::shortconv_push(unsized, bx.data()) && !onebit::npu::shortconv_push(base, nullptr), + "push refuses an unsized state and a null row" + tag); + } + // the history really is the last L-1 values of B*X, oldest first bool hist = true; for (int j = 0; j < L - 1 && hist; ++j) diff --git a/tests/npu_design_grid_test.c b/tests/npu_design_grid_test.c index c39539a75..5adbf52dd 100644 --- a/tests/npu_design_grid_test.c +++ b/tests/npu_design_grid_test.c @@ -831,7 +831,43 @@ static void test_head_tile_cache(void) { free(codes); free(ref); free(scales); free(sref); } +/* LFM2-1.2B: a hybrid whose conv layers need one more projection, in_proj (K=H, N=3*H). It is the fused-projection + * role at a new N and nothing else: no Z, the other four designs untouched, and off when the model has GDN layers. */ +static void shortconv_geometry(void) { + Q4nxModelDims d; + memset(&d, 0, sizeof d); + d.H = 2048; d.NH = 32; d.NKV = 8; d.HD = 64; d.IM = 8192; + d.fused_gu = 1; d.M = Q4NX_DESIGN_M_PREFILL; + Q4nxDesignGeom plain[Q4NX_MAX_DESIGNS], g[Q4NX_MAX_DESIGNS]; + const int n0 = q4nx_derive_designs(&d, plain); + d.shortconv_n = 3 * d.H; + check_int(q4nx_num_designs(&d), n0 + 1, "short-convolution layers add exactly one design to the count"); + const int n = q4nx_derive_designs(&d, g); + check_int(n, n0 + 1, "and exactly one to the derived table"); + Q4nxDesignGeom* in = find(g, n, Q4NX_DESIGN_QKVLIN); + check(in != NULL, "it is the fused-projection role"); + if (in) { + check_int(in->K, 2048, "in_proj K is the hidden size"); + check_int(in->N, 6144, "in_proj N is 3 * hidden: [B | C | X]"); + check(strcmp(in->elf_name, "full_i8_QKV_K2048_N6144.elf") == 0, "its artifact name is geometry-derived, like every other"); + } + check(find(g, n, Q4NX_DESIGN_Z) == NULL, "no Z design: a conv layer has no gate projection"); + int same = 1; + for (int i = 0; i < n0; i++) + if (g[i].role != plain[i].role || g[i].K != plain[i].K || g[i].N != plain[i].N || + strcmp(g[i].elf_name, plain[i].elf_name) != 0) same = 0; + check(same, "the attention and feed-forward designs are the ones the same model derives without it"); + Q4nxDesignGeom* o = find(g, n, Q4NX_DESIGN_O); + check(o && o->K == 2048 && o->N == 2048, "out_proj (H -> H) is the O design when NH*HD == H"); + d.lin_kd = 1024; d.lin_vd = 2048; + Q4nxDesignGeom both[Q4NX_MAX_DESIGNS]; + const int nb = q4nx_derive_designs(&d, both); + check_int(nb, n0 + 2, "with GDN dims set the GDN pair is derived and shortconv_n adds nothing"); + check_int(q4nx_num_designs(&d), nb, "and the count agrees with the table"); +} + int main(void) { + shortconv_geometry(); test_head_image(); test_head_tile_cache(); oga_geometry(); From d64e4f016a9c8508a75fb75fee3703c70984039b Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 07:10:23 -0400 Subject: [PATCH 3/7] npu: read the symmetric 4-bit tile kind, and LFM2's q/k norms The LFM2 packs store their projections as symmetric tiles: the same 5120-byte layout with every zero-point 0, a signed scale and a two's-complement code. Read as the asymmetric kind they correlate -0.39 with the upstream weights; the tile decoder now tells the kinds apart by the zero-point block. lfm2 also gets per-head q/k RMSNorm, which its attention layers carry. Co-Authored-By: Claude Fable 5.2 --- npu/forward/include/q4nx_dequant.h | 16 ++++++- npu/forward/src/q4nx_dequant.c | 20 +++++++-- npu/forward/src/q4nx_semantics.c | 9 ++-- tests/npu_dequant_bf16_test.c | 68 ++++++++++++++++++++++++++++++ tests/q4nx_layer_kinds_test.c | 16 +++++++ 5 files changed, 122 insertions(+), 7 deletions(-) diff --git a/npu/forward/include/q4nx_dequant.h b/npu/forward/include/q4nx_dequant.h index 6b2968b44..e19c44869 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/npu/forward/include/q4nx_dequant.h @@ -31,6 +31,17 @@ // and // value[row][col] = code * scale[g][row] + zero_point[g][row] // +// A tile of that size comes in two kinds, and the tile itself says which. The one above is the asymmetric kind +// (FastFlowLM lists such packs as "Q4_1"): codes 0..15, scales >= 0, and a zero-point that is the group's +// minimum. The symmetric kind (listed as "Q4_0"; the LFM2 packs) has the SAME byte layout with every zero-point +// stored as 0, a scale that carries a sign, and the code as a two's-complement nibble: +// value[row][col] = (code >= 8 ? code - 16 : code) * scale[g][row] +// Nothing in the container names the kind, so it is read off the tile: all 256 zero-points zero means symmetric. +// Measured 2026-11-11 on every 5120-byte tile of the fifteen packs on the box: 1,889,236 asymmetric tiles, none +// with an all-zero zero-point block; 285,696 tiles in the two LFM2 packs, every one of them all-zero. Read as +// asymmetric, an LFM2 projection correlates -0.39 with the same tensor of the upstream GGUF; read as above, 0.99. +// (An asymmetric tile whose 256 group minima are all exactly 0 would be a matrix with no negative weight.) +// // NOTE: this is NOT the layout asserted by the older tests/verify_i4_dequant.py // and tests/verify_q4nx_format.py (per-row 20-byte [scale][zp][16B nibbles] // groups with signed nibbles). That layout drives the reconstruction to NaN @@ -123,7 +134,10 @@ int q4nx_packed_grid(size_t nbytes, int rows, int cols, int* cb_out, int* rp_out int q4nx_dequant_tensor_packed(const uint8_t* data, size_t nbytes, int rows, int cols, int rows_packed, int cols_packed, int chunk_bytes, float* out); -/* Dequantize one 5120-byte tile into out[32][ld] (row-major, strided by ld). */ +/* 1 when the 5120-byte tile is the symmetric kind (every zero-point is +-0), 0 for the asymmetric kind. */ +int q4nx_tile_is_symmetric(const uint8_t* tile); + +/* Dequantize one 5120-byte tile, of either kind, into out[32][ld] (row-major, strided by ld). */ void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld); /* Dequantize a whole INT4 tensor of logical shape [rows, cols] whose packed diff --git a/npu/forward/src/q4nx_dequant.c b/npu/forward/src/q4nx_dequant.c index 98b95d212..21f0684a1 100644 --- a/npu/forward/src/q4nx_dequant.c +++ b/npu/forward/src/q4nx_dequant.c @@ -29,10 +29,19 @@ static inline uint16_t le16(const uint8_t* p) { return (uint16_t)((uint16_t)p[0] | ((uint16_t)p[1] << 8)); } +int q4nx_tile_is_symmetric(const uint8_t* tile) { + const uint8_t* zp_b = tile + Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW * 2; /* +512 */ + for (int i = 0; i < Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW; i++) + if ((le16(zp_b + 2 * i) & 0x7FFF) != 0) return 0; /* +0.0 and -0.0 both count as zero */ + return 1; +} + void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld) { const uint8_t* sc_b = tile; const uint8_t* zp_b = tile + Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW * 2; /* +512 */ const uint8_t* pk = tile + Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW * 4; /* +1024 */ + /* The symmetric kind (see the header): signed codes, signed scales, no offset. */ + const int symmetric = q4nx_tile_is_symmetric(tile); for (int r = 0; r < Q4NX_TILE_ROWS; r++) { const int h = r / 16; @@ -42,12 +51,17 @@ void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld) { const int j = c % Q4NX_GROUP_SIZE; const int region = h * Q4NX_GROUPS_PER_ROW + g; const uint8_t byte = pk[region * 256 + 8 * j + rr / 2]; - const int code = (rr % 2 == 0) ? (byte & 0xF) : (byte >> 4); + int code = (rr % 2 == 0) ? (byte & 0xF) : (byte >> 4); /* scale/zero-point index is GROUP-MAJOR: i = g*32 + r */ const int i = g * Q4NX_TILE_ROWS + r; const float s = bf16_to_f32(le16(sc_b + 2 * i)); - const float z = bf16_to_f32(le16(zp_b + 2 * i)); - out[(size_t)r * ld + c] = (float)code * s + z; + if (symmetric) { + if (code >= 8) code -= 16; /* two's complement nibble */ + out[(size_t)r * ld + c] = (float)code * s; + } else { + const float z = bf16_to_f32(le16(zp_b + 2 * i)); + out[(size_t)r * ld + c] = (float)code * s + z; + } } } } diff --git a/npu/forward/src/q4nx_semantics.c b/npu/forward/src/q4nx_semantics.c index 4aa491a5e..55ec216cb 100644 --- a/npu/forward/src/q4nx_semantics.c +++ b/npu/forward/src/q4nx_semantics.c @@ -35,7 +35,7 @@ static const char* kFields[Q4NX_SEM_COUNT] = { /* EMBED_SCALE */ "model_type (gemma* uses sqrt(hidden_size))", /* LOGIT_SOFTCAP */ "final_logit_softcapping", /* SLIDING_WINDOW */ "sliding_window", - /* EPS */ "rms_norm_eps", + /* EPS */ "rms_norm_eps (norm_eps when absent)", /* HYBRID_LAYERS */ "layer_types (count of entries != full_attention)", /* MROPE */ "rope_parameters.mrope_interleaved (rope_type)", /* PER_LAYER_ROPE */ "rope_parameters.{full,sliding}_attention sub-blocks", @@ -219,8 +219,11 @@ int q4nx_semantics_from_config(const char* json_path, const char* model_dir, * layer-0 input_layernorm weight it loads (a config-only answer would be a guess). */ s->rmsnorm_offset_probe = (is_qwen35 && !is_gemma) ? 1 : 0; - /* per-head q/k RMSNorm: qwen3* and gemma* have q_norm/k_norm; llama/phi do not */ - s->qk_norm = (is_qwen3 || is_gemma) ? 1 : 0; + /* LFM2 / LFM2.5 (Lfm2ForCausalLM): its attention layers carry q_norm/k_norm, applied before RoPE */ + const int is_lfm2 = starts_with(model_type, "lfm2") || starts_with(arch, "Lfm2"); + + /* per-head q/k RMSNorm: qwen3*, gemma* and lfm2 have q_norm/k_norm; llama/phi do not */ + s->qk_norm = (is_qwen3 || is_gemma || is_lfm2) ? 1 : 0; /* RoPE base. Qwen3.5 nests it under rope_parameters (scalar); Gemma4 nests * it PER LAYER TYPE under rope_parameters.{full,sliding}_attention with diff --git a/tests/npu_dequant_bf16_test.c b/tests/npu_dequant_bf16_test.c index 2ab1003e8..63621b7c8 100644 --- a/tests/npu_dequant_bf16_test.c +++ b/tests/npu_dequant_bf16_test.c @@ -192,7 +192,75 @@ static void tile_row_kind_arms(void) { printf(" tile-row kind: 3 kinds accepted, 4 tables the old rule took and the reader refuses are refused\n"); } +/* The two kinds of 5,120-byte tile. Same bytes for scales and codes; only the zero-point block differs, and that is + * what selects the reading: with a zero-point present the code is 0..15 and the value is code*scale + zero_point; + * with every zero-point 0 the code is a two's-complement nibble and the value is code*scale, scale carrying a sign. */ +static void tile_kinds(void) { + uint8_t* tile = (uint8_t*)calloc(1, Q4NX_TILE_BYTES); + float* out = (float*)malloc(sizeof(float) * Q4NX_TILE_ROWS * Q4NX_TILE_COLS); + if (!tile || !out) { check(0, "tile kinds: the buffers allocate"); free(tile); free(out); return; } + /* scales: exactly representable in bf16, negative on odd groups; codes: every nibble value appears */ + for (int g = 0; g < Q4NX_GROUPS_PER_ROW; g++) + for (int r = 0; r < Q4NX_TILE_ROWS; r++) + put_bf16(tile, (size_t)(g * Q4NX_TILE_ROWS + r), (g % 2 ? -1.0f : 1.0f) * 0.25f * (float)(1 + r % 4)); + for (int i = 0; i < 4096; i++) tile[1024 + i] = (uint8_t)((i * 37 + 11) & 0xFF); + #define CODE_AT(r, c) ((((r) % 16) % 2 == 0) ? (tile[1024 + (((r) / 16) * 8 + (c) / 32) * 256 + 8 * ((c) % 32) + ((r) % 16) / 2] & 0xF) \ + : (tile[1024 + (((r) / 16) * 8 + (c) / 32) * 256 + 8 * ((c) % 32) + ((r) % 16) / 2] >> 4)) + #define SCALE_AT(r, c) (((((c) / 32) % 2) ? -1.0f : 1.0f) * 0.25f * (float)(1 + (r) % 4)) + + check(q4nx_tile_is_symmetric(tile) == 1, "tile kinds: a zero-point block of zeros is the symmetric kind"); + q4nx_dequant_tile(tile, out, Q4NX_TILE_COLS); + int bad = 0, saw_neg_code = 0; + for (int r = 0; r < Q4NX_TILE_ROWS; r++) + for (int c = 0; c < Q4NX_TILE_COLS; c++) { + const int u = CODE_AT(r, c); + const int sgn = u >= 8 ? u - 16 : u; + if (u >= 8) saw_neg_code = 1; + if (out[r * Q4NX_TILE_COLS + c] != (float)sgn * SCALE_AT(r, c)) bad++; + } + check(saw_neg_code, "tile kinds: the fixture has codes of 8 and above, where the two readings differ"); + check(bad == 0, "tile kinds: symmetric is (two's-complement nibble) * signed scale, every element"); + + /* negative zero is still zero: the kind does not flip on a sign bit */ + tile[512 + 2 * 7 + 1] = 0x80; + check(q4nx_tile_is_symmetric(tile) == 1, "tile kinds: a -0.0 zero-point is still the symmetric kind"); + tile[512 + 2 * 7 + 1] = 0; + + /* one non-zero zero-point anywhere makes it the asymmetric kind, read as it always was */ + put_bf16(tile + 512, (size_t)(3 * Q4NX_TILE_ROWS + 5), -0.5f); + check(q4nx_tile_is_symmetric(tile) == 0, "tile kinds: one non-zero zero-point makes the tile asymmetric"); + q4nx_dequant_tile(tile, out, Q4NX_TILE_COLS); + bad = 0; + for (int r = 0; r < Q4NX_TILE_ROWS; r++) + for (int c = 0; c < Q4NX_TILE_COLS; c++) { + const float z = (c / 32 == 3 && r == 5) ? -0.5f : 0.0f; + if (out[r * Q4NX_TILE_COLS + c] != (float)CODE_AT(r, c) * SCALE_AT(r, c) + z) bad++; + } + check(bad == 0, "tile kinds: asymmetric is (code 0..15) * scale + zero_point, every element, as before"); + + /* the whole-tensor and one-row readers go through the same tile decoder */ + for (int i = 0; i < 512; i++) tile[512 + i] = 0; + float* row = (float*)malloc(sizeof(float) * Q4NX_TILE_COLS); + float* whole = (float*)malloc(sizeof(float) * Q4NX_TILE_ROWS * Q4NX_TILE_COLS); + if (row && whole) { + check(q4nx_dequant_tensor_chunked(tile, Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, Q4NX_TILE_COLS, Q4NX_TILE_BYTES, whole) == 0 && + q4nx_dequant_tile_row(tile, Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, Q4NX_TILE_COLS, Q4NX_TILE_BYTES, 9, row) == 0, + "tile kinds: the tensor and row readers accept a symmetric tile"); + bad = 0; + for (int c = 0; c < Q4NX_TILE_COLS; c++) { + const int u = CODE_AT(9, c); + const float want = (float)(u >= 8 ? u - 16 : u) * SCALE_AT(9, c); + if (whole[9 * Q4NX_TILE_COLS + c] != want || row[c] != want) bad++; + } + check(bad == 0, "tile kinds: and both read it as the tile decoder does"); + } else check(0, "tile kinds: the reader buffers allocate"); + #undef CODE_AT + #undef SCALE_AT + free(row); free(whole); free(tile); free(out); +} + int main(void) { + tile_kinds(); tile_row_kind_arms(); chunked_row_arms(); embed_row_arms(); diff --git a/tests/q4nx_layer_kinds_test.c b/tests/q4nx_layer_kinds_test.c index cfd7be284..0139b4e9d 100644 --- a/tests/q4nx_layer_kinds_test.c +++ b/tests/q4nx_layer_kinds_test.c @@ -124,6 +124,22 @@ int main(void) { "a substring match is quoted by name, so the container is refused instead of served"); check(s.hybrid_layers == 4, "all four refused kinds count as hybrid, as any non-full layer does"); + /* 6. LFM2 (Lfm2ForCausalLM): no layer_types at all (its hybrid is in full_attn_idxs and in the container's + * own tensors), q/k norms on its attention layers, and an epsilon spelled norm_eps. Served with qk_norm off + * and eps at the default it answered "France" to a question whose answer is "Paris". */ + write_config(dir, "{\"architectures\": [\"Lfm2ForCausalLM\"], \"model_type\": \"lfm2\", \"num_hidden_layers\": 16, " + "\"full_attn_idxs\": [2, 5, 8, 10, 12, 14], \"norm_eps\": 1e-05, \"block_norm_eps\": 1e-05, " + "\"conv_L_cache\": 3, \"rope_theta\": 1000000.0}", cfg, sizeof cfg); + q4nx_semantics_from_config(cfg, dir, &s); + check(s.qk_norm == 1, "lfm2: per-head q/k RMSNorm is on"); + check(s.eps > 0.9e-5f && s.eps < 1.1e-5f, "lfm2: epsilon is read from norm_eps when there is no rms_norm_eps"); + check(s.n_layer_types == 0 && s.unknown_layer_type[0] == '\0', "lfm2: no layer_types, so nothing is refused here"); + check(s.rmsnorm_offset == 0, "lfm2: plain x/rms * w norms"); + write_config(dir, "{\"model_type\": \"llama\", \"num_hidden_layers\": 2, \"rms_norm_eps\": 1e-06, \"norm_eps\": 1e-05}", cfg, sizeof cfg); + q4nx_semantics_from_config(cfg, dir, &s); + check(s.qk_norm == 0, "llama: still no q/k norm"); + check(s.eps > 0.9e-6f && s.eps < 1.1e-6f, "rms_norm_eps wins over norm_eps when a config has both"); + unlink(cfg); rmdir(dir); printf("\n%s (%d failure%s)\n", failures ? "FAILED" : "passed", failures, failures == 1 ? "" : "s"); From 0f6c891d6fe8cd9eb486c8a7f25b1804c59286fe Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 07:46:13 -0400 Subject: [PATCH 4/7] docs: LFM2 on the NPU route, the symmetric tile kind, and the row against FastFlowLM's runtime Co-Authored-By: Claude Fable 5.2 --- docs/npu.md | 46 ++++++++++++++++++++++++------ npu/forward/include/q4nx_dequant.h | 4 +++ 2 files changed, 41 insertions(+), 9 deletions(-) diff --git a/docs/npu.md b/docs/npu.md index 140a9a995..3286c1027 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -1820,7 +1820,7 @@ engine#587 measured the route on it): | `gemma4` | `FastFlowLM/Nemotron-3.5-Bolt-E4B-NPU2`, `FastFlowLM/Gemma4-12B-IT-NPU2` | hrx | not served: no gemma4 model type or kernels in the route | | `hunyuan-dense` | `FastFlowLM/Hy-MT2-1.8B-NPU2` | hrx | not served | | `gemma3` | `FastFlowLM/medgemma-1.5-4b-it-NPU2` | hrx | not served | -| `lfm2` | `FastFlowLM/LFM2-2.6B-NPU2`, `FastFlowLM/LFM2-1.2B-NPU2` | hrx | not served | +| `lfm2` | `FastFlowLM/LFM2-2.6B-NPU2`, `FastFlowLM/LFM2-1.2B-NPU2`, `FastFlowLM/LFM2.5-1.2B-NPU2` | hrx, npu | `LFM2.5-1.2B-NPU2` **served and measured** on the per-op route ("FastFlowLM's own containers" below); `LFM2-1.2B-NPU2` has the same tensor grid and tile kind and was not run; `LFM2-2.6B-NPU2` was not run | | `gpt-oss` | `FastFlowLM/GPT-OSS-20B-NPU2`, `FastFlowLM/GPT-OSS-Safeguard-20b-NPU2` | hrx | not served | | `hy_v3` | `FastFlowLM/Hy-Memory-30B-A3B-NPU2` | hrx | not served; the first table above records this one | @@ -1875,6 +1875,7 @@ does on a geometry that has a sealed set: `Qwen3-1.7B-NPU2` is 3.0 tok/s here an | `Llama-3.2-3B-NPU2` | `llama` | identical, `Paris.` | 1.7 to 1.8 tok/s | | `Phi4-mini-Instruct-NPU2` | `phi3` | identical, `Paris`, once the turn end is read from the template (below) | 1.4 to 1.5 tok/s | | `Qwen3-VL-4B-Instruct-NPU2`, text tower only | `qwen3` | identical, `Paris` | 1.4 to 1.5 tok/s | +| `LFM2.5-1.2B-NPU2` | `lfm2`, hybrid: 10 conv layers of 16 | identical, `Paris`; against FastFlowLM's runtime below | 3.5 to 4.6 tok/s | `D3-Nano-NPU2` is a decision pack: it loads on the `Qwen3.5-2B` kernel set, its readout converts with `tools/safetensors_to_q4nx.py decision`, and `/v1/systemone` returned three identical answers to the documented @@ -1887,18 +1888,45 @@ a stream that was sent after the generation had ended; reasoning that the templa decision-only pack. "What ends an answer" and [npu-chat-template.md](npu-chat-template.md) describe the rules that came out of them. -**Refused by name at load: the `lfm2` packs.** `LFM2.5-1.2B-NPU2` has all four dense kernel shapes on the box -and its container loads, but ten of its sixteen layers are gated short convolutions (`model.layers.N.shortconv.*`), -which the loader reads and the forward does not run. The serve exits at load with `forward init failed: -model.q4nx has 10 gated short-convolution layer(s) of 16 (the first is model.layers.0.shortconv.*, LFM2's hybrid), -a layer kind this forward does not implement …`. Before that refusal existed it ran those layers as attention, -failed its warm-up with a size error that named neither the layer nor the kind, and announced itself ready. +**The `lfm2` packs: a hybrid, and a second kind of weight tile.** Ten of `LFM2.5-1.2B-NPU2`'s sixteen layers +have no attention: they are gated short convolutions (`model.layers.N.shortconv.*`). The forward runs them +(engine#1786): `in_proj` (hidden to three times hidden, the fused-projection kernel at `K2048_N6144`) and `out_proj` +(the attention output kernel) on the array, and between them the gate and a three-tap causal depthwise convolution +on the host, whose only state is the last two positions of one product per channel. Which layers are convolutions +and how long the kernel is are read from the container's own tensors, not from `full_attn_idxs` or `conv_L_cache`. +The kernel directory is the `Llama-3.2-1B` set's four shapes plus that one `QKV_K2048_N6144` pair. + +Two things in these packs were read wrong before, and neither is about convolutions. Their 4-bit weights are a +**symmetric** kind of the 5,120-byte tile (FastFlowLM's list calls the pack `Q4_0` where the others are `Q4_1`): +the same byte layout with every zero-point stored as 0, a scale that carries a sign, and a two's-complement code. +Nothing in the container names the kind, so the decoder reads it off each tile: 256 zero-points that are all zero. +Of 1,889,236 tiles in the other packs on the box none has that block all zero; of 285,696 in the two LFM2 packs +every one has. Read as the usual kind, each LFM2 projection correlated -0.39 with the same tensor of the upstream +GGUF, and the serve answered in fluent-looking fragments of unrelated words; read as symmetric, 0.99, and the +pack's quantised `lm_head` sits within one step of its own bf16 embedding (mean 0.24 of a step). And the +attention layers carry per-head `q_norm`/`k_norm`, which the `lfm2` model type did not switch on: with only the +tiles fixed, "What is the capital of France? Answer with one word." was answered `France`. + +Measured on `LFM2.5-1.2B-NPU2` on 2026-11-11, same array, same `model.q4nx` (sha256 `b341d42021e617a5…`), five +greedy prompts of up to 40 tokens, FastFlowLM 1.0.12's runtime first and this route second in one ticket: +three answers byte-identical (`Paris`; the count from one to ten; a German sentence), two not. One leaves +FastFlowLM's text after 181 characters, the other at its first token (`Here's…` here, `Certainly!…` there). At +both points this route's choice is not a near-tie: the first token is `Here` -0.64 against `Certain` -1.26 on the +array and -0.66 against -1.23 with host float GEMMs on the same weights, and the upstream GGUF on the CPU opens +with `Here's` too. FastFlowLM's own `logprobs` for that token are `Certain` -0.81, `Here` -1.06, `S` -1.81: the +same candidates, the first two in the other order, on a 0.25 grid; its answer was the same three times running. +So the two runtimes agree on the model and differ in arithmetic, and which of them a reader should call right is +not something this row decides. Also on this pack: batched prefill equals token-by-token on every chunk checked +(`NPU_BATCH_CHECK=1`, chunks of 1 to 26 rows, all `MATCH`); a request answered from a restored prefix (351 of 352 +prompt tokens cached, conv state included) returns the bytes it returns with the prefix cache off; and decode was +4.6 tok/s at load 3 to 4 and 3.5 to 3.7 at load 8 to 10, against FastFlowLM's 48 to 83 tok/s on the same replies, +the gap "Per-op route" above describes. Not established by any row here: that a pack goes through the importer (`npu-pack` still reports these containers as recognised, not importable), anything about a vision or audio tower, and any pack whose geometry has no kernel set on the box. Of the 57 `*-NPU2` repos with a readable geometry on 2026-11-11, 19 had all four -dense kernel shapes already built in the blocked family; the families the forward may not model at all (`lfm2`, -`hunyuan_v1_dense`, the multimodal towers) are a separate question from the shapes. +dense kernel shapes already built in the blocked family; the families the forward may not model at all +(`hunyuan_v1_dense`, the multimodal towers) are a separate question from the shapes. ### Sealed kernel sets: the manifest is what says a set is ours diff --git a/npu/forward/include/q4nx_dequant.h b/npu/forward/include/q4nx_dequant.h index e19c44869..ea340ad97 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/npu/forward/include/q4nx_dequant.h @@ -41,6 +41,10 @@ // with an all-zero zero-point block; 285,696 tiles in the two LFM2 packs, every one of them all-zero. Read as // asymmetric, an LFM2 projection correlates -0.39 with the same tensor of the upstream GGUF; read as above, 0.99. // (An asymmetric tile whose 256 group minima are all exactly 0 would be a matrix with no negative weight.) +// The same tie test that settled the asymmetric layout settles this reading: LFM2.5-1.2B's lm_head is the +// quantised copy of its own bf16 token_embd, and over the first 2048 rows |dequant - embedding| / |scale| has +// mean 0.243 and maximum 1.000 (0.48 % of elements above 0.5, all of them the clip of +8 steps to +7). The +// other two candidate readings, (code - 8) * scale and unsigned code * scale, are 8 and 7 steps off on average. // // NOTE: this is NOT the layout asserted by the older tests/verify_i4_dequant.py // and tests/verify_q4nx_format.py (per-row 20-byte [scale][zp][16B nibbles] From 551db57db5c392cc86ba3ef8bec1a95695130a46 Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 08:01:43 -0400 Subject: [PATCH 5/7] docs: LFM2-1.2B-NPU2 row, one paired ticket against FastFlowLM's runtime Co-Authored-By: Claude Fable 5.2 --- docs/npu.md | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/docs/npu.md b/docs/npu.md index 696d6561c..37348df3f 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -1845,7 +1845,7 @@ engine#587 measured the route on it): | `gemma4` | `FastFlowLM/Nemotron-3.5-Bolt-E4B-NPU2`, `FastFlowLM/Gemma4-12B-IT-NPU2` | hrx | not served: no gemma4 model type or kernels in the route | | `hunyuan-dense` | `FastFlowLM/Hy-MT2-1.8B-NPU2` | hrx | not served | | `gemma3` | `FastFlowLM/medgemma-1.5-4b-it-NPU2` | hrx | not served | -| `lfm2` | `FastFlowLM/LFM2-2.6B-NPU2`, `FastFlowLM/LFM2-1.2B-NPU2`, `FastFlowLM/LFM2.5-1.2B-NPU2` | hrx, npu | `LFM2.5-1.2B-NPU2` **served and measured** on the per-op route ("FastFlowLM's own containers" below); `LFM2-1.2B-NPU2` has the same tensor grid and tile kind and was not run; `LFM2-2.6B-NPU2` was not run | +| `lfm2` | `FastFlowLM/LFM2-2.6B-NPU2`, `FastFlowLM/LFM2-1.2B-NPU2`, `FastFlowLM/LFM2.5-1.2B-NPU2` | hrx, npu | `LFM2.5-1.2B-NPU2` **served and measured** on the per-op route ("FastFlowLM's own containers" below); `LFM2-1.2B-NPU2` serves too (one paired ticket, same section); `LFM2-2.6B-NPU2` was not run | | `gpt-oss` | `FastFlowLM/GPT-OSS-20B-NPU2`, `FastFlowLM/GPT-OSS-Safeguard-20b-NPU2` | hrx | not served | | `hy_v3` | `FastFlowLM/Hy-Memory-30B-A3B-NPU2` | hrx | not served; the first table above records this one | @@ -1901,6 +1901,7 @@ does on a geometry that has a sealed set: `Qwen3-1.7B-NPU2` is 3.0 tok/s here an | `Phi4-mini-Instruct-NPU2` | `phi3` | identical, `Paris`, once the turn end is read from the template (below) | 1.4 to 1.5 tok/s | | `Qwen3-VL-4B-Instruct-NPU2`, text tower only | `qwen3` | identical, `Paris` | 1.4 to 1.5 tok/s | | `LFM2.5-1.2B-NPU2` | `lfm2`, hybrid: 10 conv layers of 16 | identical, `Paris`; against FastFlowLM's runtime below | 3.5 to 4.6 tok/s | +| `LFM2-1.2B-NPU2` | `lfm2`, hybrid: 10 conv layers of 16 | `Paris`; one request per prompt, against FastFlowLM's runtime below | 3.7 tok/s | `D3-Nano-NPU2` is a decision pack: it loads on the `Qwen3.5-2B` kernel set, its readout converts with `tools/safetensors_to_q4nx.py decision`, and `/v1/systemone` returned three identical answers to the documented @@ -1947,6 +1948,13 @@ prompt tokens cached, conv state included) returns the bytes it returns with the 4.6 tok/s at load 3 to 4 and 3.5 to 3.7 at load 8 to 10, against FastFlowLM's 48 to 83 tok/s on the same replies, the gap "Per-op route" above describes. +`LFM2-1.2B-NPU2` (sha256 `2755ab40cd927e3c…`) went through the same paired ticket on the same kernel directory, +at the merged head of that change: two of five answers byte-identical (`Paris`; the count from one to ten), three +that share 21, 58 and 53 characters with FastFlowLM's and then continue differently, each a fluent answer to the +question (a Fibonacci function that returns `0` for `n <= 0` where theirs returns a message). Decode 3.7 tok/s +against 42 to 47. Its logprobs at the three points were not taken, and it was asked each prompt once, so this row +does not carry the three-identical-answers condition. + Not established by any row here: that a pack goes through the importer (`npu-pack` still reports these containers as recognised, not importable), anything about a vision or audio tower, and any pack whose geometry has no kernel set on the box. Of the 57 `*-NPU2` repos with a readable geometry on 2026-11-11, 19 had all four From 7583389d4dc919dd8cb28b00186bf9176f20d8eb Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 08:33:48 -0400 Subject: [PATCH 6/7] npu: the symmetric tile kind needs a signed scale, and a tensor is one kind A zero offset alone decided the kind per tile, and npu_moe_container's fixture (scale 1.0, zero-point 0) was read as symmetric. A tile is now symmetric only with all zero-points zero and a scale below zero; the tensor-level readers take the kind from the whole tensor, undecided tiles follow it, and a tensor with both kinds is refused. Co-Authored-By: Claude Fable 5.2 --- docs/npu.md | 15 ++++-- npu/forward/include/q4nx_dequant.h | 43 ++++++++++++----- npu/forward/src/q4nx_dequant.c | 69 ++++++++++++++++++++++---- tests/npu_dequant_bf16_test.c | 77 ++++++++++++++++++++++++++++-- 4 files changed, 176 insertions(+), 28 deletions(-) diff --git a/docs/npu.md b/docs/npu.md index 37348df3f..baedf4d8e 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -1925,9 +1925,13 @@ The kernel directory is the `Llama-3.2-1B` set's four shapes plus that one `QKV_ Two things in these packs were read wrong before, and neither is about convolutions. Their 4-bit weights are a **symmetric** kind of the 5,120-byte tile (FastFlowLM's list calls the pack `Q4_0` where the others are `Q4_1`): the same byte layout with every zero-point stored as 0, a scale that carries a sign, and a two's-complement code. -Nothing in the container names the kind, so the decoder reads it off each tile: 256 zero-points that are all zero. -Of 1,889,236 tiles in the other packs on the box none has that block all zero; of 285,696 in the two LFM2 packs -every one has. Read as the usual kind, each LFM2 projection correlated -0.39 with the same tensor of the upstream +Nothing in the container names the kind (the name is in FastFlowLM's model list, which is not part of a pack), so +the decoder reads it from the weights: a tile is symmetric when its 256 zero-points are all zero **and** some scale +is negative, a tensor is one kind (tiles that decide nothing follow their tensor), and a tensor holding both kinds +is refused. A zero offset alone decides nothing: an asymmetric tile may have one, and the engine's own converter +test writes such a tile. Of 1,889,236 tiles in the other packs on the box none has an all-zero zero-point block or +a negative scale; of 285,696 in the two LFM2 packs every one has both. +Read as the usual kind, each LFM2 projection correlated -0.39 with the same tensor of the upstream GGUF, and the serve answered in fluent-looking fragments of unrelated words; read as symmetric, 0.99, and the pack's quantised `lm_head` sits within one step of its own bf16 embedding (mean 0.24 of a step). And the attention layers carry per-head `q_norm`/`k_norm`, which the `lfm2` model type did not switch on: with only the @@ -1955,6 +1959,11 @@ question (a Fibonacci function that returns `0` for `n <= 0` where theirs return against 42 to 47. Its logprobs at the three points were not taken, and it was asked each prompt once, so this row does not carry the three-identical-answers condition. +`LFM2.5-1.2B-Thinking-NPU2` loads on the same directory too: its template is read from the pack's +`chat_template.jinja`, the reasoning is split out, and at 96 tokens every one of the five replies was still +reasoning, so there is no answer text to compare and no row; the reasoning differs from FastFlowLM's from its +first words. + Not established by any row here: that a pack goes through the importer (`npu-pack` still reports these containers as recognised, not importable), anything about a vision or audio tower, and any pack whose geometry has no kernel set on the box. Of the 57 `*-NPU2` repos with a readable geometry on 2026-11-11, 19 had all four diff --git a/npu/forward/include/q4nx_dequant.h b/npu/forward/include/q4nx_dequant.h index ea340ad97..fd1d7a615 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/npu/forward/include/q4nx_dequant.h @@ -31,16 +31,26 @@ // and // value[row][col] = code * scale[g][row] + zero_point[g][row] // -// A tile of that size comes in two kinds, and the tile itself says which. The one above is the asymmetric kind -// (FastFlowLM lists such packs as "Q4_1"): codes 0..15, scales >= 0, and a zero-point that is the group's -// minimum. The symmetric kind (listed as "Q4_0"; the LFM2 packs) has the SAME byte layout with every zero-point -// stored as 0, a scale that carries a sign, and the code as a two's-complement nibble: +// A tile of that size comes in two kinds. The one above is the asymmetric kind (FastFlowLM lists such packs as +// "Q4_1"): codes 0..15 and a zero-point that is the group's offset. The symmetric kind (listed as "Q4_0"; the +// LFM2 packs) has the SAME byte layout with every zero-point stored as 0, a scale that carries a sign, and the +// code as a two's-complement nibble: // value[row][col] = (code >= 8 ? code - 16 : code) * scale[g][row] -// Nothing in the container names the kind, so it is read off the tile: all 256 zero-points zero means symmetric. -// Measured 2026-11-11 on every 5120-byte tile of the fifteen packs on the box: 1,889,236 asymmetric tiles, none -// with an all-zero zero-point block; 285,696 tiles in the two LFM2 packs, every one of them all-zero. Read as -// asymmetric, an LFM2 projection correlates -0.39 with the same tensor of the upstream GGUF; read as above, 0.99. -// (An asymmetric tile whose 256 group minima are all exactly 0 would be a matrix with no negative weight.) +// Nothing in the container names the kind (no metadata, no config key; the name is in FastFlowLM's model list, +// which is not part of a pack), so it is read from the weights, with three rules: +// - a tile is SYMMETRIC on its own evidence only when all 256 zero-points are zero AND some scale is below +// zero. The sign is what an asymmetric tile with no offset does not have: a zero offset with positive scales +// is a legitimate asymmetric tile (weights 0..15 * scale; tests/npu_moe_container_test.cpp writes one), and it +// reads as it always did. A symmetric scale is max/-8 of its group, so it is negative whenever the group's +// largest weight is positive: about half of them, and never none of 256 in a real tile. +// - a TENSOR is one kind. The tensor-level readers look at all of its tiles first: any symmetric tile makes +// the tensor symmetric, and its undecided tiles (all zero-points zero, no negative scale: padding, mostly) +// follow it; a tensor with both a symmetric and an asymmetric tile is refused, by name, on stderr. +// - q4nx_dequant_tile() on one tile, with no tensor around it, uses that tile's own evidence. +// Measured 2026-11-11 on every 5120-byte tile of the fifteen packs on the box: 1,889,236 tiles in the other +// packs, none with an all-zero zero-point block and none with a negative scale; 285,696 tiles in the two LFM2 +// packs, every one with both. Read as asymmetric, an LFM2 projection correlates -0.39 with the same tensor of +// the upstream GGUF; read as above, 0.99. // The same tie test that settled the asymmetric layout settles this reading: LFM2.5-1.2B's lm_head is the // quantised copy of its own bf16 token_embd, and over the first 2048 rows |dequant - embedding| / |scale| has // mean 0.243 and maximum 1.000 (0.48 % of elements above 0.5, all of them the clip of +8 steps to +7). The @@ -138,11 +148,20 @@ int q4nx_packed_grid(size_t nbytes, int rows, int cols, int* cb_out, int* rp_out int q4nx_dequant_tensor_packed(const uint8_t* data, size_t nbytes, int rows, int cols, int rows_packed, int cols_packed, int chunk_bytes, float* out); -/* 1 when the 5120-byte tile is the symmetric kind (every zero-point is +-0), 0 for the asymmetric kind. */ -int q4nx_tile_is_symmetric(const uint8_t* tile); +/* What one 5120-byte tile says about its kind (see the layout notes above). */ +#define Q4NX_TILE_ASYMMETRIC 0 /* a zero-point is present */ +#define Q4NX_TILE_SYMMETRIC 1 /* every zero-point is +-0 and some scale is below zero */ +#define Q4NX_TILE_UNDECIDED 2 /* every zero-point is +-0, no scale below zero: asymmetric unless its tensor says otherwise */ +int q4nx_tile_kind(const uint8_t* tile); -/* Dequantize one 5120-byte tile, of either kind, into out[32][ld] (row-major, strided by ld). */ +/* The kind of `ntiles` 5120-byte tiles of ONE tensor, `stride` bytes apart: Q4NX_TILE_SYMMETRIC when any tile is + * symmetric and none asymmetric, Q4NX_TILE_ASYMMETRIC when none is symmetric, -1 when both occur. */ +int q4nx_tiles_kind(const uint8_t* data, size_t ntiles, size_t stride); + +/* Dequantize one 5120-byte tile into out[32][ld] (row-major, strided by ld), on the tile's own evidence. */ void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld); +/* The same with the kind given by the caller, which knows the tensor the tile belongs to (symmetric != 0). */ +void q4nx_dequant_tile_as(const uint8_t* tile, float* out, int ld, int symmetric); /* Dequantize a whole INT4 tensor of logical shape [rows, cols] whose packed * bytes are a row-major grid of 32x256 tiles, or a PLAIN bf16 tensor of that diff --git a/npu/forward/src/q4nx_dequant.c b/npu/forward/src/q4nx_dequant.c index 21f0684a1..25acac18b 100644 --- a/npu/forward/src/q4nx_dequant.c +++ b/npu/forward/src/q4nx_dequant.c @@ -16,6 +16,7 @@ // q4nx_dequant.c — see q4nx_dequant.h for the layout and its verification. #include "q4nx_dequant.h" +#include #include static inline float bf16_to_f32(uint16_t bf) { @@ -29,19 +30,33 @@ static inline uint16_t le16(const uint8_t* p) { return (uint16_t)((uint16_t)p[0] | ((uint16_t)p[1] << 8)); } -int q4nx_tile_is_symmetric(const uint8_t* tile) { +int q4nx_tile_kind(const uint8_t* tile) { + const uint8_t* sc_b = tile; const uint8_t* zp_b = tile + Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW * 2; /* +512 */ for (int i = 0; i < Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW; i++) - if ((le16(zp_b + 2 * i) & 0x7FFF) != 0) return 0; /* +0.0 and -0.0 both count as zero */ - return 1; + if ((le16(zp_b + 2 * i) & 0x7FFF) != 0) return Q4NX_TILE_ASYMMETRIC; /* +0.0 and -0.0 both count as zero */ + for (int i = 0; i < Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW; i++) { + const uint16_t sc = le16(sc_b + 2 * i); + if ((sc & 0x8000) && (sc & 0x7FFF)) return Q4NX_TILE_SYMMETRIC; /* a scale below zero */ + } + return Q4NX_TILE_UNDECIDED; } -void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld) { +int q4nx_tiles_kind(const uint8_t* data, size_t ntiles, size_t stride) { + int sym = 0, asym = 0; + for (size_t t = 0; t < ntiles && !(sym && asym); t++) { + const int k = q4nx_tile_kind(data + t * stride); + if (k == Q4NX_TILE_SYMMETRIC) sym = 1; + else if (k == Q4NX_TILE_ASYMMETRIC) asym = 1; + } + if (sym && asym) return -1; + return sym ? Q4NX_TILE_SYMMETRIC : Q4NX_TILE_ASYMMETRIC; +} + +void q4nx_dequant_tile_as(const uint8_t* tile, float* out, int ld, int symmetric) { const uint8_t* sc_b = tile; const uint8_t* zp_b = tile + Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW * 2; /* +512 */ const uint8_t* pk = tile + Q4NX_TILE_ROWS * Q4NX_GROUPS_PER_ROW * 4; /* +1024 */ - /* The symmetric kind (see the header): signed codes, signed scales, no offset. */ - const int symmetric = q4nx_tile_is_symmetric(tile); for (int r = 0; r < Q4NX_TILE_ROWS; r++) { const int h = r / 16; @@ -66,6 +81,22 @@ void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld) { } } +void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld) { + /* One tile with nothing around it: symmetric only on its own evidence, asymmetric otherwise. */ + q4nx_dequant_tile_as(tile, out, ld, q4nx_tile_kind(tile) == Q4NX_TILE_SYMMETRIC); +} + +/* The kind of a run of 5,120-byte tiles that belong to one tensor, for the tensor-level readers below: 1 symmetric, + * 0 asymmetric, or -1 with the refusal on stderr when the tensor holds tiles of both kinds. */ +static int tensor_kind_or_refuse(const uint8_t* data, size_t ntiles, size_t stride) { + const int k = q4nx_tiles_kind(data, ntiles, stride); + if (k < 0) + fprintf(stderr, "q4nx: a tensor of %zu 5120-byte tiles holds symmetric tiles (zero-points all 0, a scale below " + "zero) and asymmetric ones (a zero-point present); one tensor is one kind, so it is refused\n", + ntiles); + return k < 0 ? -1 : (k == Q4NX_TILE_SYMMETRIC); +} + size_t q4nx_tensor_bytes(int rows, int cols) { if (rows % Q4NX_TILE_ROWS || cols % Q4NX_TILE_COLS) return 0; return (size_t)(rows / Q4NX_TILE_ROWS) * (size_t)(cols / Q4NX_TILE_COLS) * @@ -114,6 +145,11 @@ int q4nx_dequant_tensor_chunked(const uint8_t* data, size_t nbytes, int rows, return -1; const int ntr = rows / Q4NX_TILE_ROWS, ntc = cols / Q4NX_TILE_COLS; if (nbytes < (size_t)ntr * (size_t)ntc * (size_t)chunk_bytes) return -1; + int sym = 0; + if (chunk_bytes == Q4NX_TILE_BYTES) { + sym = tensor_kind_or_refuse(data, (size_t)ntr * (size_t)ntc, Q4NX_TILE_BYTES); + if (sym < 0) return -1; + } // Each tile writes a disjoint output region, so the tile grid parallelizes. // `tmp` is declared per-tile (inside the innermost body) so it is not shared. #pragma omp parallel for collapse(2) schedule(static) if(ntr * ntc >= 4) @@ -126,7 +162,7 @@ int q4nx_dequant_tensor_chunked(const uint8_t* data, size_t nbytes, int rows, else if (chunk_bytes == Q4NX_CHUNK_Q8) q4nx_dequant_tile_q8(c, tmp, Q4NX_TILE_COLS); else - q4nx_dequant_tile(c, tmp, Q4NX_TILE_COLS); + q4nx_dequant_tile_as(c, tmp, Q4NX_TILE_COLS, sym); for (int r = 0; r < Q4NX_TILE_ROWS; r++) memcpy(out + (size_t)(tr * Q4NX_TILE_ROWS + r) * cols + (size_t)tc * Q4NX_TILE_COLS, @@ -151,6 +187,12 @@ int q4nx_dequant_tile_row(const uint8_t* data, size_t nbytes, int rows, int cols if (nbytes < (size_t)ntr * (size_t)ntc * (size_t)chunk_bytes) return -1; const int tr = r / Q4NX_TILE_ROWS, rr = r % Q4NX_TILE_ROWS; float tmp[Q4NX_TILE_ROWS * Q4NX_TILE_COLS]; + /* Per token, so the kind is taken from this tile-row's own tiles (ntc of them), not from the whole table. */ + int sym = 0; + if (chunk_bytes == Q4NX_TILE_BYTES) { + sym = tensor_kind_or_refuse(data + (size_t)tr * (size_t)ntc * Q4NX_TILE_BYTES, (size_t)ntc, Q4NX_TILE_BYTES); + if (sym < 0) return -1; + } for (int tc = 0; tc < ntc; tc++) { const uint8_t* c = data + (size_t)(tr * ntc + tc) * (size_t)chunk_bytes; if (chunk_bytes == Q4NX_CHUNK_Q4K) @@ -158,7 +200,7 @@ int q4nx_dequant_tile_row(const uint8_t* data, size_t nbytes, int rows, int cols else if (chunk_bytes == Q4NX_CHUNK_Q8) q4nx_dequant_tile_q8(c, tmp, Q4NX_TILE_COLS); else - q4nx_dequant_tile(c, tmp, Q4NX_TILE_COLS); + q4nx_dequant_tile_as(c, tmp, Q4NX_TILE_COLS, sym); memcpy(out + (size_t)tc * Q4NX_TILE_COLS, tmp + (size_t)rr * Q4NX_TILE_COLS, Q4NX_TILE_COLS * sizeof(float)); } @@ -227,6 +269,11 @@ int q4nx_dequant_tensor_packed(const uint8_t* data, size_t nbytes, int rows, int const int ntr = rows / Q4NX_TILE_ROWS, ntc = cols / Q4NX_TILE_COLS; const int ntr_p = rows_packed / Q4NX_TILE_ROWS, ntc_p = cols_packed / Q4NX_TILE_COLS; if (nbytes < (size_t)ntr_p * (size_t)ntc_p * (size_t)chunk_bytes) return -1; + int sym = 0; + if (chunk_bytes == Q4NX_TILE_BYTES) { + sym = tensor_kind_or_refuse(data, (size_t)ntr_p * (size_t)ntc_p, Q4NX_TILE_BYTES); + if (sym < 0) return -1; + } // Read geometry is the container's grid (ntc_p is the tile-row stride); write geometry is the // model's, so a caller's buffer and its row stride stay exactly what they were before. #pragma omp parallel for collapse(2) schedule(static) if(ntr * ntc >= 4) @@ -239,7 +286,7 @@ int q4nx_dequant_tensor_packed(const uint8_t* data, size_t nbytes, int rows, int else if (chunk_bytes == Q4NX_CHUNK_Q8) q4nx_dequant_tile_q8(c, tmp, Q4NX_TILE_COLS); else - q4nx_dequant_tile(c, tmp, Q4NX_TILE_COLS); + q4nx_dequant_tile_as(c, tmp, Q4NX_TILE_COLS, sym); for (int r = 0; r < Q4NX_TILE_ROWS; r++) memcpy(out + (size_t)(tr * Q4NX_TILE_ROWS + r) * cols + (size_t)tc * Q4NX_TILE_COLS, @@ -278,6 +325,8 @@ int q4nx_dequant_tensor(const uint8_t* data, size_t nbytes, int rows, int cols, const int ntr = rows / Q4NX_TILE_ROWS; const int ntc = cols / Q4NX_TILE_COLS; + const int sym = tensor_kind_or_refuse(data, (size_t)ntr * (size_t)ntc, Q4NX_TILE_BYTES); + if (sym < 0) return -1; #pragma omp parallel for collapse(2) schedule(static) if(ntr * ntc >= 4) for (int tr = 0; tr < ntr; tr++) { @@ -285,7 +334,7 @@ int q4nx_dequant_tensor(const uint8_t* data, size_t nbytes, int rows, int cols, float tmp[Q4NX_TILE_ROWS * Q4NX_TILE_COLS]; const uint8_t* tile = data + (size_t)(tr * ntc + tc) * (size_t)Q4NX_TILE_BYTES; - q4nx_dequant_tile(tile, tmp, Q4NX_TILE_COLS); + q4nx_dequant_tile_as(tile, tmp, Q4NX_TILE_COLS, sym); for (int r = 0; r < Q4NX_TILE_ROWS; r++) { memcpy(out + (size_t)(tr * Q4NX_TILE_ROWS + r) * (size_t)cols + (size_t)tc * Q4NX_TILE_COLS, diff --git a/tests/npu_dequant_bf16_test.c b/tests/npu_dequant_bf16_test.c index 63621b7c8..0840e21a9 100644 --- a/tests/npu_dequant_bf16_test.c +++ b/tests/npu_dequant_bf16_test.c @@ -208,7 +208,7 @@ static void tile_kinds(void) { : (tile[1024 + (((r) / 16) * 8 + (c) / 32) * 256 + 8 * ((c) % 32) + ((r) % 16) / 2] >> 4)) #define SCALE_AT(r, c) (((((c) / 32) % 2) ? -1.0f : 1.0f) * 0.25f * (float)(1 + (r) % 4)) - check(q4nx_tile_is_symmetric(tile) == 1, "tile kinds: a zero-point block of zeros is the symmetric kind"); + check(q4nx_tile_kind(tile) == Q4NX_TILE_SYMMETRIC, "tile kinds: zero-points all zero and a scale below zero is the symmetric kind"); q4nx_dequant_tile(tile, out, Q4NX_TILE_COLS); int bad = 0, saw_neg_code = 0; for (int r = 0; r < Q4NX_TILE_ROWS; r++) @@ -223,12 +223,12 @@ static void tile_kinds(void) { /* negative zero is still zero: the kind does not flip on a sign bit */ tile[512 + 2 * 7 + 1] = 0x80; - check(q4nx_tile_is_symmetric(tile) == 1, "tile kinds: a -0.0 zero-point is still the symmetric kind"); + check(q4nx_tile_kind(tile) == Q4NX_TILE_SYMMETRIC, "tile kinds: a -0.0 zero-point is still the symmetric kind"); tile[512 + 2 * 7 + 1] = 0; /* one non-zero zero-point anywhere makes it the asymmetric kind, read as it always was */ put_bf16(tile + 512, (size_t)(3 * Q4NX_TILE_ROWS + 5), -0.5f); - check(q4nx_tile_is_symmetric(tile) == 0, "tile kinds: one non-zero zero-point makes the tile asymmetric"); + check(q4nx_tile_kind(tile) == Q4NX_TILE_ASYMMETRIC, "tile kinds: one non-zero zero-point makes the tile asymmetric"); q4nx_dequant_tile(tile, out, Q4NX_TILE_COLS); bad = 0; for (int r = 0; r < Q4NX_TILE_ROWS; r++) @@ -259,8 +259,79 @@ static void tile_kinds(void) { free(row); free(whole); free(tile); free(out); } +/* The kind is not decided from a zero offset alone. An asymmetric tile is allowed to have every zero-point 0 (weights + * that are exactly code * scale: tests/npu_moe_container_test.cpp builds its fixture that way, and a first version + * of this rule read its codes 8..15 as -8..-1). What marks the symmetric kind is the signed scale, and a tensor is + * one kind: its undecided tiles follow it, and a tensor with both kinds is refused. */ +static void tile_kind_is_per_tensor(void) { + enum { T = 4 }; /* four tiles: one tile-row of a [32][1024] tensor */ + uint8_t* t = (uint8_t*)calloc(T, Q4NX_TILE_BYTES); + float* out = (float*)malloc(sizeof(float) * Q4NX_TILE_ROWS * Q4NX_TILE_COLS * T); + float* row = (float*)malloc(sizeof(float) * Q4NX_TILE_COLS * T); + if (!t || !out || !row) { check(0, "per tensor: the buffers allocate"); free(t); free(out); free(row); return; } + const int cols = Q4NX_TILE_COLS * T; + /* every tile: scale 1.0, zero-point 0, every byte of codes 0xF9 (low nibble 9, high nibble 15) */ + for (int k = 0; k < T; k++) { + uint8_t* tile = t + (size_t)k * Q4NX_TILE_BYTES; + for (int i = 0; i < 256; i++) put_bf16(tile, (size_t)i, 1.0f); + memset(tile + 1024, 0xF9, 4096); + } + check(q4nx_tile_kind(t) == Q4NX_TILE_UNDECIDED, "per tensor: zero offsets with positive scales decide nothing"); + check(q4nx_tiles_kind(t, T, Q4NX_TILE_BYTES) == Q4NX_TILE_ASYMMETRIC, "per tensor: a tensor of undecided tiles is asymmetric"); + check(q4nx_dequant_tensor_chunked(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, cols, Q4NX_TILE_BYTES, out) == 0, + "per tensor: it decodes"); + int bad = 0; + for (int r = 0; r < Q4NX_TILE_ROWS; r++) + for (int c = 0; c < cols; c++) + if (out[r * cols + c] != ((r % 2) ? 15.0f : 9.0f)) bad++; + check(bad == 0, "per tensor: and reads codes 9 and 15 as 9 and 15, the asymmetric reading it always had"); + q4nx_dequant_tile(t, out, Q4NX_TILE_COLS); + check(out[0] == 9.0f && out[Q4NX_TILE_COLS] == 15.0f, "per tensor: the one-tile decoder agrees"); + + /* give ONE tile a negative scale: that tile is symmetric, and the other three now follow their tensor */ + put_bf16(t + 2 * (size_t)Q4NX_TILE_BYTES, 40, -1.0f); + check(q4nx_tile_kind(t + 2 * (size_t)Q4NX_TILE_BYTES) == Q4NX_TILE_SYMMETRIC, "per tensor: a negative scale makes a tile symmetric"); + check(q4nx_tiles_kind(t, T, Q4NX_TILE_BYTES) == Q4NX_TILE_SYMMETRIC, "per tensor: one symmetric tile makes the tensor symmetric"); + check(q4nx_dequant_tensor_chunked(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, cols, Q4NX_TILE_BYTES, out) == 0 && + q4nx_dequant_tensor(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, cols, out + 0) == 0, + "per tensor: both whole-tensor readers decode it"); + bad = 0; + for (int r = 0; r < Q4NX_TILE_ROWS; r++) + for (int c = 0; c < cols; c++) { + float want = (r % 2) ? -1.0f : -7.0f; /* 15 -> -1, 9 -> -7 */ + if (c / Q4NX_TILE_COLS == 2 && (c % Q4NX_TILE_COLS) / 32 == 1 && r == 8) want = -want; /* index 40 = g 1, r 8 */ + if (out[r * cols + c] != want) bad++; + } + check(bad == 0, "per tensor: every tile, the undecided ones included, reads as two's complement * scale"); + check(q4nx_dequant_tile_row(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, cols, Q4NX_TILE_BYTES, 3, row) == 0 && + row[0] == -1.0f && row[cols - 1] == -1.0f, + "per tensor: the one-row reader takes the kind from the row's tiles, undecided ones following"); + /* the padded-grid reader: the model's tensor is the first tile only, the container's grid is all four, and the + * kind is the grid's -- tile 0 is undecided on its own and symmetric as part of this tensor */ + check(q4nx_dequant_tensor_packed(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, Q4NX_TILE_COLS, Q4NX_TILE_ROWS, cols, + Q4NX_TILE_BYTES, out) == 0 && out[0] == -7.0f && out[Q4NX_TILE_COLS] == -1.0f, + "per tensor: the padded-grid reader decodes it, taking the kind from the container's grid"); + q4nx_dequant_tile(t, out, Q4NX_TILE_COLS); + check(out[0] == 9.0f, "per tensor: an undecided tile decoded alone, with no tensor around it, stays asymmetric"); + + /* now an offset in another tile: one tensor, both kinds -> refused by every tensor-level reader */ + put_bf16(t + 512, 3, 0.5f); + check(q4nx_tiles_kind(t, T, Q4NX_TILE_BYTES) == -1, "per tensor: a symmetric and an asymmetric tile in one tensor is neither kind"); + check(q4nx_dequant_tensor_chunked(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, cols, Q4NX_TILE_BYTES, out) == -1, + "per tensor: the chunked reader refuses it"); + check(q4nx_dequant_tensor(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, cols, out) == -1, + "per tensor: the whole-tensor reader refuses it"); + check(q4nx_dequant_tile_row(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, cols, Q4NX_TILE_BYTES, 0, row) == -1, + "per tensor: the one-row reader refuses it"); + check(q4nx_dequant_tensor_packed(t, (size_t)T * Q4NX_TILE_BYTES, Q4NX_TILE_ROWS, Q4NX_TILE_COLS, Q4NX_TILE_ROWS, cols, + Q4NX_TILE_BYTES, out) == -1, + "per tensor: the padded-grid reader refuses it (the kind is the container grid's, not the model's)"); + free(t); free(out); free(row); +} + int main(void) { tile_kinds(); + tile_kind_is_per_tensor(); tile_row_kind_arms(); chunked_row_arms(); embed_row_arms(); From a95c68eacc3fc56400fba04bc19c2839d3c6d5d3 Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 08:46:43 -0400 Subject: [PATCH 7/7] npu: the fused int8 packer takes the tile kind from the tensor too q4nx_pack_q4_1_tiles_to_int8() decoded each tile as a one-tile tensor, so inside a symmetric tensor a tile that decides nothing on its own was read as asymmetric there and as symmetric by the tensor readers. It now takes the kind once over all tiles, decodes with q4nx_dequant_tile_as, and refuses a tensor with both kinds, as the readers do. Co-Authored-By: Claude Fable 5.2 --- npu/forward/include/q4nx_dequant.h | 3 +++ npu/forward/src/q4nx_pack.c | 25 ++++++++++++------- tests/npu_design_grid_test.c | 39 ++++++++++++++++++++++++++++++ 3 files changed, 58 insertions(+), 9 deletions(-) diff --git a/npu/forward/include/q4nx_dequant.h b/npu/forward/include/q4nx_dequant.h index fd1d7a615..adf2736c3 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/npu/forward/include/q4nx_dequant.h @@ -46,6 +46,9 @@ // - a TENSOR is one kind. The tensor-level readers look at all of its tiles first: any symmetric tile makes // the tensor symmetric, and its undecided tiles (all zero-points zero, no negative scale: padding, mostly) // follow it; a tensor with both a symmetric and an asymmetric tile is refused, by name, on stderr. +// The fused int8 packer (q4nx_pack_q4_1_tiles_to_int8) takes the kind the same way, over the whole tensor. +// - the one-row reader (q4nx_dequant_tile_row) takes it from the tiles of the row's own tile-row: it runs once +// per token on an embedding it exists to avoid paging in, so it does not scan the table. // - q4nx_dequant_tile() on one tile, with no tensor around it, uses that tile's own evidence. // Measured 2026-11-11 on every 5120-byte tile of the fifteen packs on the box: 1,889,236 tiles in the other // packs, none with an all-zero zero-point block and none with a negative scale; 285,696 tiles in the two LFM2 diff --git a/npu/forward/src/q4nx_pack.c b/npu/forward/src/q4nx_pack.c index 1a53600b1..cda1659c8 100644 --- a/npu/forward/src/q4nx_pack.c +++ b/npu/forward/src/q4nx_pack.c @@ -689,7 +689,7 @@ float q4nx_pack_B_int8(const float* B, int K, int N, int8_t* out, int ld, * * This emits the int8 image STRAIGHT FROM THE TILES: one 32x256 tile is dequantised into a 32 KB scratch, folded into * the per-column amax, then a second pass dequantises the same tiles again and writes the codes. Same dequantiser - * (q4nx_dequant_tile), same arithmetic, same amax floor -- so the result is byte-identical to the two-step path, which + * (q4nx_dequant_tile_as, with the tensor's tile kind), same arithmetic, same amax floor -- so the result is byte-identical to the two-step path, which * is what the fixture in tests/npu_design_grid_test.c asserts by running both. * * Only the q4_1 tile kind (5120 B, 32x256) has a one-tile decoder; any other chunk kind returns -2 so the caller keeps @@ -701,18 +701,26 @@ int q4nx_pack_q4_1_tiles_to_int8(const uint8_t* tiles, size_t nbytes, int K, int if (K % Q4NX_TILE_ROWS || N % Q4NX_TILE_COLS) return -1; const size_t ntr = (size_t)(K / Q4NX_TILE_ROWS), ntc = (size_t)(N / Q4NX_TILE_COLS); if (nbytes != ntr * ntc * (size_t)Q4NX_TILE_BYTES) return -2; // not the q4_1 tile kind: a caller keeps its path + // The kind of 5,120-byte tile is the TENSOR's (q4nx_dequant.h): taken once over all of its tiles, exactly as the + // tensor readers of the two-step path take it, so a tile that decides nothing on its own is read here as it is + // read there. Decoding tile by tile as one-tile tensors read such a tile as asymmetric inside a symmetric tensor + // (8192 of 8192 elements different on a two-tile probe). A tensor holding both kinds is refused, as it is there. + const int kind = q4nx_tiles_kind(tiles, ntr * ntc, (size_t)Q4NX_TILE_BYTES); + if (kind < 0) { + fprintf(stderr, "q4nx: a tensor of %zu 5120-byte tiles holds symmetric and asymmetric tiles; one tensor is one " + "kind, so it is not packed\n", ntr * ntc); + return -1; + } + const int sym = (kind == Q4NX_TILE_SYMMETRIC); float* amax = (float*)calloc((size_t)N, sizeof(float)); if (!amax) return -1; float tilef[Q4NX_TILE_ROWS * Q4NX_TILE_COLS]; // pass 1: per-column amax, one tile at a time for (size_t r = 0; r < ntr; r++) { for (size_t c = 0; c < ntc; c++) { - // the SAME decoder the two-step reference path uses, applied to one tile: q4nx_dequant_tile() is a - // different function, and on the fixture it disagreed with the tensor path on 316 of 32768 codes -- rounding - // detail at the .5 boundary rather than layout (the scales matched exactly). Byte-exactness needs one - // decoder, not two that agree to 99%. - q4nx_dequant_tensor_chunked(tiles + (r * ntc + c) * (size_t)Q4NX_TILE_BYTES, Q4NX_TILE_BYTES, - Q4NX_TILE_ROWS, Q4NX_TILE_COLS, Q4NX_TILE_BYTES, tilef); + // the SAME one-tile decoder the tensor readers of the two-step path call (q4nx_dequant_tile_as), with + // the tensor's kind: byte-exactness needs one decoder and one decision, not two that agree to 99%. + q4nx_dequant_tile_as(tiles + (r * ntc + c) * (size_t)Q4NX_TILE_BYTES, tilef, Q4NX_TILE_COLS, sym); for (int i = 0; i < Q4NX_TILE_ROWS; i++) { const float* row = tilef + (size_t)i * Q4NX_TILE_COLS; float* am = amax + c * Q4NX_TILE_COLS; @@ -735,8 +743,7 @@ int q4nx_pack_q4_1_tiles_to_int8(const uint8_t* tiles, size_t nbytes, int K, int // pass 2: codes, same tile order for (size_t r = 0; r < ntr; r++) { for (size_t c = 0; c < ntc; c++) { - q4nx_dequant_tensor_chunked(tiles + (r * ntc + c) * (size_t)Q4NX_TILE_BYTES, Q4NX_TILE_BYTES, - Q4NX_TILE_ROWS, Q4NX_TILE_COLS, Q4NX_TILE_BYTES, tilef); + q4nx_dequant_tile_as(tiles + (r * ntc + c) * (size_t)Q4NX_TILE_BYTES, tilef, Q4NX_TILE_COLS, sym); const float* tis = amax + c * Q4NX_TILE_COLS; for (int i = 0; i < Q4NX_TILE_ROWS; i++) { const float* row = tilef + (size_t)i * Q4NX_TILE_COLS; diff --git a/tests/npu_design_grid_test.c b/tests/npu_design_grid_test.c index 5adbf52dd..880d196f2 100644 --- a/tests/npu_design_grid_test.c +++ b/tests/npu_design_grid_test.c @@ -284,6 +284,45 @@ static void oga_geometry(void) { "fused: a non-q4_1 byte count returns -2 so the caller keeps the two-step path"); printf(" fused chunks->int8: K=%d N=%d %d codes and %d scales byte-identical to dequant->pack\n", K, N, K * N, N); + + /* THE SAME COMPARISON ON A SYMMETRIC TENSOR WITH AN UNDECIDED TILE (the LFM2 packs' tile kind). Three tiles + * are symmetric on their own evidence (zero-points 0, scales of both signs); tile 1 has every zero-point 0 + * and every scale positive, which decides nothing on its own. The tensor is symmetric, so both paths must + * read tile 1 as two's-complement codes: decoded as a one-tile tensor it comes out unsigned, and the fused + * path did exactly that. */ + for (size_t t = 0; t < ntr * ntc; t++) { + uint8_t* tile = tiles + t * 5120; + for (int i = 0; i < 256; i++) { + const int neg = (t != 1) && (i % 3 == 0); + tile[2 * i] = 0x00; + tile[2 * i + 1] = neg ? 0xBE : 0x3E; /* bf16 -0.125 / +0.125 */ + tile[512 + 2 * i] = 0; tile[512 + 2 * i + 1] = 0; + } + for (size_t i = 0; i < 4096; i++) tile[1024 + i] = (uint8_t)((i * 53 + t * 29 + 0xC8) & 0xFF); + } + check_int(q4nx_tile_kind(tiles + 5120), Q4NX_TILE_UNDECIDED, "fused, symmetric: tile 1 decides nothing on its own"); + check_int(q4nx_tiles_kind(tiles, ntr * ntc, 5120), Q4NX_TILE_SYMMETRIC, "fused, symmetric: the tensor is symmetric"); + check_int(q4nx_dequant_tensor_chunked(tiles, ntr * ntc * 5120, K, N, 5120, B), 0, + "fused, symmetric: the reference path dequantises it"); + int neg_in_tile1 = 0; + for (int r = 0; r < 32; r++) + for (int c = 256; c < 512; c++) + if (B[(size_t)r * N + c] < 0) neg_in_tile1++; + check(neg_in_tile1 > 0, "fused, symmetric: the reference reads tile 1 with negative weights (signed codes, positive scales)"); + q4nx_pack_B_int8(B, K, N, ref, ld, sc_ref); + check_int(q4nx_pack_q4_1_tiles_to_int8(tiles, ntr * ntc * 5120, K, N, fused, ld, sc_fused), 0, + "fused, symmetric: the fused path accepts it"); + bad_code = 0; bad_scale = 0; + for (int i = 0; i < K * N; i++) if (fused[i] != ref[i]) bad_code++; + for (int j = 0; j < N; j++) if (sc_fused[j] != sc_ref[j]) bad_scale++; + check_int(bad_code, 0, "fused, symmetric: every code is byte-identical to dequant -> pack, the undecided tile's included"); + check_int(bad_scale, 0, "fused, symmetric: every column scale as well"); + /* one zero-point in tile 3: a tensor with both kinds, refused by both paths */ + tiles[3 * 5120 + 512 + 1] = 0x3F; + check_int(q4nx_dequant_tensor_chunked(tiles, ntr * ntc * 5120, K, N, 5120, B), -1, + "fused, mixed: the reference path refuses a tensor with both kinds"); + check_int(q4nx_pack_q4_1_tiles_to_int8(tiles, ntr * ntc * 5120, K, N, fused, ld, sc_fused), -1, + "fused, mixed: and so does the fused path"); } free(tiles); free(B); free(ref); free(fused); free(sc_ref); free(sc_fused); }