From 3a064df5e538f8f3c1760595527bd87705ce60fb Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 21:58:03 -0400 Subject: [PATCH 1/3] npu: the per-op forward reads gemma3 text decoders (Refs #1786) FastFlowLM's Medgemma-4B-NPU2 was refused at load (rope scaling "linear"), and four other gemma3 rules were missing: the layer mix from sliding_window_pattern when there is no layer_types list, the sliding layers' rope_local_base_freq (a sliding layer with no base of its own had theta 0), query_pre_attn_scalar, and the sandwich norms without Gemma4's post-PLE norm and layer scalar. Linear scaling divides the frequency on the full layers only when a local base is stated, and on every layer otherwise. The gemma3 containers carry the norm's 1 already (they equal llama.cpp's GGUF norms to bf16), so gemma3 takes the plain form; with (1 + w) the known answer was "*,,,,,,,", now it is "Paris". Co-Authored-By: Claude Opus 5.5 --- CMakeLists.txt | 5 ++ docs/npu.md | 35 ++++++++- npu/forward/include/q4nx_forward.h | 1 + npu/forward/include/q4nx_semantics.h | 10 ++- npu/forward/src/q4nx_forward.cpp | 49 ++++++++++-- npu/forward/src/q4nx_semantics.c | 46 +++++++++-- tests/q4nx_gemma3_semantics_test.c | 113 +++++++++++++++++++++++++++ 7 files changed, 243 insertions(+), 16 deletions(-) create mode 100644 tests/q4nx_gemma3_semantics_test.c diff --git a/CMakeLists.txt b/CMakeLists.txt index a6a5c3dcb..8d5185da8 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/docs/npu.md b/docs/npu.md index 057da210a..43efa0808 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -1949,6 +1949,37 @@ 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. + +Not yet a device row: the pack's four shapes are `QKV K2560 N4096`, `O K2048 N2560`, `GU K2560 N20480` and +`D K10240 N2560`, and `QKV K2560 N4096` is not built in any kernel family on the box (see +[arch-gaps.md](arch-gaps.md)). Gemma3 layers also take the per-token path for a prefill (like Gemma4), so a +prefill costs one `step()` per prompt token. Gemma3-270M-NPU2 still fails at load: its tiles are on an unpadded +K = 640 grid that the decoder does not read yet. + **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 +2071,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 +2099,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 | diff --git a/npu/forward/include/q4nx_forward.h b/npu/forward/include/q4nx_forward.h index 3448703d2..227344fb7 100644 --- a/npu/forward/include/q4nx_forward.h +++ b/npu/forward/include/q4nx_forward.h @@ -517,6 +517,7 @@ class Q4nxNpuForward { std::vector 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 diff --git a/npu/forward/include/q4nx_semantics.h b/npu/forward/include/q4nx_semantics.h index 1a5af0c9b..e7bb68e1f 100644 --- a/npu/forward/include/q4nx_semantics.h +++ b/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; diff --git a/npu/forward/src/q4nx_forward.cpp b/npu/forward/src/q4nx_forward.cpp index fd8529883..d0b7d2d42 100644 --- a/npu/forward/src/q4nx_forward.cpp +++ b/npu/forward/src/q4nx_forward.cpp @@ -234,6 +234,26 @@ Q4nxRopeDynamic read_rope_dynamic(const std::string& cfg_path) { return d; } +// rope_type "linear": the factor of the config's rope_scaling block when that block declares "linear" with a factor +// above zero, else 0 (no scaling; a "linear" block without a usable factor is still refused by the caller). +static double read_rope_linear(const std::string& cfg_path) { + std::ifstream f(cfg_path); + if (!f) return 0.0; + std::string t((std::istreambuf_iterator(f)), std::istreambuf_iterator()); + char type[64]; + q4nx_rope_declared_type(t.c_str(), type, (int)sizeof type); + if (std::strcmp(type, "linear")) return 0.0; + const size_t rs = t.find("\"rope_scaling\""); + const size_t at = rs == std::string::npos ? rs : t.find("\"factor\"", rs); + // the key must sit inside the block: before the first closing brace after it (the block holds no objects) + if (at == std::string::npos || at > t.find('}', rs)) return 0.0; + const size_t colon = t.find(':', at); + if (colon == std::string::npos) return 0.0; + char* end = nullptr; + const double v = std::strtod(t.c_str() + colon + 1, &end); + return (end != t.c_str() + colon + 1 && v > 0.0) ? v : 0.0; +} + static Q4nxRopeLlama3 read_rope_llama3(const std::string& cfg_path) { Q4nxRopeLlama3 s{}; std::ifstream f(cfg_path); @@ -803,6 +823,7 @@ bool Q4nxNpuForward::init_impl(xrt::device* dev, const char* model_path, rope_scale_ = read_rope_short_factor(cfg_path); // LongRoPE (MiniCPM4) rope_llama3_ = read_rope_llama3(cfg_path); // rope_type "llama3" (engine#1048) rope_dynamic_ = read_rope_dynamic(cfg_path); // rope_type "dynamic" (NTK by base, engine#1740 W2) + rope_linear_factor_ = read_rope_linear(cfg_path); // rope_type "linear" (Gemma3's full layers, engine#1786) { // A declared scaling this forward does not apply is refused BY NAME (engine#1087): serving it // would run plain rope, stay fluent and be wrong past the original window, with nothing in the @@ -811,15 +832,17 @@ bool Q4nxNpuForward::init_impl(xrt::device* dev, const char* model_path, std::string t((std::istreambuf_iterator(f)), std::istreambuf_iterator()); char declared[64]; q4nx_rope_declared_type(t.c_str(), declared, (int)sizeof declared); - if (!q4nx_rope_type_supported(declared, &rope_llama3_, rope_scale_.empty() ? 0 : 1, rope_dynamic_.on)) { + const bool linear_ok = !std::strcmp(declared, "linear") && rope_linear_factor_ > 0.0; + if (!linear_ok && + !q4nx_rope_type_supported(declared, &rope_llama3_, rope_scale_.empty() ? 0 : 1, rope_dynamic_.on)) { if (!std::strcmp(declared, "llama3")) err_ = "config.json declares rope_type \"llama3\" but its numbers are not usable (factor > 0, " "original_max_position_embeddings > 0, high_freq_factor > low_freq_factor), so the " "scaling cannot be applied; serving it would run plain RoPE"; else err_ = std::string("config.json declares rope scaling \"") + declared + - "\", which this forward does not implement (only llama3 and longrope short_factor " - "are), so it would be served plain RoPE and reply wrongly past the model's " + "\", which this forward does not implement (only llama3, linear with a factor and " + "longrope short_factor are), so it would be served plain RoPE and reply wrongly past the model's " "original window (engine#1048, #1087)"; return false; } @@ -1866,8 +1889,12 @@ bool Q4nxNpuForward::load_norms() { pre_ff_norm_[l].data(), H)) return false; if (!rd("model.layers.%d.post_feedforward_layernorm.weight", post_ff_norm_[l].data(), H)) return false; - if (!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; + // 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; + if (!rd("model.layers.%d.layer_output_scale.weight", &layer_scale_[l], 1)) return false; + } } if (ple_dim_ > 0) { inp_gate_[l].assign((size_t)ple_dim_ * (size_t)H, 0.0f); @@ -4232,7 +4259,9 @@ bool Q4nxNpuForward::step_batch_impl(const std::vector& tokens, int start_p } { LhScope _t_attn(LH_ATTN); - const float scale = sem_.attn_scaling_one ? 1.0f : 1.0f / std::sqrt((float)l_hd); + const float scale = sem_.attn_scaling_one ? 1.0f + : 1.0f / std::sqrt(sem_.query_pre_attn_scalar > 0 ? sem_.query_pre_attn_scalar + : (float)l_hd); const int nt = ((long)M * (start_pos + M) * l_hd >= 8192) ? omp_heads_threads() : 1; q4nx_attn::rows_context(qall.data(), kcache_[l].data(), vcache_[l].data(), M, c_.NH, l_nkv, l_hd, Tm.data(), t0m.data(), scale, CTX.data(), nt); @@ -5081,6 +5110,10 @@ bool Q4nxNpuForward::attention(const std::vector& qkv, int pos, const int // llama3 (Llama-3.1/3.2/3.3): the wavelength-band rule, on the FREQUENCY, so it costs // nothing per token. No-op unless the config declares rope_type "llama3". inv = q4nx_rope_llama3_inv_freq(inv, &rope_llama3_); + // rope_type "linear": positions divided by the factor, i.e. the frequency is. Gemma3 applies its + // block to the full layers only; its sliding layers have their own base (rope_local_base_freq) and + // no scaling. A model without that split scales every layer, as HF does. + if (rope_linear_factor_ > 0.0 && (window == 0 || !sem_.rope_local_base)) inv /= rope_linear_factor_; } // Interleaved MRoPE: each rotary frequency rotates by its own axis's position (t, h or w). With the // three planes equal this is exactly `pos * inv`, which is why text-only prompts are bit-identical @@ -5121,7 +5154,9 @@ bool Q4nxNpuForward::attention(const std::vector& qkv, int pos, const int const int T = kv_len > 0 ? kv_len : pos + 1; // Gemma4 sets attention scaling to 1.0 (q/k carry an RMSNorm); everyone else // uses the usual 1/sqrt(head_dim). - const float scale = sem_.attn_scaling_one ? 1.0f : 1.0f / std::sqrt((float)HD); + // Gemma2/3 scale by query_pre_attn_scalar, which equals head_dim on most of the family but not all (Gemma3-27B). + const float scale = sem_.attn_scaling_one ? 1.0f + : 1.0f / std::sqrt(sem_.query_pre_attn_scalar > 0 ? sem_.query_pre_attn_scalar : (float)HD); // Sliding-window attention: ignore positions older than the window. const int t0 = (window > 0 && T > window) ? T - window : 0; diff --git a/npu/forward/src/q4nx_semantics.c b/npu/forward/src/q4nx_semantics.c index 56a20bbbf..671069218 100644 --- a/npu/forward/src/q4nx_semantics.c +++ b/npu/forward/src/q4nx_semantics.c @@ -25,7 +25,7 @@ #include static const char* kFields[Q4NX_SEM_COUNT] = { - /* RMSNORM_OFFSET */ "model_type/architectures (gemma* uses the 1+w form)", + /* RMSNORM_OFFSET */ "model_type/architectures (gemma2 uses the 1+w form; gemma3 containers carry the 1)", /* QK_NORM */ "use_qk_norm when stated, else model_type/architectures", /* ROPE_THETA */ "rope_theta", /* ROTARY_DIM */ "partial_rotary_factor * head_dim", @@ -209,12 +209,20 @@ int q4nx_semantics_from_config(const char* json_path, const char* model_dir, * initialised to ones, whereas gemma2/gemma3 do `x/rms * (1 + w)`. */ const int is_gemma4 = strstr(model_type, "gemma4") != NULL || strstr(arch, "Gemma4") != NULL; + /* Gemma3 (gemma3, gemma3_text, FastFlowLM's gemma3_text_only): the sandwich norms, a local RoPE base for the sliding + * layers and the 5:1 sliding pattern. gemma3n is another architecture (AltUp, laurel) and is not this. */ + const int is_gemma3 = (starts_with(model_type, "gemma3") && !starts_with(model_type, "gemma3n")) || + (starts_with(arch, "Gemma3") && !starts_with(arch, "Gemma3n")); const int is_qwen3 = starts_with(model_type, "qwen3") || starts_with(arch, "Qwen3"); const int is_phi = starts_with(model_type, "phi") || starts_with(arch, "Phi"); const int is_qwen35 = starts_with(model_type, "qwen3_5") || starts_with(arch, "Qwen3_5"); - /* RMSNorm convention: gemma2/gemma3 scale by (1 + w); gemma4 by w. */ - s->rmsnorm_offset = (is_gemma && !is_gemma4) ? 1 : 0; + /* RMSNorm convention: gemma2 scales by (1 + w); gemma4 by w. Gemma3 is (1 + w) in HF, but the containers this + * forward reads (FastFlowLM's gemma3 NPU2 packs) store the norm weights with the 1 already added, as llama.cpp's + * GGUF does: on Medgemma-4B-NPU2 every norm kind (input, post-attention, pre/post feedforward, q, final) equals + * the GGUF's to bf16 rounding, and adding 1 again answered the known-answer prompt "*,,,,,," where the plain + * form answers "Paris" (engine#1786). */ + s->rmsnorm_offset = (is_gemma && !is_gemma4 && !is_gemma3) ? 1 : 0; /* qwen3_5 ships both conventions under the same model_type, so the forward resolves it from the * 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; @@ -366,6 +374,13 @@ int q4nx_semantics_from_config(const char* json_path, const char* model_dir, } else if (get_num(json, "rope_theta", &d) && d > 0) { s->rope_theta = (float)d; } + /* Gemma3 gives the sliding layers their own base (10000 against the full layers' 1e6). Without it a sliding + * layer was rotated with theta 0, i.e. no usable rotation at all (engine#1786). */ + if (!s->per_layer_rope && get_num(json, "rope_local_base_freq", &d) && d > 0) { + s->rope_theta_sliding = (float)d; + s->rope_local_base = 1; + } + if (get_num(json, "query_pre_attn_scalar", &d) && d > 0) s->query_pre_attn_scalar = (float)d; /* partial rotary: top-level key, or nested in rope_parameters (scalar form), * or nested per layer type (Gemma4's full_attention.partial_rotary_factor). */ @@ -476,6 +491,24 @@ int q4nx_semantics_from_config(const char* json_path, const char* model_dir, else s->sliding_window = 0; + /* Gemma3 configs that predate layer_types state the pattern instead: every `sliding_window_pattern`-th layer is + * full attention, the rest slide (HF: is_sliding = bool((layer_idx + 1) % sliding_window_pattern)). */ + if (s->n_layer_types == 0 && s->sliding_window > 0 && get_num(json, "sliding_window_pattern", &d) && d >= 1) { + double nl = 0; + const int pat = (int)(d + 0.5); + if (get_num(json, "num_hidden_layers", &nl) && nl > 0 && nl <= Q4NX_MAX_LAYERS) { + int hybrid = 0; + for (int i = 0; i < (int)nl; i++) { + s->layer_types[i] = ((i + 1) % pat) ? Q4NX_LAYER_SLIDING : Q4NX_LAYER_FULL; + if (s->layer_types[i] != Q4NX_LAYER_FULL) hybrid++; + } + s->n_layer_types = (int)nl; + s->hybrid_layers = hybrid; + } + } + /* A sliding layer with no base of its own uses the model's (theta 0 is not a base). */ + if (s->rope_theta_sliding <= 0) s->rope_theta_sliding = s->rope_theta; + /* 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) */ @@ -489,7 +522,7 @@ int q4nx_semantics_from_config(const char* json_path, const char* model_dir, if (get_num(json, "num_kv_shared_layers", &d) && d > 0) s->num_kv_shared_layers = (int)(d + 0.5); /* Gemma4 adds pre/post feedforward norms and a post-PLE norm. */ - s->extra_layer_norms = is_gemma4 ? 1 : 0; + s->extra_layer_norms = is_gemma4 ? 1 : (is_gemma3 ? 2 : 0); /* `use_double_wide_mlp` is a JSON boolean, so read it as "present in text". */ s->double_wide_mlp = (strstr(json, "use_double_wide_mlp\":true") != NULL || strstr(json, "use_double_wide_mlp\": true") != NULL) ? 1 : 0; @@ -602,9 +635,10 @@ void q4nx_semantics_dump(const Q4nxSemantics* s) { fprintf(stderr, " [semantics] layer mix: %s <- %s\n", mix, q4nx_semantics_field(Q4NX_SEM_LAYER_TYPES)); fprintf(stderr, " [semantics] per-type rope: full(theta=%.6g rot=%d) " - "sliding(theta=%.6g rot=%d) rope_type_named=%d\n", + "sliding(theta=%.6g rot=%d%s) rope_type_named=%d query_pre_attn_scalar=%.6g\n", s->rope_theta, s->rotary_dim, s->rope_theta_sliding, - s->rotary_dim_sliding, s->rope_type_named); + s->rotary_dim_sliding, s->rope_local_base ? ", rope_local_base_freq" : "", + s->rope_type_named, s->query_pre_attn_scalar); fprintf(stderr, " [semantics] ple=%d (dim=%d vocab=%d) kv_shared_layers=%d " "extra_layer_norms=%d double_wide_mlp=%d\n", q4nx_semantics_uses_ple(s), s->hidden_size_per_layer_input, diff --git a/tests/q4nx_gemma3_semantics_test.c b/tests/q4nx_gemma3_semantics_test.c new file mode 100644 index 000000000..7eafddcd0 --- /dev/null +++ b/tests/q4nx_gemma3_semantics_test.c @@ -0,0 +1,113 @@ +/* Copyright 2026 bong-water-water-bong + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +/* q4nx_gemma3_semantics_test -- what the forward reads from a gemma3 config (engine#1786). + * + * Medgemma-4B-NPU2's config.json states its layer mix as `sliding_window_pattern: 6` (no layer_types), its sliding + * layers' RoPE base as `rope_local_base_freq` and its attention scale as `query_pre_attn_scalar`. Before this, none + * of the three was read: the 29 sliding layers ran as full attention with no window, and a sliding layer of a config + * that does list layer_types (Gemma3-270M-NPU2) was rotated with theta 0. The gemma3 containers also carry the norm's + * `1 +` already, so the plain form is the right one for them, where gemma2 keeps (1 + w). + * + * The strings are cut down from the configs on this box to the fields the rules read. No model, no device. */ +#include "q4nx_semantics.h" + +#include +#include +#include +#include +#include + +static int failures = 0; + +static void check(int ok, const char* what) { + printf("%s %s\n", ok ? "ok " : "FAIL", what); + if (!ok) failures++; +} + +static int sem_of(const char* dir, const char* json, Q4nxSemantics* s) { + char path[512]; + snprintf(path, sizeof path, "%s/config.json", dir); + int fd = open(path, O_WRONLY | O_CREAT | O_TRUNC, 0600); + FILE* f = fd >= 0 ? fdopen(fd, "w") : NULL; + if (!f) { printf("FAIL cannot write %s\n", path); failures++; return -1; } + fputs(json, f); + fclose(f); + if (q4nx_semantics_from_config(path, dir, s) != 0) { printf("FAIL cannot read %s\n", path); failures++; return -1; } + return 0; +} + +static const char* kMedgemma = + "{\"head_dim\": 256, \"hidden_size\": 2560, \"model_type\": \"gemma3_text\", \"num_attention_heads\": 8," + " \"num_hidden_layers\": 34, \"num_key_value_heads\": 4, \"query_pre_attn_scalar\": 256," + " \"rope_local_base_freq\": 10000.0, \"rope_scaling\": {\"factor\": 8.0, \"rope_type\": \"linear\"}," + " \"rope_theta\": 1000000.0, \"sliding_window\": 1024, \"sliding_window_pattern\": 6}"; + +int main(void) { + char dir[] = "/tmp/q4nx-gemma3-sem-XXXXXX"; + if (!mkdtemp(dir)) { printf("FAIL mkdtemp\n"); return 1; } + Q4nxSemantics s; + + if (sem_of(dir, kMedgemma, &s) == 0) { + int full = 0, sliding = 0, pattern_ok = 1; + for (int i = 0; i < s.n_layer_types; i++) { + if (s.layer_types[i] == Q4NX_LAYER_FULL) full++; + else if (s.layer_types[i] == Q4NX_LAYER_SLIDING) sliding++; + const int want = ((i + 1) % 6) ? Q4NX_LAYER_SLIDING : Q4NX_LAYER_FULL; + if (s.layer_types[i] != want) pattern_ok = 0; + } + check(s.n_layer_types == 34 && full == 5 && sliding == 29, + "Medgemma: sliding_window_pattern 6 over 34 layers gives 5 full and 29 sliding layers"); + check(pattern_ok, "... every 6th layer (5, 11, ..., 29) is the full one, as HF's (i + 1) % 6 says"); + check(s.hybrid_layers == 29, "... and the hybrid count says 29"); + check(s.rope_theta == 1000000.0f && s.rope_theta_sliding == 10000.0f && s.rope_local_base, + "Medgemma: full layers theta 1e6, sliding layers rope_local_base_freq 10000"); + check(s.query_pre_attn_scalar == 256.0f, "Medgemma: query_pre_attn_scalar is read"); + check(s.extra_layer_norms == 2, "gemma3: the two feedforward norms, without Gemma4's post-PLE norm and scalar"); + check(s.rmsnorm_offset == 0, "gemma3 containers carry the norm's 1: the plain form"); + check(s.sliding_window == 1024, "the window is unchanged"); + } + + /* Gemma3-270M-NPU2 lists layer_types and has no local base key: the list wins, and the sliding layers take the + * model's theta rather than 0. */ + if (sem_of(dir, "{\"model_type\": \"gemma3_text_only\", \"num_hidden_layers\": 3, \"rope_theta\": 1000000.0," + " \"sliding_window\": 512, \"sliding_window_pattern\": 6," + " \"layer_types\": [\"sliding_attention\", \"full_attention\", \"sliding_attention\"]}", + &s) == 0) { + check(s.n_layer_types == 3 && s.layer_types[1] == Q4NX_LAYER_FULL && s.layer_types[0] == Q4NX_LAYER_SLIDING, + "a stated layer_types list is not replaced by the pattern"); + check(s.rope_theta_sliding == 1000000.0f && !s.rope_local_base, + "with no rope_local_base_freq a sliding layer takes the model's theta, not 0"); + check(s.extra_layer_norms == 2 && s.rmsnorm_offset == 0, "gemma3_text_only is gemma3 too"); + } + + /* 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"); + if (sem_of(dir, "{\"model_type\": \"gemma3n_text\", \"hidden_size\": 2048}", &s) == 0) + check(s.extra_layer_norms == 0, "gemma3n is another architecture, not gemma3"); + if (sem_of(dir, "{\"model_type\": \"qwen3\", \"rope_theta\": 1000000.0, \"sliding_window\": 4096," + " \"num_hidden_layers\": 4}", &s) == 0) + check(s.n_layer_types == 0 && s.query_pre_attn_scalar == 0.0f && !s.rope_local_base, + "a config with a window but no pattern gets no layer list and no gemma keys"); + + char path[512]; + snprintf(path, sizeof path, "%s/config.json", dir); + unlink(path); + rmdir(dir); + printf(failures ? "q4nx_gemma3_semantics: %d FAILED\n" : "q4nx_gemma3_semantics: ok\n", failures); + return failures ? 1 : 0; +} From de3a5d4adf29acc7166f28a6035c42ef6f4d4c5a Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 22:58:33 -0400 Subject: [PATCH 2/3] npu: gemma3 prefills batched, and the serve adds the BOS tokenizer_config.json states step_batch now runs gemma3 (extra_layer_norms 2) with the post-attention, pre- and post-feedforward norms instead of falling back to one step() per token; batch_check_state_test on Medgemma-4B: batched = sequential, rel-L1 0. forward-serve took bos_token_id/add_bos_token from config.json only. Medgemma's states neither, its tokenizer_config.json states both, and the served reply lost its BOS ("The sky is blue because of the way the light is."). What config.json does not state now comes from tokenizer_config.json; the device reply is the host forward's ("...due to a phenomenon called Rayleigh scattering"). Co-Authored-By: Claude Opus 5.5 --- app/forward_serve.cpp | 16 ++++++++++++++ app/npu_stop_tokens.h | 23 +++++++++++++++++--- docs/npu.md | 19 +++++++++++----- npu/forward/src/q4nx_forward.cpp | 37 +++++++++++++++++++++++++------- tests/npu_stop_tokens_test.cpp | 19 ++++++++++++++++ 5 files changed, 98 insertions(+), 16 deletions(-) diff --git a/app/forward_serve.cpp b/app/forward_serve.cpp index 27dede566..209f003ba 100644 --- a/app/forward_serve.cpp +++ b/app/forward_serve.cpp @@ -1115,6 +1115,22 @@ int run_forward_serve(int argc, char** argv) { e.tok = std::make_unique(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|>"); diff --git a/app/npu_stop_tokens.h b/app/npu_stop_tokens.h index 281adf045..4f357e514 100644 --- a/app/npu_stop_tokens.h +++ b/app/npu_stop_tokens.h @@ -58,14 +58,31 @@ inline std::vector 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 ("" or {"content": ""}). +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(); if (e.is_object() && e.contains("content") && e["content"].is_string()) return e["content"].get(); 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() ? 1 : 0; + if (a.is_number_integer()) return a.get() != 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. diff --git a/docs/npu.md b/docs/npu.md index 43efa0808..c3f6d9a68 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -1974,11 +1974,20 @@ again, the known-answer prompt answered `*,,,,,,,`; with the plain form it answe form, and gemma2 keeps `(1 + w)` (not measured here). `tests/q4nx_gemma3_semantics_test.c` pins each rule with Medgemma's own keys. -Not yet a device row: the pack's four shapes are `QKV K2560 N4096`, `O K2048 N2560`, `GU K2560 N20480` and -`D K10240 N2560`, and `QKV K2560 N4096` is not built in any kernel family on the box (see -[arch-gaps.md](arch-gaps.md)). Gemma3 layers also take the per-token path for a prefill (like Gemma4), so a -prefill costs one `step()` per prompt token. Gemma3-270M-NPU2 still fails at load: its tiles are on an unpadded -K = 640 grid that the decoder does not read yet. +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 (``, 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 still fails at load: its tiles are on an unpadded K = 640 grid that the decoder +does not read yet. **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*`. diff --git a/npu/forward/src/q4nx_forward.cpp b/npu/forward/src/q4nx_forward.cpp index d0b7d2d42..0a372ff85 100644 --- a/npu/forward/src/q4nx_forward.cpp +++ b/npu/forward/src/q4nx_forward.cpp @@ -3936,8 +3936,11 @@ bool Q4nxNpuForward::step_batch_impl(const std::vector& tokens, int start_p // recurrence), so lin_kd_ no longer disqualifies a model, and attn_output_gate // is handled in the attention block below. MLA, per-layer embeddings, Gemma4 // and shared-KV stacks still fall back. + // Gemma3 (extra_layer_norms 2) is Gemma4's sandwich without PLE or KV sharing, so the batched body + // below runs it with the three extra norms; Gemma4 itself still falls back. + const bool sandwich = gemma4_layer_ && ple_dim_ == 0 && sem_.extra_layer_norms == 2; const bool simple_dense = - !mla_ && ple_dim_ == 0 && !gemma4_layer_ && kv_self; + !mla_ && ple_dim_ == 0 && (!gemma4_layer_ || sandwich) && kv_self; if (!simple_dense || getenv("NPU_NO_BATCH")) { if (inject_embeds_) { err_ = "step_batch_embeds: this model's batched body falls back to per-token step(), which needs " @@ -4313,14 +4316,23 @@ bool Q4nxNpuForward::step_batch_impl(const std::vector& tokens, int start_p { LhScope _t_res(LH_RES); const int nrow = q4nx_pad_n(H); - for (int m = 0; m < M; m++) - for (int k = 0; k < H; k++) - X[(size_t)m * H + k] += O[(size_t)m * nrow + k] * residual_scale_; + if (sandwich) { + // as step(): the attention output is normalised (post_attention_layernorm) before the add + std::vector on((size_t)H); + for (int m = 0; m < M; m++) { + rmsnorm(O.data() + (size_t)m * nrow, post_norm_[l].data(), H, on.data()); + for (int k = 0; k < H; k++) X[(size_t)m * H + k] += on[(size_t)k]; + } + } else { + for (int m = 0; m < M; m++) + for (int k = 0; k < H; k++) + X[(size_t)m * H + k] += O[(size_t)m * nrow + k] * residual_scale_; + } } // ---- MLP ---- for (int m = 0; m < M; m++) - rmsnorm(X.data() + (size_t)m * H, post_norm_[l].data(), H, + rmsnorm(X.data() + (size_t)m * H, (sandwich ? pre_ff_norm_ : post_norm_)[l].data(), H, XN.data() + (size_t)m * H); if (n_expert_ > 0) { // MoE FFN: the router/top-k pick different experts per row, so this @@ -4409,9 +4421,18 @@ bool Q4nxNpuForward::step_batch_impl(const std::vector& tokens, int start_p { LhScope _t_res(LH_RES); const int nrow = q4nx_pad_n(H); // as above: the design's N, identity when H is on the grid - for (int m = 0; m < M; m++) - for (int k = 0; k < H; k++) - X[(size_t)m * H + k] += D[(size_t)m * nrow + k] * residual_scale_; + if (sandwich) { + // post_feedforward_layernorm on the down projection, then the add (as step()) + std::vector dn((size_t)H); + for (int m = 0; m < M; m++) { + rmsnorm(D.data() + (size_t)m * nrow, post_ff_norm_[l].data(), H, dn.data()); + for (int k = 0; k < H; k++) X[(size_t)m * H + k] += dn[(size_t)k]; + } + } else { + for (int m = 0; m < M; m++) + for (int k = 0; k < H; k++) + X[(size_t)m * H + k] += D[(size_t)m * nrow + k] * residual_scale_; + } } // Same post-residual hidden-state dump as step() emits as [L], one row per diff --git a/tests/npu_stop_tokens_test.cpp b/tests/npu_stop_tokens_test.cpp index 9ebb686f6..6daff22eb 100644 --- a/tests/npu_stop_tokens_test.cpp +++ b/tests/npu_stop_tokens_test.cpp @@ -152,6 +152,25 @@ int main() { "a marker the tokenizer does not have as one token adds nothing"); } + // engine#1786: the BOS a container's tokenizer_config.json states, for a config.json that states none + // (FastFlowLM's Medgemma-4B-NPU2: no bos_token_id, no add_bos_token in config.json). + { + using onebit::npu::add_bos_token_flag; + using onebit::npu::special_token_text; + const json med = json::parse(R"({"bos_token": "", "add_bos_token": true, "eos_token": ""})"); + check(special_token_text(med, "bos_token") == "" && add_bos_token_flag(med) == 1, + "Medgemma's tokenizer_config.json: BOS , added"); + check(special_token_text(json::parse(R"({"bos_token": {"content": "", "lstrip": false}})"), "bos_token") == "", + "bos_token in the object spelling"); + check(add_bos_token_flag(json::parse(R"({"add_bos_token": false})")) == 0 && + add_bos_token_flag(json::parse(R"({"add_bos_token": 1})")) == 1, + "add_bos_token false and 1"); + check(add_bos_token_flag(json::parse("{}")) == -1 && add_bos_token_flag(json::parse(R"({"add_bos_token": null})")) == -1 && + add_bos_token_flag(json(nullptr)) == -1 && special_token_text(json::parse(R"({"bos_token": null})"), "bos_token").empty(), + "an absent or null key says nothing"); + check(onebit::npu::eos_token_text(med) == "", "eos_token_text is unchanged by the shared reader"); + } + if (failures == 0) std::printf("all stop-token checks passed\n"); return failures == 0 ? 0 : 1; } From 5b698165b6560016a87ed6282c736d1860b13908 Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Wed, 11 Nov 2026 23:22:27 -0400 Subject: [PATCH 3/3] npu: read the 1,280-byte quarter tile, so Gemma3-270M-NPU2 loads (Refs #1786) Gemma3-270M's K = 640 is off the 5,120-byte tile's 256-column grid, and FastFlowLM stores every I8 tensor of it as 16 x 128 quarter tiles: scales and zero-points at g*16 + r, codes at c*8 + r/2, grid row-major. Both kinds cost 0.625 bytes a weight, so the declared row width (1280 against 5120) names the kind, and dequant_any() reads it into the forward's padded K with the pad columns 0. Settled against an independent GGUF of the model (v_proj vs its Q5_0 attn_v: corr 0.9962; the tile reading 0.09) and the pack's own embedding (lm_head: corr 0.9963, max 0.567 of a step). On the host the pack answers "Paris" where it failed at load ("chunk 4266 B is not a kind"). Co-Authored-By: Claude Opus 5.5 --- docs/npu.md | 13 +++++++-- npu/forward/include/q4nx_dequant.h | 21 +++++++++++++ npu/forward/src/q4nx_dequant.c | 47 ++++++++++++++++++++++++++++++ npu/forward/src/q4nx_forward.cpp | 32 ++++++++++++++------ tests/npu_dequant_bf16_test.c | 41 ++++++++++++++++++++++++++ 5 files changed, 143 insertions(+), 11 deletions(-) diff --git a/docs/npu.md b/docs/npu.md index c3f6d9a68..f8449e847 100644 --- a/docs/npu.md +++ b/docs/npu.md @@ -1986,8 +1986,17 @@ phenomenon called Rayleigh scattering", the same as the host forward with the BO 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 still fails at load: its tiles are on an unpadded K = 640 grid that the decoder -does not read yet. +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*`. diff --git a/npu/forward/include/q4nx_dequant.h b/npu/forward/include/q4nx_dequant.h index 9476d9dbc..35b7d5a04 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/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 diff --git a/npu/forward/src/q4nx_dequant.c b/npu/forward/src/q4nx_dequant.c index dca7810ce..ecb2438db 100644 --- a/npu/forward/src/q4nx_dequant.c +++ b/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; diff --git a/npu/forward/src/q4nx_forward.cpp b/npu/forward/src/q4nx_forward.cpp index 0a372ff85..bc7a4e443 100644 --- a/npu/forward/src/q4nx_forward.cpp +++ b/npu/forward/src/q4nx_forward.cpp @@ -186,8 +186,22 @@ static bool valid_sidecar_dir(const std::string& d) { // and a 3-D [ntr, ntc, N] tensor both work. (src/q4nx_dequant.c; the Q4_K/Q8 // layouts are in docs/qwen35-4b-weight-format.md.) -static int dequant_any(const void* data, size_t nbytes, int rows, int cols, float* out) { +// 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. +static int64_t declared_row_bytes(const TensorDesc* d) { + return (d && d->ndim == 2) ? d->shape[1] : 0; +} + +static int dequant_any(const void* data, size_t nbytes, int rows, int cols, float* out, int64_t decl_row_bytes = 0) { if (!data || !out || rows <= 0 || cols <= 0) return -1; + if (decl_row_bytes == Q4NX_CHUNK_Q4H) { + // Quarter tiles on the container's own K (Gemma3-270M: 640), read into the forward's padded width. + if (rows % Q4NX_Q4H_ROWS) return -1; + const size_t chunks = nbytes / Q4NX_CHUNK_Q4H, row_tiles = (size_t)(rows / Q4NX_Q4H_ROWS); + if (chunks % row_tiles) return -1; + return q4nx_dequant_tensor_q4h((const uint8_t*)data, nbytes, rows, + (int)(chunks / row_tiles) * Q4NX_Q4H_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; @@ -1033,7 +1047,7 @@ bool Q4nxNpuForward::init_impl(xrt::device* dev, const char* model_path, } if (!embed_tile_ && !embed_i8_) { // no element-wise read: keep the whole-table buffer (unchanged path) embed_f_.resize((size_t)NV * (size_t)kpad_); - if (dequant_any((const uint8_t*)e, esz, NV, kpad_, embed_f_.data()) != 0) { + if (dequant_any((const uint8_t*)e, esz, NV, kpad_, embed_f_.data(), declared_row_bytes(etd)) != 0) { err_ = std::string("embedding dtype ") + dt + " but dequant failed"; return false; } @@ -1620,7 +1634,7 @@ bool Q4nxNpuForward::init_impl(xrt::device* dev, const char* model_path, head_f32.assign((size_t)head_rows * (size_t)kpad_, 0.0f); head_out = head_f32.data(); } - if (!lp || dequant_any(lp, (size_t)ltd->data_size, head_rows, kpad_, head_out) != 0) { + if (!lp || dequant_any(lp, (size_t)ltd->data_size, head_rows, kpad_, head_out, declared_row_bytes(ltd)) != 0) { err_ = "lm_head dequant failed"; return false; } @@ -1977,7 +1991,7 @@ bool Q4nxNpuForward::load_layer(int l) { if (!d) { fprintf(stderr, " [mla] missing " ); snprintf(nm,sizeof(nm),fmt,l); fprintf(stderr,"%s\n", nm); return false; } v.resize((size_t)rows * (size_t)cols); const uint8_t* raw = (const uint8_t*)model_tensor_data(mw_, d); - return raw && dequant_any(raw, (size_t)d->data_size, rows, cols, v.data()) == 0; + return raw && dequant_any(raw, (size_t)d->data_size, rows, cols, v.data(), declared_row_bytes(d)) == 0; }; // A bf16 tensor (norms, router) is stored raw, not in 5120-byte tiles, so // it bypasses dequant_any. @@ -2059,7 +2073,7 @@ bool Q4nxNpuForward::load_layer(int l) { auto deq = [&](TensorDesc* d, float* out, int rows, int cols) -> bool { const uint8_t* raw = (const uint8_t*)model_tensor_data(mw_, d); if (!raw) { err_ = std::string("MLP tensor ") + d->name + " has no data"; return false; } - if (dequant_any(raw, (size_t)d->data_size, rows, cols, out) != 0) { + if (dequant_any(raw, (size_t)d->data_size, rows, cols, out, declared_row_bytes(d)) != 0) { // Name it: the bare "dequant failed" this used to leave cost an afternoon on engine#1097. const size_t tiles = (size_t)(rows / Q4NX_TILE_ROWS) * (size_t)(cols / Q4NX_TILE_COLS); char why[256]; @@ -2162,7 +2176,7 @@ bool Q4nxNpuForward::load_layer(int l) { if (!d) return false; v.resize((size_t)rows * (size_t)cols); const uint8_t* raw = (const uint8_t*)model_tensor_data(mw_, d); - return raw && dequant_any(raw, (size_t)d->data_size, rows, cols, v.data()) == 0; + return raw && dequant_any(raw, (size_t)d->data_size, rows, cols, v.data(), declared_row_bytes(d)) == 0; }; if (!load_share(mn.share_gate, share_g_, expert_im_, H)) share_g_.clear(); if (!load_share(mn.share_up, share_u_, expert_im_, H)) { share_g_.clear(); share_u_.clear(); } @@ -2195,7 +2209,7 @@ bool Q4nxNpuForward::load_gdn_layer(int l) { if (!d) { err_ = std::string("GDN: no tensor ") + nm; return false; } const uint8_t* raw = (const uint8_t*)model_tensor_data(mw_, d); if (!raw) { err_ = std::string("GDN: no data ") + nm; return false; } - if (dequant_any(raw, (size_t)d->data_size, rows, cols, out) != 0) { + if (dequant_any(raw, (size_t)d->data_size, rows, cols, out, declared_row_bytes(d)) != 0) { char m[64]; snprintf(m, sizeof(m), " (rows=%d cols=%d bytes=%llu)", rows, cols, (unsigned long long)d->data_size); @@ -2272,7 +2286,7 @@ bool Q4nxNpuForward::load_gdn_layer(int l) { auto deqT = [&](TensorDesc* d, float* out, int rows, int cols) -> bool { const uint8_t* raw = (const uint8_t*)model_tensor_data(mw_, d); if (!raw || d->data_size == 0) { err_ = std::string("GDN: missing MLP tensor ") + d->name; return false; } - if (dequant_any(raw, (size_t)d->data_size, rows, cols, out) != 0) { + if (dequant_any(raw, (size_t)d->data_size, rows, cols, out, declared_row_bytes(d)) != 0) { err_ = std::string("GDN: MLP dequant ") + d->name; return false; } @@ -2318,7 +2332,7 @@ bool Q4nxNpuForward::load_gdn_layer(int l) { if (!d) return false; v.resize((size_t)rows * (size_t)cols); const uint8_t* raw = (const uint8_t*)model_tensor_data(mw_, d); - return raw && dequant_any(raw, (size_t)d->data_size, rows, cols, v.data()) == 0; + return raw && dequant_any(raw, (size_t)d->data_size, rows, cols, v.data(), declared_row_bytes(d)) == 0; }; if (!load_share(mn.share_gate, share_g_, expert_im_, H)) share_g_.clear(); if (!load_share(mn.share_up, share_u_, expert_im_, H)) { share_g_.clear(); share_u_.clear(); } diff --git a/tests/npu_dequant_bf16_test.c b/tests/npu_dequant_bf16_test.c index 86626d70c..046643026 100644 --- a/tests/npu_dequant_bf16_test.c +++ b/tests/npu_dequant_bf16_test.c @@ -379,8 +379,49 @@ static void q4s_chunk(void) { free(whole); free(t); free(out); } +/* The 1,280-byte quarter tile (Gemma3-270M-NPU2): 16 rows x 128 columns, scales and zero-points at g*16 + r, codes at + * c*8 + r/2. A [32 x 256] tensor of 2 x 2 chunks, read into a 384-wide buffer (the forward pads K to 256). */ +static void q4h_chunks(void) { + const int rows = 32, k = 256, cols = 384, ntc = k / Q4NX_Q4H_COLS; + const size_t n = (size_t)(rows / Q4NX_Q4H_ROWS) * (size_t)ntc; + uint8_t* t = (uint8_t*)calloc(n, Q4NX_CHUNK_Q4H); + float* out = (float*)malloc(sizeof(float) * (size_t)rows * cols); + if (!t || !out) { check(0, "q4h: the buffers allocate"); free(t); free(out); return; } + for (size_t ci = 0; ci < n; ci++) { + uint8_t* c = t + ci * Q4NX_CHUNK_Q4H; + for (int i = 0; i < 64; i++) { + put_bf16(c, (size_t)i, 0.25f * (float)(1 + (int)((i + ci) % 7))); /* scale */ + put_bf16(c + 128, (size_t)i, -1.0f * (float)(1 + (int)((i + 3 * ci) % 5))); /* zero-point */ + } + for (int b = 0; b < 1024; b++) c[256 + b] = (uint8_t)((b * 7 + 3 * ci) & 0xff); + } + for (size_t i = 0; i < (size_t)rows * cols; i++) out[i] = 99.0f; + check(q4nx_dequant_tensor_q4h(t, n * Q4NX_CHUNK_Q4H, rows, k, cols, out) == 0, "q4h: 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 / 128, rr = r % 16, cc = col % 128; + const size_t ci = (size_t)tr * ntc + tc; + const int i = (cc / 32) * 16 + rr; + const uint8_t byte = (uint8_t)(((cc * 8 + rr / 2) * 7 + 3 * ci) & 0xff); + const int code = (rr % 2 == 0) ? (byte & 0xF) : (byte >> 4); + const float s = 0.25f * (float)(1 + (int)((i + ci) % 7)), z = -1.0f * (float)(1 + (int)((i + 3 * ci) % 5)); + if (got != (float)code * s + z) bad++; + } + check(bad == 0, "q4h: every element is code (at c*8 + r/2) * scale + zero (at g*16 + r), grid row-major"); + check(pad_bad == 0, "q4h: the columns past the container's K are written 0"); + check(q4nx_dequant_tensor_q4h(t, n * Q4NX_CHUNK_Q4H - 1, rows, k, cols, out) == -1, "q4h: one byte short is refused"); + check(q4nx_dequant_tensor_q4h(t, n * Q4NX_CHUNK_Q4H, rows, k, k - 128, out) == -1, "q4h: a buffer narrower than K is refused"); + check(q4nx_dequant_tensor_q4h(t, n * Q4NX_CHUNK_Q4H, 24, k, cols, out) == -1, "q4h: rows off the 16-row grid are refused"); + free(t); + free(out); +} + int main(void) { q4s_chunk(); + q4h_chunks(); tile_kinds(); tile_kind_is_per_tensor(); tile_row_kind_arms();