From 16b9d8ae8b4a93f0870eae229bea5304ce50280b Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Thu, 12 Nov 2026 04:11:09 -0400 Subject: [PATCH] npu: Gemma4-12B-IT-NPU2 loads and answers: the 16x256 wide chunks, attention_k_eq_v, no PLE (Refs #1786) Three things stood between FastFlowLM's dense Gemma4 and an answer: - init failed with an EMPTY error: load_norms read post_layernorm (the per-layer-embedding block's norm) for every Gemma4, and this one has no PLE. Read only when ple_dim_ > 0; the lookups now name what is missing. - its I8 tensors are declared [rows/16, K/256, bytes], 2304 (Q4W) or 4352 (Q8W, the head): 16 x 256 chunks, 128 bf16 scales at g*16 + r (symmetric), then 4-bit two's-complement codes at c*8 + r/2 or int8 at c*16 + r. A Q4W tensor has a 4608-byte Q4S grid's byte count and was read as one (non-finite at layer 0); the declared shape now routes it. Settled against google's QAT Q4_0 GGUF (corr 0.9986, 7 tensors, layers 0/5/47) and the pack's own bf16 embedding (Q8W head, 1.00000). - attention_k_eq_v: global layers have no v_proj; v is the k projection (our llama.cpp fork's gemma4.cpp: Vcur = wv ? wv x : Kcur), then Gemma4's scale-free v norm. A missing v_proj without the key is refused. Host: '<|channel>thought\nParis' (the template's empty-thought prefix, then the answer); Gemma4-E2B (PLE) still 'Paris'. Tests: wide_chunks in npu_dequant_bf16_test (both kinds, refusals, the Q4S byte-count collision), Gemma4-12B keys in q4nx_gemma3_semantics_test. ctest 268/268. Co-Authored-By: Claude Opus 5.5 --- docs/npu.md | 22 ++++++++++++ npu/forward/include/q4nx_dequant.h | 15 ++++++++ npu/forward/include/q4nx_semantics.h | 4 +++ npu/forward/src/q4nx_dequant.c | 35 ++++++++++++++++++ npu/forward/src/q4nx_forward.cpp | 32 ++++++++++++++--- npu/forward/src/q4nx_semantics.c | 2 ++ tests/npu_dequant_bf16_test.c | 53 ++++++++++++++++++++++++++++ tests/q4nx_gemma3_semantics_test.c | 11 ++++++ 8 files changed, 170 insertions(+), 4 deletions(-) diff --git a/docs/npu.md b/docs/npu.md index 2b675926b..d8c874d26 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -2052,6 +2052,28 @@ RoPE. `Qwen2.5-VL-3B-Instruct-NPU2`'s text tower now answers `Paris<|im_end|>` o loaded here. `batch_check_state_qwen2vl` is the Qwen2 above declaring that block: the same logits as the plain-RoPE reference (corr 1.0000000). With `"mrope"` refused again (main's behaviour) the arm fails. +**Gemma4-12B-IT-NPU2: a dense Gemma4 and two new chunk kinds (engine#1786).** The text tower loads and answers on the +host after three changes: + +- **The wide chunks.** Every I8 tensor is declared `[rows/16, K/256, bytes]`, with `bytes` either 2304 (`Q4NX_CHUNK_Q4W`) + or 4352 (`Q4NX_CHUNK_Q8W`, the tied head). A chunk is 16 rows x 256 columns. It starts with 128 bf16 scales at + `g*16 + r` (one per row and 32-column group, symmetric, no zero-point). Then come either 4-bit two's-complement + codes at `c*8 + r/2` (even row in the low nibble) or int8 codes at `c*16 + r`. The reading was settled against + independent weights: Google's own QAT Q4_0 GGUF of the model agrees at corr 0.9986 on seven tensors over layers 0, + 5 and 47 (the 5% rel-rms is FastFlowLM's re-quantization, the same in each), and the Q8W head reproduces the + pack's bf16 embedding at 1.00000. A Q4W tensor has exactly the byte count of a 4608-byte Q4S grid, which is how it + was read before. Layer 0 then went non-finite at its first attention. The declared shape now picks the kind. +- **`attention_k_eq_v`.** The global layers have no `v_proj`: their values are the k projection, taken before + `k_norm` and then given Gemma4's scale-free v norm. This is our llama.cpp fork's rule (`gemma4.cpp`: + `Vcur = wv ? wv x : Kcur`). The forward fills those v rows from k. A layer with no `v_proj` and no such key is + refused by name. +- **No per-layer embedding.** `post_layernorm` belongs to the PLE block, so it is read only when the model has one. + Before, init failed on its absence with an empty error. Those lookups now name the missing tensor. + +On the host the known-answer prompt gives `<|channel>thought\nParis`: the model writes the +template's own empty-thought prefix, then answers. Gemma4-E2B, the PLE model, still answers `Paris`. No kernel set +for the 12B's shapes was built or run (the NPU is off). + **Per-head q/k normalisation is read from the container (engine#1788).** Whether the forward normalises q and k per head used to be decided by `config.json`'s `model_type` alone: on for `qwen3*` and `gemma*`. A pack of the same architecture under another name got it wrong in silence. FastFlowLM's MiniCPM-V 4.7 diff --git a/npu/forward/include/q4nx_dequant.h b/npu/forward/include/q4nx_dequant.h index 35b7d5a04..165e99dd2 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/npu/forward/include/q4nx_dequant.h @@ -122,6 +122,21 @@ void q4nx_dequant_chunk_q4h(const uint8_t* chunk, float* out, int ld, int symmet * the tensor mixes the symmetric and asymmetric kinds. */ int q4nx_dequant_tensor_q4h(const uint8_t* data, size_t nbytes, int rows, int k, int cols, float* out); +/* The 16 x 256 "wide" chunks of FastFlowLM's Gemma4-12B-IT-NPU2 (engine#1786), declared [rows/16, K/256, bytes]: + * bytes 0..255 128 bf16 scales, slot g*16 + r (g the 32-column group, r the row), symmetric (no zero-point) + * Q4W (2304 B) then 2048 B of 4-bit codes at byte c*8 + r/2, even row in the low nibble, two's complement + * Q8W (4352 B) then 4096 B of int8 codes at byte c*16 + r + * w[r][c] = code * scale[c/32][r]. The Q4W reading matches google's own QAT Q4_0 GGUF of the model at corr 0.9986 + * (seven tensors over layers 0, 5 and 47); the Q8W head matches the pack's bf16 embedding at 1.00000. A Q4W + * tensor's byte count equals a 4608-byte Q4S grid's, which is why the declared shape, not the size, picks it. */ +#define Q4NX_CHUNK_Q4W 2304 +#define Q4NX_CHUNK_Q8W 4352 +#define Q4NX_WIDE_ROWS 16 +#define Q4NX_WIDE_COLS 256 +void q4nx_dequant_chunk_wide(const uint8_t* chunk, int bits, float* out, int ld); +/* Whole tensor: ntr x ntc chunks, row-major, into a rows x cols buffer (cols >= k; the pad columns are zeroed). */ +int q4nx_dequant_tensor_wide(const uint8_t* data, size_t nbytes, int chunk_bytes, int rows, int k, int cols, float* out); + /* Q4_K chunk: 32 rows x 256 K, 4736 B. * scales[8][32] uint8 @[0,256) index g*32 + r * mins [8][32] uint8 @[256,512) index g*32 + r diff --git a/npu/forward/include/q4nx_semantics.h b/npu/forward/include/q4nx_semantics.h index bfc7d3e32..7511f754b 100644 --- a/npu/forward/include/q4nx_semantics.h +++ b/npu/forward/include/q4nx_semantics.h @@ -190,6 +190,10 @@ typedef struct { * both. Sources: global_head_dim, num_global_key_value_heads. */ int global_head_dim; int num_global_kv_heads; + /* Gemma4 attention_k_eq_v (Gemma4-12B-IT): a layer with no v_proj takes its values from the k projection, before + * k_norm (our llama.cpp fork's gemma4.cpp: `Vcur = wv ? wv x : Kcur`, then v's scale-free RMSNorm). Source: + * attention_k_eq_v. */ + int attention_k_eq_v; /* Gemma4 sets attention scaling to 1.0 (q/k carry a per-head RMSNorm), * normalises v with a SCALE-FREE RMSNorm, and uses "proportional" RoPE on * its full layers. Sources: model_type/architectures, rope_type. */ diff --git a/npu/forward/src/q4nx_dequant.c b/npu/forward/src/q4nx_dequant.c index ecb2438db..be8657634 100644 --- a/npu/forward/src/q4nx_dequant.c +++ b/npu/forward/src/q4nx_dequant.c @@ -197,6 +197,41 @@ int q4nx_dequant_tensor_q4h(const uint8_t* data, size_t nbytes, int rows, int k, return 0; } +void q4nx_dequant_chunk_wide(const uint8_t* chunk, int bits, float* out, int ld) { + float sc[Q4NX_WIDE_ROWS * 8]; + for (int i = 0; i < Q4NX_WIDE_ROWS * 8; i++) sc[i] = bf16_to_f32(le16(chunk + 2 * i)); + const uint8_t* codes = chunk + 256; + for (int r = 0; r < Q4NX_WIDE_ROWS; r++) + for (int c = 0; c < Q4NX_WIDE_COLS; c++) { + int q; + if (bits == 4) { + const uint8_t b = codes[c * 8 + r / 2]; + const int v = (r & 1) ? (b >> 4) : (b & 15); + q = v >= 8 ? v - 16 : v; + } else { + q = (int8_t)codes[c * 16 + r]; + } + out[(size_t)r * ld + c] = (float)q * sc[(c / 32) * Q4NX_WIDE_ROWS + r]; + } +} + +int q4nx_dequant_tensor_wide(const uint8_t* data, size_t nbytes, int chunk_bytes, int rows, int k, int cols, float* out) { + if (!data || !out || rows <= 0 || k <= 0 || cols < k) return -1; + const int bits = chunk_bytes == Q4NX_CHUNK_Q4W ? 4 : (chunk_bytes == Q4NX_CHUNK_Q8W ? 8 : 0); + if (!bits || rows % Q4NX_WIDE_ROWS || k % Q4NX_WIDE_COLS) return -1; + const int ntr = rows / Q4NX_WIDE_ROWS, ntc = k / Q4NX_WIDE_COLS; + if (nbytes != (size_t)ntr * (size_t)ntc * (size_t)chunk_bytes) return -1; + #pragma omp parallel for collapse(2) schedule(static) if((size_t)ntr * ntc >= 4) + for (int tr = 0; tr < ntr; tr++) + for (int tc = 0; tc < ntc; tc++) + q4nx_dequant_chunk_wide(data + ((size_t)tr * ntc + tc) * (size_t)chunk_bytes, bits, + out + (size_t)tr * Q4NX_WIDE_ROWS * cols + (size_t)tc * Q4NX_WIDE_COLS, cols); + if (cols > k) + for (int r = 0; r < rows; r++) + memset(out + (size_t)r * cols + k, 0, (size_t)(cols - k) * sizeof(float)); + return 0; +} + int q4nx_dequant_tensor_chunked(const uint8_t* data, size_t nbytes, int rows, int cols, int chunk_bytes, float* out) { if (!data || !out || rows <= 0 || cols <= 0) return -1; diff --git a/npu/forward/src/q4nx_forward.cpp b/npu/forward/src/q4nx_forward.cpp index 55ebf170a..f04b29699 100644 --- a/npu/forward/src/q4nx_forward.cpp +++ b/npu/forward/src/q4nx_forward.cpp @@ -188,7 +188,10 @@ static bool valid_sidecar_dir(const std::string& d) { // The row width a tensor declares, in bytes: 1280 names the quarter-tile kind (Q4NX_CHUNK_Q4H), which a byte count // cannot tell from the 5120-byte tile. 0 for a tensor that is not two-dimensional. +// A three-dimensional [rows/16, K/256, bytes] declares its chunk in the last dimension: 2304 and 4352 are the wide kinds +// (Q4NX_CHUNK_Q4W / Q8W), and a Q4W tensor's byte count would otherwise pass for a 4608-byte Q4S grid. static int64_t declared_row_bytes(const TensorDesc* d) { + if (d && d->ndim == 3 && (d->shape[2] == Q4NX_CHUNK_Q4W || d->shape[2] == Q4NX_CHUNK_Q8W)) return d->shape[2]; return (d && d->ndim == 2) ? d->shape[1] : 0; } @@ -202,6 +205,13 @@ static int dequant_any(const void* data, size_t nbytes, int rows, int cols, floa return q4nx_dequant_tensor_q4h((const uint8_t*)data, nbytes, rows, (int)(chunks / row_tiles) * Q4NX_Q4H_COLS, cols, out); } + if (decl_row_bytes == Q4NX_CHUNK_Q4W || decl_row_bytes == Q4NX_CHUNK_Q8W) { + if (rows % Q4NX_WIDE_ROWS) return -1; + const size_t chunks = nbytes / (size_t)decl_row_bytes, row_tiles = (size_t)(rows / Q4NX_WIDE_ROWS); + if (chunks % row_tiles) return -1; + return q4nx_dequant_tensor_wide((const uint8_t*)data, nbytes, (int)decl_row_bytes, rows, + (int)(chunks / row_tiles) * Q4NX_WIDE_COLS, cols, out); + } if (rows % Q4NX_TILE_ROWS || cols % Q4NX_TILE_COLS) return -1; const size_t n = (size_t)(rows / Q4NX_TILE_ROWS) * (size_t)(cols / Q4NX_TILE_COLS); if (n == 0) return -1; @@ -1941,9 +1951,10 @@ bool Q4nxNpuForward::load_norms() { auto rd = [&](const char* fmt, float* dst, int n) -> bool { snprintf(nm, sizeof(nm), fmt, l); TensorDesc* td = model_tensor_by_name(mw_, nm); - if (!td || n < 0) return false; + // named, not silent: init_host used to fail here with an empty error (Gemma4-12B-IT, engine#1786) + if (!td || n < 0) { err_ = std::string("model.q4nx has no ") + nm; return false; } const uint16_t* p = bf16_data(mw_, td, (size_t)n); - if (!p) return false; + if (!p) { err_ = std::string(nm) + " is not bf16 of " + std::to_string(n) + " values"; return false; } for (int i = 0; i < n; i++) dst[i] = bf16_to_f32(p[i]); return true; }; @@ -1972,7 +1983,10 @@ bool Q4nxNpuForward::load_norms() { // Gemma3 (extra_layer_norms 2) has the two feedforward norms only: no post-PLE norm and no // layer scalar, so post_ln_ stays unused (ple_dim_ is 0) and layer_scale_ stays 1. if (sem_.extra_layer_norms == 1) { - if (!rd("model.layers.%d.post_layernorm.weight", post_ln_[l].data(), H)) return false; + // post_layernorm is the per-layer-embedding block's norm (post_ln_ is read only there): a dense + // Gemma4 (Gemma4-12B-IT, hidden_size_per_layer_input 0) has no PLE and no such tensor. + if (ple_dim_ > 0 && + !rd("model.layers.%d.post_layernorm.weight", post_ln_[l].data(), H)) return false; if (!rd("model.layers.%d.layer_output_scale.weight", &layer_scale_[l], 1)) return false; } } @@ -2205,7 +2219,17 @@ bool Q4nxNpuForward::load_layer(int l) { NH * hd * (attn_gate_ ? 2 : 1), kpad_)) return false; // q|gate rows if (!kv_shared) { if (!deq(&lw->k_proj_weight, wk_.data(), nkv * hd, kpad_)) return false; - if (!deq(&lw->v_proj_weight, wv_.data(), nkv * hd, kpad_)) return false; + const bool v_absent = lw->v_proj_weight.name[0] == '\0' || lw->v_proj_weight.data_size == 0; + if (v_absent && sem_.attention_k_eq_v) { + // attention_k_eq_v: v is the k projection (normalised scale-free in attention(), as Gemma4's v is) + std::copy(wk_.begin(), wk_.begin() + (ptrdiff_t)((size_t)nkv * hd * kpad_), wv_.begin()); + } else if (v_absent) { + err_ = "model.layers." + std::to_string(l) + " has no self_attn.v_proj.weight and config.json does not " + "declare attention_k_eq_v"; + return false; + } else if (!deq(&lw->v_proj_weight, wv_.data(), nkv * hd, kpad_)) { + return false; + } } if (!deq(&lw->o_proj_weight, wo_.data(), H, q4nx_pad_k(NH * hd))) return false; } diff --git a/npu/forward/src/q4nx_semantics.c b/npu/forward/src/q4nx_semantics.c index 681fcc696..d833aad8c 100644 --- a/npu/forward/src/q4nx_semantics.c +++ b/npu/forward/src/q4nx_semantics.c @@ -402,6 +402,8 @@ int q4nx_semantics_from_config(const char* json_path, const char* model_dir, * no num_global_key_value_heads the head count is the same for both types. */ if (get_num(json, "global_head_dim", &d) && d > 0) s->global_head_dim = (int)d; + s->attention_k_eq_v = (strstr(json, "\"attention_k_eq_v\": true") != NULL || + strstr(json, "\"attention_k_eq_v\":true") != NULL) ? 1 : 0; if (get_num(json, "num_global_key_value_heads", &d) && d > 0) s->num_global_kv_heads = (int)d; const double full_hd = s->global_head_dim > 0 ? (double)s->global_head_dim : hd; diff --git a/tests/npu_dequant_bf16_test.c b/tests/npu_dequant_bf16_test.c index 046643026..1c6f9c967 100644 --- a/tests/npu_dequant_bf16_test.c +++ b/tests/npu_dequant_bf16_test.c @@ -419,9 +419,62 @@ static void q4h_chunks(void) { free(out); } +/* engine#1786: Gemma4-12B-IT-NPU2's 16 x 256 wide chunks (q4nx_dequant.h). Settled against google's QAT Q4_0 GGUF + * of the model (corr 0.9986) and the pack's own bf16 embedding (the Q8W head, 1.00000); this pins the reading. */ +static void wide_chunks(void) { + const int rows = 32, k = 512, cols = 768, ntc = k / Q4NX_WIDE_COLS; + const size_t n = (size_t)(rows / Q4NX_WIDE_ROWS) * (size_t)ntc; + for (int bits = 4; bits <= 8; bits += 4) { + const int cb = bits == 4 ? Q4NX_CHUNK_Q4W : Q4NX_CHUNK_Q8W, ncode = cb - 256; + uint8_t* t = (uint8_t*)calloc(n, (size_t)cb); + float* out = (float*)malloc(sizeof(float) * (size_t)rows * cols); + if (!t || !out) { check(0, "wide: the buffers allocate"); free(t); free(out); return; } + for (size_t ci = 0; ci < n; ci++) { + uint8_t* c = t + ci * (size_t)cb; + for (int i = 0; i < 128; i++) + put_bf16(c, (size_t)i, (i % 3 ? 0.125f : -0.25f) * (float)(1 + (int)((i + ci) % 5))); /* signed scale */ + for (int b = 0; b < ncode; b++) c[256 + b] = (uint8_t)((b * 13 + 5 * ci) & 0xff); + } + for (size_t i = 0; i < (size_t)rows * cols; i++) out[i] = 99.0f; + check(q4nx_dequant_tensor_wide(t, n * (size_t)cb, cb, rows, k, cols, out) == 0, + bits == 4 ? "q4w: a 2 x 2 grid reads" : "q8w: a 2 x 2 grid reads"); + int bad = 0, pad_bad = 0; + for (int r = 0; r < rows; r++) + for (int col = 0; col < cols; col++) { + const float got = out[(size_t)r * cols + col]; + if (col >= k) { if (got != 0.0f) pad_bad++; continue; } + const int tr = r / 16, tc = col / 256, rr = r % 16, cc = col % 256; + const size_t ci = (size_t)tr * ntc + tc; + const int i = (cc / 32) * 16 + rr; + const float sc = (i % 3 ? 0.125f : -0.25f) * (float)(1 + (int)((i + ci) % 5)); + int code; + if (bits == 4) { + const uint8_t byte = (uint8_t)(((cc * 8 + rr / 2) * 13 + 5 * ci) & 0xff); + const int v = (rr % 2 == 0) ? (byte & 0xF) : (byte >> 4); + code = v >= 8 ? v - 16 : v; + } else { + code = (int8_t)(uint8_t)(((cc * 16 + rr) * 13 + 5 * ci) & 0xff); + } + if (got != (float)code * sc) bad++; + } + check(bad == 0, bits == 4 ? "q4w: code (two's complement, c*8 + r/2, even row low) * scale (g*16 + r), row-major grid" + : "q8w: int8 code (c*16 + r) * scale (g*16 + r), row-major grid"); + check(pad_bad == 0, "wide: the columns past the container's K are written 0"); + check(q4nx_dequant_tensor_wide(t, n * (size_t)cb - 1, cb, rows, k, cols, out) == -1, "wide: one byte short is refused"); + check(q4nx_dequant_tensor_wide(t, n * (size_t)cb, cb, 24, k, cols, out) == -1, "wide: rows off the 16-row grid are refused"); + check(q4nx_dequant_tensor_wide(t, n * (size_t)cb, 4608, rows, k, cols, out) == -1, "wide: another chunk size is not a wide kind"); + free(t); + free(out); + } + /* why the declared shape decides: a Q4W tensor is exactly as many bytes as a Q4S (4608) grid of the same size */ + check((size_t)(4096 / 16) * (3840 / 256) * Q4NX_CHUNK_Q4W == (size_t)(4096 / 32) * (3840 / 256) * 4608, + "a 4096 x 3840 Q4W tensor has a 4608-byte Q4S grid's byte count, so only the declared shape tells them apart"); +} + int main(void) { q4s_chunk(); q4h_chunks(); + wide_chunks(); tile_kinds(); tile_kind_is_per_tensor(); tile_row_kind_arms(); diff --git a/tests/q4nx_gemma3_semantics_test.c b/tests/q4nx_gemma3_semantics_test.c index 7eafddcd0..96a47cb9e 100644 --- a/tests/q4nx_gemma3_semantics_test.c +++ b/tests/q4nx_gemma3_semantics_test.c @@ -94,6 +94,17 @@ int main(void) { check(s.extra_layer_norms == 2 && s.rmsnorm_offset == 0, "gemma3_text_only is gemma3 too"); } + /* Gemma4-12B-IT-NPU2 (engine#1786): a dense Gemma4. Its global layers have no v_proj, and attention_k_eq_v says + * their values are the k projection; no per-layer embedding, so no post_layernorm to read. */ + if (sem_of(dir, "{\"model_type\": \"gemma4_unified\", \"num_hidden_layers\": 6, \"head_dim\": 256," + " \"global_head_dim\": 512, \"num_global_key_value_heads\": 1, \"attention_k_eq_v\": true," + " \"hidden_size_per_layer_input\": 0, \"sliding_window\": 1024}", &s) == 0) + check(s.attention_k_eq_v == 1 && s.extra_layer_norms == 1 && s.hidden_size_per_layer_input == 0 && + s.global_head_dim == 512 && s.num_global_kv_heads == 1, + "Gemma4-12B: attention_k_eq_v read, Gemma4 norms, no PLE, global head 512 x 1 kv head"); + if (sem_of(dir, "{\"model_type\": \"gemma4_text\", \"num_hidden_layers\": 2}", &s) == 0) + check(s.attention_k_eq_v == 0, "a Gemma4 config without the key keeps its own v_proj"); + /* The neighbours keep what they had. */ if (sem_of(dir, "{\"model_type\": \"gemma2\", \"hidden_size\": 2304}", &s) == 0) check(s.rmsnorm_offset == 1 && s.extra_layer_norms == 0, "gemma2 keeps (1 + w) and is not gemma3");