forward: read the hub Qwen3.5 packs' Q8 code order from the codes (refs #1786)

docs/npu.md

@@ -371,6 +371,20 @@ apart. Decoding tile (0, 0) of each lm_head and comparing with a GGUF of the sam
 gives correlation 0.9998 with the right layout and about 0.01 with the other. `npu/q4nx`
 decodes the first layout.
 
+The forward's decoder (`npu/forward/src/q4nx_dequant.c`) reads the layout from the codes
+themselves (`q4nx_q8_split16`, engine#1786). Each (row, 32-column group) of a symmetric Q8 tensor is
+scaled by its own absmax, so in the right layout every group that holds a nonzero code holds one
+of magnitude 127; read in the other layout the groups mix rows, and most miss it. Measured on 64
+chunks: 1.000 against 0.50-0.77 on Qwen3.5-0.8B-NPU2's `lm_head`, `ssm_alpha_proj` and
+`ssm_out_proj`, and the reverse on the local Qwen3.5-2B-NPU2. The second layout is chosen only
+when it alone holds every group at 127; otherwise the decoder keeps the first, as before. The
+hub's Qwen3.5 packs (`flm_version` 0.9.45: 0.8B, 2B, 4B, 9B) are in the second layout and also
+store `ssm_out_proj` as Q8. A FastFlowLM install rewrites a pack it fetches (the local 2B copy
+reports 1.0.3, stores `ssm_out_proj` as Q4_K and uses the first layout), so a row measured on
+`~/.config/flm/models` is not a row on the hub's files. Before this, Qwen3.5-0.8B-NPU2 from the
+hub answered "What is the capital of France?" with noise on the host and on the NPU. Its
+`blk.0.ssm_out` against the BF16 GGUF of the same checkpoint: corr 0.0054 before, 0.9993 after.
+
 `npu/q4nx.h` decodes all five (the two q4_1 parameter forms included), and repacks Q4_K into q4_1
 (exact apart from rounding `S * scale` and `M * min` to bf16), so code that reads q4_1 runs Q4_K
 weights unchanged. Q8 does

npu/forward/include/q4nx_dequant.h

@@ -169,6 +169,17 @@ void q4nx_dequant_tile_q4k(const uint8_t* chunk, float* out, int ld);
  *   q[8192] i8  @[512,8704) byte k*32 + r
  *   value(r,k) = d[g*32+r] * q[k*32+r] */
 void q4nx_dequant_tile_q8(const uint8_t* chunk, float* out, int ld);
+/* The same 8704-byte Q8 chunk as FastFlowLM's 0.9.x packs write it (the Qwen3.5 packs on the hub, flm_version 0.9.45):
+ * the scales are as above, but the 8192 codes are two 16-row halves, each column-major -- byte h*4096 + c*16 + r%16
+ * for row r = 16h + r%16 -- where the 1.0.x writer stores one 32-row column-major block (byte c*32 + r). The size is
+ * the same, so the order is read from the codes (q4nx_q8_split16). Settled on Qwen3.5-0.8B-NPU2 against the BF16 GGUF
+ * of the same checkpoint: blk.0.ssm_alpha/ssm_beta corr 0.99998, blk.0.ssm_out 0.9993 (the 32-row reading: 0.01). */
+void q4nx_dequant_tile_q8_split16(const uint8_t* chunk, float* out, int ld);
+/* Which order a run of 8704-byte Q8 chunks is in: 1 for split16, 0 for the 32-row block. A symmetric Q8 group is
+ * scaled by its own absmax, so in the right order every (row, 32-column group) with a nonzero code holds a code of
+ * magnitude 127; the wrong order mixes rows and most groups miss it (0.50 to 0.77 measured). split16 is chosen only
+ * when it alone holds every group at 127; anything else keeps the 32-row reading this decoder always used. */
+int q4nx_q8_split16(const uint8_t* data, size_t ntiles, size_t stride);
 
 /* Dequantize a whole [rows, cols] tensor whose tiles are `chunk_bytes` each
  * (5120 = q4_1, 4736 = Q4_K, 8704 = Q8).  Picks the matching tile decoder and

npu/forward/src/q4nx_dequant.c

@@ -137,6 +137,48 @@ void q4nx_dequant_tile_q8(const uint8_t* t, float* out, int ld) {
     }
 }
 
+void q4nx_dequant_tile_q8_split16(const uint8_t* t, float* out, int ld) {
+    const uint8_t* d_b = t;                      /* [256] bf16, index g*32+r */
+    const int8_t* q = (const int8_t*)(t + 512);  /* byte (r/16)*4096 + k*16 + r%16 */
+    for (int r = 0; r < Q4NX_TILE_ROWS; r++) {
+        const int8_t* qh = q + (size_t)(r / 16) * 4096 + (r % 16);
+        for (int k = 0; k < Q4NX_TILE_COLS; k++) {
+            const int g = k / Q4NX_GROUP_SIZE;
+            const float d = bf16_to_f32(le16(d_b + 2 * (g * Q4NX_TILE_ROWS + r)));
+            out[(size_t)r * ld + k] = d * (float)qh[k * 16];
+        }
+    }
+}
+
+/* Groups (row, 32 columns) of one chunk that hold a nonzero code, and those of them that reach |127|, read in
+ * `split` order (0: byte c*32 + r, 1: byte (r/16)*4096 + c*16 + r%16). */
+static void q8_group_hits(const uint8_t* t, int split, long* nonzero, long* hits) {
+    const int8_t* q = (const int8_t*)(t + 512);
+    for (int r = 0; r < Q4NX_TILE_ROWS; r++)
+        for (int g = 0; g < Q4NX_TILE_COLS / Q4NX_GROUP_SIZE; g++) {
+            int mx = 0;
+            for (int j = 0; j < Q4NX_GROUP_SIZE; j++) {
+                const int c = g * Q4NX_GROUP_SIZE + j;
+                const int v = split ? q[(size_t)(r / 16) * 4096 + (size_t)c * 16 + (r % 16)] : q[(size_t)c * 32 + r];
+                const int a = v < 0 ? -v : v;
+                if (a > mx) mx = a;
+            }
+            if (mx) { (*nonzero)++; if (mx >= 127) (*hits)++; }
+        }
+}
+
+int q4nx_q8_split16(const uint8_t* data, size_t ntiles, size_t stride) {
+    if (!data || ntiles == 0) return 0;
+    const size_t samples = ntiles < 64 ? ntiles : 64;   /* evenly spread; the order is one per tensor */
+    long nz[2] = {0, 0}, hit[2] = {0, 0};
+    for (size_t i = 0; i < samples; i++) {
+        const uint8_t* t = data + (i * ntiles / samples) * stride;
+        q8_group_hits(t, 0, &nz[0], &hit[0]);
+        q8_group_hits(t, 1, &nz[1], &hit[1]);
+    }
+    return nz[1] > 0 && hit[1] == nz[1] && hit[0] < nz[0];
+}
+
 void q4nx_dequant_tile_q4s(const uint8_t* t, float* out, int ld) {
     const uint8_t* sc_b = t;          /* [256] bf16, index g*32 + r */
     const uint8_t* qs = t + 512;      /* [256][16], byte c*16 + r/2 */
@@ -309,6 +351,7 @@ int q4nx_dequant_tensor_chunked(const uint8_t* data, size_t nbytes, int rows,
         sym = tensor_kind_or_refuse(data, (size_t)ntr * (size_t)ntc, Q4NX_TILE_BYTES);
         if (sym < 0) return -1;
     }
+    const int q8s = chunk_bytes == Q4NX_CHUNK_Q8 && q4nx_q8_split16(data, (size_t)ntr * (size_t)ntc, Q4NX_CHUNK_Q8);
     // Each tile writes a disjoint output region, so the tile grid parallelizes.
     // `tmp` is declared per-tile (inside the innermost body) so it is not shared.
     #pragma omp parallel for collapse(2) schedule(static) if(ntr * ntc >= 4)
@@ -319,7 +362,7 @@ int q4nx_dequant_tensor_chunked(const uint8_t* data, size_t nbytes, int rows,
             if (chunk_bytes == Q4NX_CHUNK_Q4K)
                 q4nx_dequant_tile_q4k(c, tmp, Q4NX_TILE_COLS);
             else if (chunk_bytes == Q4NX_CHUNK_Q8)
-                q4nx_dequant_tile_q8(c, tmp, Q4NX_TILE_COLS);
+                (q8s ? q4nx_dequant_tile_q8_split16 : q4nx_dequant_tile_q8)(c, tmp, Q4NX_TILE_COLS);
             else if (chunk_bytes == Q4NX_CHUNK_Q4S)
                 q4nx_dequant_tile_q4s(c, tmp, Q4NX_TILE_COLS);
             else
@@ -354,12 +397,15 @@ int q4nx_dequant_tile_row(const uint8_t* data, size_t nbytes, int rows, int cols
         sym = tensor_kind_or_refuse(data + (size_t)tr * (size_t)ntc * Q4NX_TILE_BYTES, (size_t)ntc, Q4NX_TILE_BYTES);
         if (sym < 0) return -1;
     }
+    /* Per token, so as with the 5120 kind above the order is read from this tile-row's own chunks. */
+    const int q8s = chunk_bytes == Q4NX_CHUNK_Q8 &&
+                    q4nx_q8_split16(data + (size_t)tr * (size_t)ntc * Q4NX_CHUNK_Q8, (size_t)ntc, Q4NX_CHUNK_Q8);
     for (int tc = 0; tc < ntc; tc++) {
         const uint8_t* c = data + (size_t)(tr * ntc + tc) * (size_t)chunk_bytes;
         if (chunk_bytes == Q4NX_CHUNK_Q4K)
             q4nx_dequant_tile_q4k(c, tmp, Q4NX_TILE_COLS);
         else if (chunk_bytes == Q4NX_CHUNK_Q8)
-            q4nx_dequant_tile_q8(c, tmp, Q4NX_TILE_COLS);
+            (q8s ? q4nx_dequant_tile_q8_split16 : q4nx_dequant_tile_q8)(c, tmp, Q4NX_TILE_COLS);
         else if (chunk_bytes == Q4NX_CHUNK_Q4S)
             q4nx_dequant_tile_q4s(c, tmp, Q4NX_TILE_COLS);
         else
@@ -437,6 +483,7 @@ int q4nx_dequant_tensor_packed(const uint8_t* data, size_t nbytes, int rows, int
         sym = tensor_kind_or_refuse(data, (size_t)ntr_p * (size_t)ntc_p, Q4NX_TILE_BYTES);
         if (sym < 0) return -1;
     }
+    const int q8s = chunk_bytes == Q4NX_CHUNK_Q8 && q4nx_q8_split16(data, (size_t)ntr_p * (size_t)ntc_p, Q4NX_CHUNK_Q8);
     // Read geometry is the container's grid (ntc_p is the tile-row stride); write geometry is the
     // model's, so a caller's buffer and its row stride stay exactly what they were before.
     #pragma omp parallel for collapse(2) schedule(static) if(ntr * ntc >= 4)
@@ -447,7 +494,7 @@ int q4nx_dequant_tensor_packed(const uint8_t* data, size_t nbytes, int rows, int
             if (chunk_bytes == Q4NX_CHUNK_Q4K)
                 q4nx_dequant_tile_q4k(c, tmp, Q4NX_TILE_COLS);
             else if (chunk_bytes == Q4NX_CHUNK_Q8)
-                q4nx_dequant_tile_q8(c, tmp, Q4NX_TILE_COLS);
+                (q8s ? q4nx_dequant_tile_q8_split16 : q4nx_dequant_tile_q8)(c, tmp, Q4NX_TILE_COLS);
             else if (chunk_bytes == Q4NX_CHUNK_Q4S)
                 q4nx_dequant_tile_q4s(c, tmp, Q4NX_TILE_COLS);
             else

npu/q4nx.h

@@ -53,6 +53,8 @@
 //     (k%32)*16 + r%16 with g = k/32                         (corr 0.9998; k*32 + r: 0.011)
 // dequant_tile implements the first. The 35B's q8 tensors are read by the private 35B route,
 // which implements the second. Do not run this decoder on a Q8 tensor of an unverified model.
+// The forward's decoder (npu/forward/src/q4nx_dequant.c) reads both and picks one per tensor
+// from the codes (q4nx_q8_split16: in the right layout every group's absmax code is 127).
 //
 // Checked bit for bit against 1bit-MONSTER's verified decoders (npu-infer/tools/
 // q4nx_dequant.py) on real Qwen3.5-4B and Qwen3-0.6B tiles: tests/npu_q4nx_test.cpp.

tests/npu_padded_grid_test.cpp

@@ -60,6 +60,61 @@ static std::vector<uint8_t> make_grid(int ntr, int ntc) {
     return g;
 }
 
+// engine#1786: the two code orders of the 8704-byte Q8 chunk (q4nx_dequant.h, q4nx_dequant_tile_q8_split16). A symmetric
+// Q8 group holds one code of magnitude 127 (its absmax), so each (row, 32-column group) here gets one +-127 and 31 codes
+// below it, with a bf16 scale of its own. The same values are written in each order: each must be detected as its own
+// order and decode to exactly the values written (bf16 scale x int8 code is exact in float).
+static uint16_t bf16_of(float f) { uint32_t u; std::memcpy(&u, &f, 4); return (uint16_t)(u >> 16); }
+static float f_of_bf16(uint16_t b) { const uint32_t u = (uint32_t)b << 16; float f; std::memcpy(&f, &u, 4); return f; }
+
+static void test_q8_split16() {
+    const int ntr = 2, ntc = 2, rows = ntr * Q4NX_TILE_ROWS, cols = ntc * Q4NX_TILE_COLS;
+    std::vector<float> want((size_t)rows * cols);
+    std::vector<uint8_t> old_order((size_t)ntr * ntc * Q4NX_CHUNK_Q8, 0), split((size_t)ntr * ntc * Q4NX_CHUNK_Q8, 0);
+    uint32_t seed = 1786;
+    auto rnd = [&] { seed = seed * 1664525u + 1013904223u; return seed >> 8; };
+    for (int tr = 0; tr < ntr; tr++)
+        for (int tc = 0; tc < ntc; tc++) {
+            uint8_t* to = &old_order[((size_t)tr * ntc + tc) * Q4NX_CHUNK_Q8];
+            uint8_t* ts = &split[((size_t)tr * ntc + tc) * Q4NX_CHUNK_Q8];
+            for (int r = 0; r < Q4NX_TILE_ROWS; r++)
+                for (int g = 0; g < 8; g++) {
+                    const uint16_t d = bf16_of(0.001f * (float)(1 + rnd() % 50));
+                    std::memcpy(to + 2 * (g * Q4NX_TILE_ROWS + r), &d, 2);
+                    std::memcpy(ts + 2 * (g * Q4NX_TILE_ROWS + r), &d, 2);
+                    const int peak = (int)(rnd() % 32);
+                    for (int j = 0; j < 32; j++) {
+                        const int k = g * 32 + j;
+                        int code = j == peak ? ((rnd() & 1) ? 127 : -127) : (int)(rnd() % 253) - 126;
+                        to[512 + k * 32 + r] = (uint8_t)(int8_t)code;
+                        ts[512 + (r / 16) * 4096 + k * 16 + (r % 16)] = (uint8_t)(int8_t)code;
+                        want[(size_t)(tr * 32 + r) * cols + (size_t)tc * 256 + k] = f_of_bf16(d) * (float)code;
+                    }
+                }
+        }
+    check(q4nx_q8_split16(split.data(), (size_t)ntr * ntc, Q4NX_CHUNK_Q8) == 1, "q8: split16 codes are detected as split16");
+    check(q4nx_q8_split16(old_order.data(), (size_t)ntr * ntc, Q4NX_CHUNK_Q8) == 0,
+          "q8: 32-row codes are detected as the 32-row order");
+    std::vector<float> got((size_t)rows * cols);
+    check(q4nx_dequant_tensor_chunked(split.data(), split.size(), rows, cols, Q4NX_CHUNK_Q8, got.data()) == 0 && got == want,
+          "q8: a split16 tensor decodes to the values written");
+    check(q4nx_dequant_tensor_chunked(old_order.data(), old_order.size(), rows, cols, Q4NX_CHUNK_Q8, got.data()) == 0 &&
+              got == want,
+          "q8: a 32-row tensor still decodes to the values written");
+    // The per-token row path reads the same order from its tile-row's chunks.
+    std::vector<float> row(cols);
+    bool rows_ok = true;
+    for (int r : {0, 15, 16, 31, 32, 63}) {
+        rows_ok &= q4nx_dequant_tile_row(split.data(), split.size(), rows, cols, Q4NX_CHUNK_Q8, r, row.data()) == 0 &&
+                   std::memcmp(row.data(), &want[(size_t)r * cols], (size_t)cols * sizeof(float)) == 0;
+    }
+    check(rows_ok, "q8: the row path decodes split16 rows on both 16-row halves");
+    // A tile that reaches 127 in both orders (constant codes) is not evidence for split16: the 32-row reading stays.
+    std::vector<uint8_t> flat(Q4NX_CHUNK_Q8);
+    make_q8_tile(flat.data(), 127);
+    check(q4nx_q8_split16(flat.data(), 1, Q4NX_CHUNK_Q8) == 0, "q8: a tile that fits both orders keeps the 32-row order");
+}
+
 int main() {
     const int rows = 2 * Q4NX_TILE_ROWS;    // 64 rows  -> 2 tile-rows
     const int cols = 2 * Q4NX_TILE_COLS;    // 512 cols -> 2 tile-cols
@@ -129,6 +184,8 @@ int main() {
     check(q4nx_packed_grid(7, 1536, 8960, &cb, &rp_g, &cp_g) == 0,
           "a size no chunk kind divides is refused");
 
+    test_q8_split16();
+
     if (failures) {
         std::printf("%d failure(s)\n", failures);
         return 1;