npu: Gemma4-12B-IT-NPU2 loads and answers: the 16x256 wide chunks, attention_k_eq_v, no PLE (Refs #1786)

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\n<channel|>Paris<turn|>`: 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

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

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. */

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;

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;
     }

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;

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();

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");