npu: read the 1,280-byte quarter tile, so Gemma3-270M-NPU2 loads (Refs #1786; stacked on #2242)

CMakeLists.txt

@@ -1732,6 +1732,11 @@ add_executable(q4nx_qk_norm_test tests/q4nx_qk_norm_test.c npu/forward/src/q4nx_
 target_include_directories(q4nx_qk_norm_test PRIVATE npu/forward/include)
 target_link_libraries(q4nx_qk_norm_test PRIVATE m)
 add_test(NAME q4nx_qk_norm COMMAND q4nx_qk_norm_test)
+add_executable(q4nx_gemma3_semantics_test tests/q4nx_gemma3_semantics_test.c npu/forward/src/q4nx_semantics.c)
+target_include_directories(q4nx_gemma3_semantics_test PRIVATE npu/forward/include)
+target_link_libraries(q4nx_gemma3_semantics_test PRIVATE m)
+add_test(NAME q4nx_gemma3_semantics COMMAND q4nx_gemma3_semantics_test)
+add_dependencies(npu_sampling_test q4nx_gemma3_semantics_test)
 add_dependencies(npu_sampling_test q4nx_qk_norm_test)
 add_executable(q4nx_layer_kinds_test tests/q4nx_layer_kinds_test.c npu/forward/src/q4nx_semantics.c)
 target_include_directories(q4nx_layer_kinds_test PRIVATE npu/forward/include)

app/forward_serve.cpp

@@ -1115,6 +1115,22 @@ int run_forward_serve(int argc, char** argv) {
     e.tok = std::make_unique<npu::Tokenizer>(tok_json);
     e.bos_id  = json_int(cfg, "bos_token_id", -1);
     e.add_bos = json_int(cfg, "add_bos_token", 0) != 0;
+    // config.json is the first source; what it does not state, tokenizer_config.json does (the BOS text, looked up
+    // in the tokenizer, and add_bos_token). Without this a Gemma3 pack was served with no BOS at all.
+    if (e.bos_id < 0 || json_int(cfg, "add_bos_token", -1) < 0) {
+        const nlohmann::json tc = npu::read_json_or_null(dir + "/tokenizer_config.json");
+        if (e.bos_id < 0) {
+            const std::string b = npu::special_token_text(tc, "bos_token");
+            if (!b.empty()) e.bos_id = e.tok->token_id(b);
+        }
+        if (json_int(cfg, "add_bos_token", -1) < 0) {
+            const int a = npu::add_bos_token_flag(tc);
+            if (a >= 0) e.add_bos = a != 0;
+        }
+        if (e.bos_id >= 0 && e.add_bos)
+            std::fprintf(stderr, "1bit forward-serve: BOS %d from tokenizer_config.json (config.json does not state it)\n",
+                         e.bos_id);
+    }
     e.eos = npu::eos_token_ids_from_config(cfg);
     const size_t eos_from_config = e.eos.size();
     e.endoftext = e.tok->token_id("<|endoftext|>");

app/npu_stop_tokens.h

@@ -58,14 +58,31 @@ inline std::vector<int> eos_token_ids(const nlohmann::json& cfg) {
 
 /// tokenizer_config.json's `eos_token` as text: a plain string, or the AddedToken object form
 /// `{"content": "<|im_end|>", ...}`.  Empty when the key is absent, null or another shape.
-inline std::string eos_token_text(const nlohmann::json& tokenizer_config) {
-    if (!tokenizer_config.is_object() || !tokenizer_config.contains("eos_token")) return {};
-    const nlohmann::json& e = tokenizer_config["eos_token"];
+/// A special token's text from tokenizer_config.json, in either spelling HF writes ("<bos>" or {"content": "<bos>"}).
+inline std::string special_token_text(const nlohmann::json& tokenizer_config, const char* key) {
+    if (!tokenizer_config.is_object() || !tokenizer_config.contains(key)) return {};
+    const nlohmann::json& e = tokenizer_config[key];
     if (e.is_string()) return e.get<std::string>();
     if (e.is_object() && e.contains("content") && e["content"].is_string()) return e["content"].get<std::string>();
     return {};
 }
 
+inline std::string eos_token_text(const nlohmann::json& tokenizer_config) {
+    return special_token_text(tokenizer_config, "eos_token");
+}
+
+/// tokenizer_config.json's add_bos_token: 1 or 0 when it states a boolean (or integer), -1 when it says nothing.
+/// A container's config.json often omits bos_token_id and add_bos_token (FastFlowLM's Medgemma-4B-NPU2 has neither)
+/// while its tokenizer_config.json states both; HF's tokenizer then adds the BOS, and Gemma without it answers
+/// "The sky is blue because of the way the light is" where it should explain Rayleigh scattering (engine#1786).
+inline int add_bos_token_flag(const nlohmann::json& tokenizer_config) {
+    if (!tokenizer_config.is_object() || !tokenizer_config.contains("add_bos_token")) return -1;
+    const nlohmann::json& a = tokenizer_config["add_bos_token"];
+    if (a.is_boolean()) return a.get<bool>() ? 1 : 0;
+    if (a.is_number_integer()) return a.get<int>() != 0 ? 1 : 0;
+    return -1;
+}
+
 /// The marker a chat template closes an assistant turn with, read off the template's own output.  `after` is the
 /// text a template rendered immediately after an assistant message's content; the marker is the `<...>` token it
 /// starts with (leading whitespace skipped), or empty when it starts with anything else.

docs/npu.md

@@ -1949,6 +1949,55 @@ the missing q/k norm above was found:
   from one teacher-forced pass. Compared with a reference's per-row choices, a low agreement that is flat
   across positions is an error in the layers; a step at some row is an error at a position or a batch edge.
 
+**Gemma3 text decoders: the forward reads them (engine#1786).** FastFlowLM's `Medgemma-4B-NPU2` (`gemma3_text`, 34
+layers) now answers correctly through the per-op forward on the host (`NPU_INFER_HOST_GEMM=1`, no device). Before this
+change it was refused at load: `config.json declares rope scaling "linear", which this forward does not implement`.
+Five things were missing. All are read from the config or decided by the family:
+
+- `sliding_window_pattern`: the layer mix, when there is no `layer_types` list. Every 6th layer is full attention
+  (HF's `(i + 1) % 6`), so Medgemma has 5 full and 29 sliding layers. Before, all 34 ran as full attention.
+- `rope_local_base_freq`: the sliding layers' own RoPE base (10000, against 1e6 for the full layers). A sliding layer
+  with no base of its own now takes the model's base. Before, it took 0, which was the case for Gemma3-270M's
+  stated `layer_types`.
+- `rope_scaling` `{"rope_type": "linear", "factor": 8}`: the frequency is divided by the factor, on the full layers
+  only, because Gemma3's sliding layers rotate unscaled. A model without a local base scales every layer, as HF does.
+  A `linear` block with no usable factor is still refused.
+- `query_pre_attn_scalar`: the attention scale is `1/sqrt(query_pre_attn_scalar)`. It equals `head_dim` on Medgemma
+  but not on every Gemma3.
+- The sandwich norms: `post_attention_layernorm` on the attention output, plus `pre_feedforward_layernorm` and
+  `post_feedforward_layernorm` (`extra_layer_norms = 2`: Gemma4's set without its post-PLE norm and layer scalar).
+
+The norm form is the one non-obvious rule. HF's Gemma3 RMSNorm is `x/rms * (1 + w)`, but these containers store
+`w` with the 1 already added. Every norm kind of the Medgemma pack (input, post-attention, pre- and post-feedforward,
+q, final) equals llama.cpp's GGUF of the same model, whose converter adds the 1, to bf16 rounding. With the 1 added
+again, the known-answer prompt answered `*,,,,,,,`; with the plain form it answers `Paris`. So gemma3 takes the plain
+form, and gemma2 keeps `(1 + w)` (not measured here). `tests/q4nx_gemma3_semantics_test.c` pins each rule with
+Medgemma's own keys.
+
+A prefill runs batched (`step_batch`) with the same three extra norms; only Gemma4 still falls back to one
+`step()` per prompt token. `batch_check_state_test` on the Medgemma pack: batched equals sequential, rel-L1 0.
+
+The serve also needs the BOS. Medgemma's `config.json` states neither `bos_token_id` nor `add_bos_token`, and its
+`tokenizer_config.json` states both (`<bos>`, true), as HF reads them. forward-serve now takes what `config.json`
+does not state from `tokenizer_config.json`. Served without it, the reply to "Explain in three sentences why the sky
+is blue." was "The sky is blue because of the way the light is." With it, the reply is "The sky appears blue due to a
+phenomenon called Rayleigh scattering", the same as the host forward with the BOS.
+
+On the device, with a locally built kernel set of the four shapes (`QKV K2560 N4096`, `O K2048 N2560`,
+`GU K2560 N20480`, `D K10240 N2560`) and their 16-row twins, in one `npu-hold` ticket, the known answer came back
+`Paris` three times, identical. That is not yet a leaderboard row: decode is 0.68 tok/s, and no kernel set for these
+shapes ships.
+
+**Gemma3-270M-NPU2 and the 1,280-byte quarter tile.** This pack's hidden size, 640, is not a multiple of the
+5,120-byte tile's 256 columns. FastFlowLM stores every I8 tensor of it in a quarter tile instead: 16 rows x 128
+columns, with scales and zero-points at `g*16 + r` and codes at `c*8 + r/2` (`Q4NX_CHUNK_Q4H`, `q4nx_dequant.h`).
+Both kinds cost 0.625 bytes a weight, so the byte count cannot tell them apart. The tensor's declared shape does:
+`[chunks, 1280]` against `[tiles, 5120]`, and the reader goes by it. The reading was settled against an independent
+GGUF of the same model (`v_proj` against its Q5_0 `attn_v`: correlation 0.9962; the 5,120-tile reading gives 0.09)
+and against the pack's own embedding (`lm_head`: 0.9963, at most 0.57 of a step off). On the host the pack answers
+the known-answer prompt `Paris` and explains Rayleigh scattering in coherent sentences. Before, it failed at load
+with `chunk 4266 B is not a kind`. No kernel set for its 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
@@ -2040,7 +2089,7 @@ engine#587 measured the route on it):
 | `dflash` | `FastFlowLM/Nemotron-3.5-Bolt-E4B-DFlash-NPU2` | hrx | not served: a draft head, not a decoder; "The draft loop, and whether to serve it" below covers what the HRX route does with 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, npu | **served and measured** on the per-op route ("FastFlowLM's own containers" below) |
-| `gemma3` | `FastFlowLM/medgemma-1.5-4b-it-NPU2` | hrx | not served |
+| `gemma3` | `FastFlowLM/medgemma-1.5-4b-it-NPU2` | hrx | not served on the device. The forward reads the family and answers on the host ("Gemma3 text decoders" above); one kernel shape, `QKV K2560 N4096`, is missing |
 | `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 |
@@ -2068,7 +2117,7 @@ pack declares (a config with no `architectures` declares its `model_type`):
 | `DFlashDraftModel` | Nemotron-3.5-Bolt-E4B-DFlash | not served: a draft head, not a decoder | engine#1780 |
 | `Gemma4ForConditionalGeneration`, `Gemma4UnifiedForConditionalGeneration` | Nemotron-3.5-Bolt-E4B, Gemma4-E4B-IT, Gemma4-E2B-IT, Gemma4-12B-IT | not served: `gemma4`, hrx only | engine#1786 |
 | `HunYuanDenseV1ForCausalLM` | Hy-MT2-1.8B | not served: `hunyuan-dense`, hrx only | engine#1786 |
-| `Gemma3ForCausalLM`, `gemma3_text` | Gemma3-270M, Gemma3-1B, Gemma3-4B, Gemma3-4B-Text, medgemma-4b-it, medgemma-1.5-4b-it, Translategemma-4B-Instruct | not served: `gemma3`, hrx only | engine#1786 |
+| `Gemma3ForCausalLM`, `gemma3_text` | Gemma3-270M, Gemma3-1B, Gemma3-4B, Gemma3-4B-Text, medgemma-4b-it, medgemma-1.5-4b-it, Translategemma-4B-Instruct | not served: `gemma3`, hrx only (the per-op forward reads the family on the host; the device needs `QKV K2560 N4096`) | engine#1786 |
 | `Lfm2ForCausalLM` | LFM2-1.2B, LFM2-2.6B, LFM2-1.2B-Tool, LFM2-2.6B-Transcript, LFM2.5-1.2B, LFM2.5-1.2B-Thinking | not in the registry's NPU set: `lfm2`, hrx only (LFM2.5-1.2B on the per-op route is engine#2170, which adds no `NPU_ARCHS` entry) | engine#1786 |
 | `GptOssForCausalLM` | GPT-OSS-20B, GPT-OSS-Safeguard-20b | not served: `gpt-oss`, hrx only | engine#1786 |
 | `Qwen3_5MoeForConditionalGeneration` | Qwen3.6-35B-A3B | not served for this class: the route serves `qwen3_5_moe` through `Qwen3_5MoeForCausalLM` (GGUF on the NPU); the VLM wrapper class is hrx only, and FastFlowLM's own pack is not measured on the route | engine#1786 |

npu/forward/include/q4nx_dequant.h

@@ -101,6 +101,27 @@ extern "C" {
 #define Q4NX_CHUNK_Q4S 4608
 void q4nx_dequant_tile_q4s(const uint8_t* chunk, float* out, int ld);
 
+/* The quarter tile (FastFlowLM's Gemma3-270M-NPU2, whose K = 640 is not on the 256-column grid): 16 rows x 128 columns
+ * in 1280 bytes, the 5120-byte tile's reading with one half and four groups:
+ *   scales      [64] bf16 @[0,128)     index g*16 + r
+ *   zero-points [64] bf16 @[128,256)   index g*16 + r
+ *   codes            @[256,1280)       byte c*8 + r/2, even row = low nibble
+ *   value[r][c] = code * scale + zero   (the symmetric kind: zero-points 0, signed scale, two's-complement code)
+ * A tensor of these declares its shape as [chunks, 1280], where the tiled kind declares [tiles, 5120]; both cost 0.625
+ * bytes a weight, so the byte count cannot tell them apart and the declared row width is what names the kind. The grid
+ * is (rows/16) x (K/128), row-major. Settled 2026-11-12 on that pack: v_proj against the Q5_0 attn_v of an
+ * independent GGUF of the same model, corr 0.9962 (the next reading 0.84, the 5120-tile reading 0.09); and lm_head
+ * against the pack's own bf16 embedding (which it stores x sqrt(H)), corr 0.9963, |dequant - embedding| / |scale|
+ * mean 0.235, max 0.567. */
+#define Q4NX_CHUNK_Q4H 1280
+#define Q4NX_Q4H_ROWS 16
+#define Q4NX_Q4H_COLS 128
+void q4nx_dequant_chunk_q4h(const uint8_t* chunk, float* out, int ld, int symmetric);
+/* A [rows, k] tensor of quarter tiles (k = the container's own K, a multiple of 128) read into `out` as [rows, cols]
+ * with cols >= k; columns k..cols-1 are written 0 (the forward's K pad). Returns 0, or -1 when the sizes disagree or
+ * 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);
+
 /* 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_forward.h

@@ -517,6 +517,7 @@ class Q4nxNpuForward {
     std::vector<float> rope_scale_;  // LongRoPE short_factor (per-half-dim), empty when none
     Q4nxRopeLlama3 rope_llama3_;     // rope_type "llama3" (Llama-3.1/3.2/3.3); on == 0 for everything else
     Q4nxRopeDynamic rope_dynamic_;   // rope_type "dynamic" (NTK base rule, engine#1740 W2); on == 0 elsewhere
+    double rope_linear_factor_ = 0.0; // rope_type "linear" factor (Gemma3's full layers); 0 elsewhere
     // MiniCPM4 scales (config "embedding_scale" / "residual_scale"); 1.0 for
     // every other architecture.  MiniCPM4 scales the embedding by 12 and each
     // residual branch by scale_depth/sqrt(NL) (0.2475 for 32 layers); skipping

npu/forward/include/q4nx_semantics.h

@@ -152,8 +152,16 @@ typedef struct {
      * Source: num_kv_shared_layers. */
     int num_kv_shared_layers;
     /* Gemma4 adds pre/post feedforward norms and a post-PLE norm that the plain
-     * layer does not have.  Source: presence of those tensors / model_type. */
+     * layer does not have (1).  Gemma3 has the same sandwich around attention and
+     * MLP but no post-PLE norm and no layer scalar (2).  Source: model_type. */
     int extra_layer_norms;
+    /* Gemma3 states the sliding layers' RoPE base on its own, as
+     * `rope_local_base_freq`, and its `rope_scaling` applies to the full layers
+     * only.  1 when that key was read into rope_theta_sliding. */
+    int rope_local_base;
+    /* Gemma2/3: attention scores are scaled by 1/sqrt(query_pre_attn_scalar)
+     * instead of 1/sqrt(head_dim).  0 when absent.  Source: query_pre_attn_scalar. */
+    float query_pre_attn_scalar;
     /* Gemma4 use_double_wide_mlp: KV-shared layers run a 2*IM-wide MLP.
      * Source: use_double_wide_mlp. */
     int double_wide_mlp;

npu/forward/src/q4nx_dequant.c

@@ -150,6 +150,53 @@ void q4nx_dequant_tile_q4s(const uint8_t* t, float* out, int ld) {
     }
 }
 
+void q4nx_dequant_chunk_q4h(const uint8_t* t, float* out, int ld, int symmetric) {
+    const uint8_t* sc_b = t;
+    const uint8_t* zp_b = t + 128;
+    const uint8_t* pk = t + 256;
+    for (int r = 0; r < Q4NX_Q4H_ROWS; r++) {
+        for (int c = 0; c < Q4NX_Q4H_COLS; c++) {
+            const int i = (c / Q4NX_GROUP_SIZE) * Q4NX_Q4H_ROWS + r;
+            const uint8_t byte = pk[c * 8 + r / 2];
+            int code = (r % 2 == 0) ? (byte & 0xF) : (byte >> 4);
+            const float s = bf16_to_f32(le16(sc_b + 2 * i));
+            if (symmetric) {
+                if (code >= 8) code -= 16;
+                out[(size_t)r * ld + c] = (float)code * s;
+            } else {
+                out[(size_t)r * ld + c] = (float)code * s + bf16_to_f32(le16(zp_b + 2 * i));
+            }
+        }
+    }
+}
+
+int q4nx_dequant_tensor_q4h(const uint8_t* data, size_t nbytes, int rows, int k, int cols, float* out) {
+    if (!data || !out || rows <= 0 || k <= 0 || cols < k) return -1;
+    if (rows % Q4NX_Q4H_ROWS || k % Q4NX_Q4H_COLS) return -1;
+    const int ntr = rows / Q4NX_Q4H_ROWS, ntc = k / Q4NX_Q4H_COLS;
+    const size_t n = (size_t)ntr * (size_t)ntc;
+    if (nbytes != n * (size_t)Q4NX_CHUNK_Q4H) return -1;
+    /* One kind per tensor, the 5120-byte rule: symmetric when every zero-point is 0 and some scale is negative. */
+    int zero_zp = 1, neg = 0;
+    for (size_t i = 0; i < n && (zero_zp || !neg); i++) {
+        const uint8_t* c = data + i * Q4NX_CHUNK_Q4H;
+        for (int j = 0; j < 64; j++) {
+            if (le16(c + 128 + 2 * j) & 0x7fff) zero_zp = 0;
+            if (le16(c + 2 * j) & 0x8000) neg = 1;
+        }
+    }
+    const int sym = zero_zp && neg;
+    #pragma omp parallel for collapse(2) schedule(static) if(n >= 4)
+    for (int tr = 0; tr < ntr; tr++)
+        for (int tc = 0; tc < ntc; tc++)
+            q4nx_dequant_chunk_q4h(data + ((size_t)tr * ntc + tc) * Q4NX_CHUNK_Q4H,
+                                   out + (size_t)tr * Q4NX_Q4H_ROWS * cols + (size_t)tc * Q4NX_Q4H_COLS, cols, sym);
+    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;