From 2058c7dffba5544545596fb08c27d4277e185047 Mon Sep 17 00:00:00 2001 From: bong-water-water-bong Date: Thu, 12 Nov 2026 16:06:27 -0400 Subject: [PATCH] forward: read the hub Qwen3.5 packs' Q8 code order from the codes (refs #1786) FastFlowLM's 8704-byte Q8 chunk has two code orders under one size: one 32-row column-major block (byte c*32 + r, what the decoder read) and two 16-row halves (byte (r/16)*4096 + c*16 + r%16). The hub's Qwen3.5 packs (flm_version 0.9.45) are in the second order and also store ssm_out_proj as Q8, so Qwen3.5-0.8B-NPU2 answered with noise on the host and on the NPU. q4nx_q8_split16() picks the order per tensor: a symmetric Q8 group is scaled by its absmax, so in the right order every nonzero (row, 32-column) group holds a code of magnitude 127. Measured 1.000 vs 0.50-0.77. split16 is chosen only when it alone fits; otherwise the 32-row reading stays. All three Q8 call sites (chunked, padded, per-row) use it. Co-Authored-By: Claude Opus 5.5 --- docs/npu.md | 14 ++++++++ npu/forward/include/q4nx_dequant.h | 11 ++++++ npu/forward/src/q4nx_dequant.c | 53 +++++++++++++++++++++++++-- npu/q4nx.h | 2 ++ tests/npu_padded_grid_test.cpp | 57 ++++++++++++++++++++++++++++++ 5 files changed, 134 insertions(+), 3 deletions(-) diff --git a/docs/npu.md b/docs/npu.md index 275f524d9..f7afab884 100644 --- a/docs/npu.md +++ b/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 diff --git a/npu/forward/include/q4nx_dequant.h b/npu/forward/include/q4nx_dequant.h index 64b464094..0232775cc 100644 --- a/npu/forward/include/q4nx_dequant.h +++ b/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 diff --git a/npu/forward/src/q4nx_dequant.c b/npu/forward/src/q4nx_dequant.c index 5c41767a2..c574bea83 100644 --- a/npu/forward/src/q4nx_dequant.c +++ b/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 diff --git a/npu/q4nx.h b/npu/q4nx.h index 160e8bbce..d0ad6f3be 100644 --- a/npu/q4nx.h +++ b/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. diff --git a/tests/npu_padded_grid_test.cpp b/tests/npu_padded_grid_test.cpp index fed2896a1..2413b305e 100644 --- a/tests/npu_padded_grid_test.cpp +++ b/tests/npu_padded_grid_test.cpp @@ -60,6 +60,61 @@ static std::vector 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 want((size_t)rows * cols); + std::vector 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 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 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 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;