diff --git a/docs/npu.md b/docs/npu.md index c17e721f8..baedf4d8e 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` | 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` 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 | @@ -1900,6 +1900,8 @@ 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 | +| `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 @@ -1912,18 +1914,61 @@ 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 (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 +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. + +`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. + +`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 -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/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_dequant.h b/npu/forward/include/q4nx_dequant.h index 6b2968b44..adf2736c3 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/npu/forward/include/q4nx_dequant.h @@ -31,6 +31,34 @@ // and // value[row][col] = code * scale[g][row] + zero_point[g][row] // +// 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 (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. +// 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 +// 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 +// 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] // groups with signed nibbles). That layout drives the reconstruction to NaN @@ -123,8 +151,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); -/* Dequantize one 5120-byte tile into out[32][ld] (row-major, strided by ld). */ +/* 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); + +/* 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/include/q4nx_forward.h b/npu/forward/include/q4nx_forward.h index 9618e07f4..1dbed49dd 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_dequant.c b/npu/forward/src/q4nx_dequant.c index 98b95d212..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,7 +30,30 @@ static inline uint16_t le16(const uint8_t* p) { return (uint16_t)((uint16_t)p[0] | ((uint16_t)p[1] << 8)); } -void q4nx_dequant_tile(const uint8_t* tile, float* out, int ld) { +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 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; +} + +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 */ @@ -42,16 +66,37 @@ 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; + } } } } +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) * @@ -100,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) @@ -112,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, @@ -137,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) @@ -144,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)); } @@ -213,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) @@ -225,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, @@ -264,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++) { @@ -271,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/npu/forward/src/q4nx_forward.cpp b/npu/forward/src/q4nx_forward.cpp index 37696df84..9dca11690 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) @@ -850,20 +851,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_); } } @@ -1164,6 +1203,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 @@ -1283,6 +1323,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); @@ -1980,7 +2041,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) { @@ -2376,6 +2442,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 @@ -2414,6 +2481,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(); @@ -2445,6 +2513,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; } @@ -2895,6 +2964,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) // -------------------------------------------------------------------------- @@ -3705,7 +3875,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]; @@ -4960,7 +5132,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..cda1659c8 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; @@ -684,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 @@ -696,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; @@ -730,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/npu/forward/src/q4nx_semantics.c b/npu/forward/src/q4nx_semantics.c index e5ee8f81d..b116a1950 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 @@ -463,6 +466,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_dequant_bf16_test.c b/tests/npu_dequant_bf16_test.c index 2ab1003e8..0840e21a9 100644 --- a/tests/npu_dequant_bf16_test.c +++ b/tests/npu_dequant_bf16_test.c @@ -192,7 +192,146 @@ 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_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++) + 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_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_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++) + 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); +} + +/* 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(); diff --git a/tests/npu_design_grid_test.c b/tests/npu_design_grid_test.c index c39539a75..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); } @@ -831,7 +870,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(); 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");