From a951800b2d325ca71a2bf2cf0ca6169461e3c27c Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 15:43:36 -0400 Subject: [PATCH 01/12] fp8: CPU read path + matmul_fp8, as fmt=7 (public ordinal, #524) Ports fmt6/fp8-passthrough-r2@150f15b (the CPU read path -- quant.h's E4M3_LUT/e4m3_decode/matmul_fp8, colibri.c's qt_bytes/qt_scale_bytes/ qt_from_disk/qt_resolve_fmt fmt=1-vs-fp8 disambiguation) onto dev 9baae9b, renumbered fmt=100 -> fmt=7 throughout (code, tests, comments): the maintainer assigned fmt=7 as this format's public ordinal on #524, so this branch stops using the private-block number the format developed under. Mechanical rename per the PRIVATE ORDINAL BLOCK convention's own "find-and-replace, zero on-disk impact" promise -- qt_resolve_fmt still infers format from byte arithmetic alone; nothing on disk ever encoded the number 100 or the number 7. Two structural differences from the source commit, both scoped by the maintainer's #524 reply (PR 1 = CPU+repack+collision/refusal+Metal; registry+stamp move to a separate PR): - qt_resolve_fmt has NO stamped_name parameter and never will in this branch -- its signature is the plain 6-argument form every call site already uses (name,O,I,nb,ns,gs). A self-describing container stamp that could resolve a genuine byte-collision instead of refusing it is a follow-up proposal, called out as such in this function's own header comment and inside THE DESIGN LANDMINE's comment, not implemented here. - fmt=7's scale ENCODING is documented as a declared property of the format, not a hardcoded constant (QT struct comment + qt_resolve_fmt's new "SCALE ENCODING IS A DECLARED PROPERTY" comment). f32 (4 bytes/ block) is what this build implements. UE8M0 (1 byte/block, power-of-two exponent) is a REAL, distinct encoding the same weight geometry ships with -- DeepSeek-V4, per the maintainer's #524 finding -- recognized by its own byte signature (ns==ceil(O/128)*ceil(I/128)) inside the fmt=1-vs-fmt=7 landmine and refused BY NAME rather than silently misread as truncated/corrupt f32 or matched to the wrong candidate. Checked for collision against fmt=1's and fmt=7-f32's own ns arithmetic: realistically distinct for GLM-sized shapes, but not categorically -- the same small-O regime that already makes the f32 landmine possible also makes a ue8m0 collision possible (documented in-code, new test case, still refuses either way). tests/test_fp8_load.c's disambiguation suite (Part A) carries forward unchanged in substance (renumbered), plus a new Part A3 (test_ue8m0_scale_refusal) covering both the clean UE8M0 signature and the small-O ue8m0-vs-fmt=1 collision. Part C (check_fp8_bytes) carries forward: it exercises qt_bytes()/qt_scale_bytes(), CPU-read-path functionality added in this same commit, not the metadata stamp it was originally bundled with upstream in the source branch's history. Does NOT yet include: the fmt=6-vs-fmt=7 "SECOND DESIGN LANDMINE" collision reconciliation (source branch's final commit) -- qt_resolve_fmt here still has dev's unconditional fmt=6 early-return, same gap the source branch's first port commit deliberately left open. Follow-up commit on this branch closes it. Does NOT yet include the Metal kernel, the repack tool, or the end-to-end loader test -- following commits on this branch. FIX ROUND (Windows CI): test_fp8_load.c's expect_refuse() forked a child to catch qt_resolve_fmt's exit(1) refusals -- MinGW has no fork()/sys/wait.h. Gated / and expect_refuse's body behind #ifndef _WIN32, mirroring tests/test_st_pread.c's own established Windows arm for exactly this problem (fork-based exit(1) checks): the Windows arm prints an explicit "skipped on Windows (no fork): " line per skipped case (visible, never silent) and returns pass-through (1) rather than failing the build; every non-refusing case (expect_fmt and its callers) still runs and asserts for real on Windows. No product code touched by this fix -- test-file-only, per the CI failure's own diagnosis (fatal error on tests/test_fp8_load.c:51, compile-time only). RAN (this commit): make glm METAL=1 (clean rebuild) -- zero warnings. ./tests/test_fp8_passthrough -> "fp8 passthrough CPU tests: ok". ./tests/test_fp8_load -> "fp8 loader-seam tests: ok" (Part A + A3 + B + C, including both new ue8m0 refusal cases and the renumbered fmt=1-vs-fmt=7 suite; POSIX refusal path confirmed still running, not skipped, on this platform -- zero "skipped on Windows" lines in the output). Co-Authored-By: Claude Fable 5 --- c/Makefile | 8 +- c/colibri.c | 282 +++++++++++++++++++++++-- c/quant.h | 105 ++++++++++ c/tests/test_fp8_load.c | 369 +++++++++++++++++++++++++++++++++ c/tests/test_fp8_passthrough.c | 165 +++++++++++++++ 5 files changed, 912 insertions(+), 17 deletions(-) create mode 100644 c/tests/test_fp8_load.c create mode 100644 c/tests/test_fp8_passthrough.c diff --git a/c/Makefile b/c/Makefile index 9247ab81..02031cc8 100644 --- a/c/Makefile +++ b/c/Makefile @@ -199,7 +199,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_mirror$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_temp_env$(EXE) tests/test_tok_o200k$(EXE) tests/test_kv_alloc$(EXE) tests/test_int3$(EXE) tests/test_int3_load$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) tests/test_router_nan$(EXE) tests/test_pipe_block$(EXE) tests/test_e8_kernel$(EXE) tests/test_rans$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_mirror$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_temp_env$(EXE) tests/test_tok_o200k$(EXE) tests/test_kv_alloc$(EXE) tests/test_int3$(EXE) tests/test_int3_load$(EXE) tests/test_fp8_passthrough$(EXE) tests/test_fp8_load$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) tests/test_router_nan$(EXE) tests/test_pipe_block$(EXE) tests/test_e8_kernel$(EXE) tests/test_rans$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -478,6 +478,12 @@ tests/test_int3$(EXE): tests/test_int3.c colibri.c st.h uring.h json.h tok.h tok tests/test_int3_load$(EXE): tests/test_int3_load.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +tests/test_fp8_passthrough$(EXE): tests/test_fp8_passthrough.c quant.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + +tests/test_fp8_load$(EXE): tests/test_fp8_load.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h sample.h kv_persist.h telemetry.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + tests/test_logit_nan$(EXE): tests/test_logit_nan.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) diff --git a/c/colibri.c b/c/colibri.c index 7783036e..0a7445c7 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -100,7 +100,11 @@ typedef struct { * fmt=1 INT8 -> q8 (1 byte/param) + scala per riga * fmt=2 INT4 -> q4 (2 valori per byte, impacchettati) + scala per riga * INT4 e' cio' che fa stare la densa residente nei 15 GB (0.5 byte/param). */ -/* fmt: 0 F32, 1 INT8, 2 INT4 (2/byte), 3 INT2 (4/byte), 4 INT4-GROUPED, 5 INT3-G64. +/* fmt: 0 F32, 1 INT8, 2 INT4 (2/byte), 3 INT2 (4/byte), 4 INT4-GROUPED, 5 INT3-G64, + * 6 E8/IQ3 lattice, 7 FP8-E4M3 (native, passthrough -- see quant.h). fmt=7 is a + * PUBLIC ordinal, assigned by the maintainer on #524; it developed under the + * PRIVATE ORDINAL BLOCK convention below as fmt=100 and graduated out of that + * block into this list at the same time this comment was written. * q4 ospita int4/int2/int3 packed. fmt=4 (grouped int4, #242): per-row nibbles + one f32 * scale per group of `gs` inputs (s has O*ceil(I/gs) entries). * fmt=6 (E8/IQ3 lattice, #452): 98B per 256 weights = 3.0625 bits/weight, grid @@ -109,7 +113,53 @@ typedef struct { * fmt=5 (int3, per-GROUP scales, group=64, see quant.h I3_*): values in [-4,3] stored per * 64-input group as 24 bytes = 16B low plane (2 bits/val, int2 layout) + 8B high plane * (1 bit/val), plus ONE f32 scale PER GROUP (s has O*ceil(I/64) entries, not O). 3.5 - * bits/weight effective — the quality/size sweet spot measured in the #132 ablation. */ + * bits/weight effective — the quality/size sweet spot measured in the #132 ablation. + * fmt=7 (native FP8-e4m3 passthrough, resident/quality-core tier -- see quant.h's + * FP8_BLOCK/e4m3_decode/matmul_fp8): q8 holds O*I raw e4m3 bytes, ONE byte per weight, + * byte-identical layout to fmt=1's weight bytes (q4 is unused/NULL, same as fmt=1) -- + * disambiguated from fmt=1 (and, at small [O,I], from fmt=6 -- see below) purely by + * scale geometry, see qt_resolve_fmt's "THE DESIGN LANDMINE" comment further down. + * s holds ONE f32 scale PER 128x128 BLOCK, ceil(O/128)*ceil(I/128) entries total + * (row-major: block-row-major then block-col), NOT O and NOT O*ceil(I/gs) -- + * qt_bytes()/qt_scale_bytes() below are the authoritative byte-count formulas. The + * scale ENCODING is itself a declared PROPERTY of this format, not a hardcoded + * constant -- f32 (4 bytes/block, what's above) is the value THIS build + * implements; qt_resolve_fmt's "SCALE ENCODING IS A DECLARED PROPERTY" comment + * documents why (a DeepSeek-V4 checkpoint ships this identical weight geometry + * with a UE8M0 scale encoding instead) and how an unimplemented encoding is + * recognized and refused by name rather than silently misread. + * gs is unused (0) for fmt=7, same as fmt 1/2/3. */ +/* ---- PRIVATE ORDINAL BLOCK CONVENTION ------------------------------------ + * fmt values 0-7 are upstream-assigned, public, stable ordinals -- do not + * reuse or renumber them (fmt=7 is the newest member: assigned by the + * maintainer on #524, after developing under this block as fmt=100 -- see + * "graduated" above). fmt values 100+ remain this repo's PRIVATE/EXPERIMENTAL + * block for any OTHER in-flight format proposal: ordinals a branch mints for + * itself during development so it can't collide with a number upstream claims + * out from under it -- exactly what happened to fmt=7's OWN earlier private + * number, fmt=6, when #465's E8/IQ3 proposal claimed that ordinal upstream + * while this branch was still developing against it, forcing a re-mint to + * fmt=100 (see upstream_contribution/FORMATS_registry_draft.md for the + * incident that prompted the rule). + * + * A PRIVATE-BLOCK ordinal is an internal enum value only -- qt_resolve_fmt + * (below) infers format purely from byte arithmetic; the container on disk + * carries no format ordinal at all. (A self-describing container stamp that + * would persist a format's NAME, not its ordinal, is a follow-up proposal -- + * see qt_resolve_fmt's own note on where that plumbing would attach -- not + * present in this build.) Nothing outside this binary's own compiled code + * ever observes a 100+ number, so renumbering one later (e.g. when a format + * is upstreamed and assigned a real public ordinal at merge, as just happened + * for fmt=7) is a pure find-and-replace with zero on-disk or cross-version + * compatibility impact. + * + * Rule for adding a new format to this branch or a future one: claim the next + * unused 100+ integer, never a number already claimed upstream (check dev + * before picking one -- this is exactly the mistake that forced fmt=7's own + * proposal to renumber away from fmt=6) or by another in-flight private + * format. Never ship a 100+ ordinal as a public default/committed-upstream + * value -- the real ordinal is assigned by the maintainer at merge time, + * exactly as fmt=7's was. */ typedef struct { int fmt; float *qf; int8_t *q8; uint8_t *q4; float *s; int O, I, gs; /* gs=group size (0=per-row, 128=grouped) */ #ifdef COLI_CUDA @@ -131,8 +181,52 @@ static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ return (int64_t)t->O*ng*24 + (int64_t)t->O*ng*4; } if(t->fmt==6) /* E8/IQ3: 98B per 256 weights, scales in-block, .qs is a 4-byte tag */ return (int64_t)t->O*(((int64_t)t->I+255)/256)*98 + 4; + if(t->fmt==7){ /* fp8-e4m3 passthrough: O*I raw e4m3 bytes (n, byte-identical layout + * to fmt=1's weight bytes) + one f32 scale per 128x128 block + * (FP8_BLOCK in quant.h, included below qt_bytes -- keep the + * arithmetic literal here, same discipline as fmt=5's comment + * above). Missing this branch would fall through to the fmt=2 + * default below (packed-nibble formula, ~half the real weight + * bytes) and undercount a resident fp8 tensor's byte footprint -- + * feeds AUTOPIN/RAM-budget math, so this branch is load-bearing + * from day one, not a later fix (contrast fmt=6 above, upstream's + * own #452 review round 1 caught this exact bug shape for E8). */ + int64_t nblkO=((int64_t)t->O+127)/128, nblkI=((int64_t)t->I+127)/128; + return n + nblkO*nblkI*4; } return (int64_t)t->O*((t->I+1)/2) + (int64_t)t->O*4; /* fmt=2 int4 per-row */ } +/* scale-array byte count only, format-aware -- split out of qt_bytes() because + * qt_wire_mmap/qt_unwire_mmap mlock the weight and scale ranges as TWO SEPARATE + * regions (separate allocations, not one contiguous buffer: qt_from_disk always + * qalloc's t->q8/t->q4 and t->s independently) and need the scale byte count on + * its own to split qt_bytes(t) into a weight half and a scale half. The two call + * sites used to hardcode scale_b=O*4 -- i.e. assume PER-ROW scale for every + * format -- which is right for fmt 0/1/2/3 but wrong for fmt=4 (grouped, + * O*ceil(I/gs) scales), fmt=5 (int3-g64, O*ceil(I/64) scales), and fmt=7 + * (per-128x128-block, ceil(O/128)*ceil(I/128) scales): any tensor in one of + * those three formats reaching qt_wire_mmap/qt_unwire_mmap would mlock/munlock + * the wrong byte ranges on BOTH halves (weight_b = qt_bytes(t)-scale_b shifts + * too). fmt=6 (E8) is deliberately NOT covered here: its .qs is a fixed 4-byte + * tag with no wire-mmap-relevant weight/scale split of its own (matches + * qt_bytes' `+4` literal above), and t->fmt==6 tensors never reach + * qt_wire_mmap/qt_unwire_mmap's mlock path in this build (int3-g64/E8 stay + * CPU-side, see qt_cuda_upload -- MMAP wiring is a CUDA/Metal-adjacent memory + * concern this build doesn't extend to them). fixed generically since the same + * bug shape applies to fmt=4/5, not just the format that surfaced it -- fmt=4 + * is the routed-expert "int4-g64" format (qt_resolve_fmt assigns it to ESlot + * g/u/d the same as any other tensor, see expert_load_impl), so this is not + * purely a dormant fmt=7-only fix: any COLI_MMAP + mem_should_wire() session + * pinning fmt=4 experts was already mlocking the wrong ranges before this + * branch existed. UNVERIFIED at runtime (no model run performed -- static/ + * arithmetic fix only, see the report). */ +static int64_t qt_scale_bytes(const QT *t){ + if(t->fmt==4){ int ng=(t->I+t->gs-1)/t->gs; return (int64_t)t->O*ng*4; } + if(t->fmt==5){ int64_t ng=((int64_t)t->I+63)/64; return (int64_t)t->O*ng*4; } + if(t->fmt==7){ int64_t nblkO=((int64_t)t->O+127)/128, nblkI=((int64_t)t->I+127)/128; return nblkO*nblkI*4; } + return (int64_t)t->O*4; /* fmt 0 (t->s NULL, guarded by callers)/1/2/3/6: per-row (or, for + * fmt=6, the 4-byte tag -- callers of this function never see fmt=6, + * see the comment above; the fallback is harmless dead code for it). */ +} typedef struct { float *in_ln, *post_ln; @@ -300,6 +394,14 @@ static int g_cuda_e8_ready; /* codebook published to the devices (see cuda_boo static int qt_cuda_upload(QT *t){ if(t->fmt==5) return 0; /* int3-g64: no CUDA kernel yet — tensor stays CPU-side */ if(t->fmt==6 && !g_cuda_e8_ready) return 0; /* E8 without its codebook would decode garbage */ + if(t->fmt==7) return 0; /* fp8-e4m3 passthrough: no CUDA kernel yet either (matches the + * fmt=5/6 idiom above; matmul_qt_ex's own caller-side fmt exclusion + * already keeps this from being reached with cuda_eligible tensors + * in the exact-kernel path, but row_bytes()/weight_at() in + * backend_cuda.cu have no fmt=7 case either -- explicit early- + * return here so a future caller doesn't have to rediscover that by + * tracing through to the CUDA source. UNVERIFIED at runtime: no CUDA + * toolchain on this macOS build host to compile-check. */ const void *weights = t->fmt==0 ? (const void*)t->qf : t->fmt==1 ? (const void*)t->q8 : (const void*)t->q4; if(t->fmt==4) /* grouped int4 (#334): scales are [O, ceil(I/gs)] — the plain @@ -572,13 +674,26 @@ static void matmul_i4_grouped_pair(float *yg, float *yu, const float *x, * (~+12% perplexity), measured. Every other prefill matmul keeps IDOT as before. */ static void matmul_qt_ex(float *y, const float *x, QT *w, int S, int allow_idot){ #ifdef COLI_METAL + /* fmt=7 (fp8 passthrough) deliberately absent from this allowlist, same as fmt=5/6: + * the S>=g_metal_gemm_min batched prefill GEMM path (coli_metal_gemm) only knows + * fmt 1/2 in this build, so it fails CLOSED to the CPU branch below (matmul_fp8) + * rather than being silently misread. coli_metal_gemm() also carries its own + * internal fmt!=1&&fmt!=2 guard, so this is belt-and-braces. */ if(g_metal_enabled && S>=g_metal_gemm_min && !spec_pinned() && (w->fmt==1||w->fmt==2) && !omp_in_parallel()){ const void *wp = w->fmt==1 ? (const void*)w->q8 : (const void*)w->q4; if(coli_metal_gemm(y,x,wp,w->s,w->fmt,S,w->I,w->O)) return; } #endif #ifdef COLI_CUDA - if(g_cuda_enabled && w->cuda_eligible && !w->cuda_failed && w->fmt!=5 && !omp_in_parallel()){ + /* fmt=7 excluded exactly like fmt=5/6: the CUDA backend's row_bytes()/weight_at() + * (backend_cuda.cu) have no case for it and fall through to `return 0` / undefined + * weight_at() behavior. row_bytes()==0 already makes coli_cuda_tensor_upload_g + * refuse (defensive fail-closed), but that path is reached only after paying an + * upload attempt and printing a "disabled after an error" message every load; + * excluding it here up front is the same fmt=5/6 idiom, silent and immediate. + * UNVERIFIED on this build (no CUDA toolchain on macOS to compile-check; traced + * from backend_cuda.cu source only). */ + if(g_cuda_enabled && w->cuda_eligible && !w->cuda_failed && w->fmt!=5 && w->fmt!=7 && !omp_in_parallel()){ const void *weights = w->fmt==0 ? (const void*)w->qf : w->fmt==1 ? (const void*)w->q8 : (const void*)w->q4; if(coli_cuda_matmul(&w->cuda,y,x,weights,w->s,w->fmt,S,w->I,w->O,w->cuda_device,w->gs)) return; @@ -602,6 +717,7 @@ static void matmul_qt_ex(float *y, const float *x, QT *w, int S, int allow_idot) if(w->fmt==1) matmul_q(y,x,w->q8,w->s,S,w->I,w->O); else if(w->fmt==3) matmul_i2(y,x,w->q4,w->s,S,w->I,w->O); else if(w->fmt==5) matmul_i3(y,x,w->q4,w->s,S,w->I,w->O); + else if(w->fmt==7) matmul_fp8(y,x,(const uint8_t*)w->q8,w->s,S,w->I,w->O); else matmul_i4(y,x,w->q4,w->s,S,w->I,w->O); } @@ -1023,26 +1139,135 @@ static int detect_group_size(int O, int I, int64_t ns){ * diventava un int2 valido e il matmul leggeva oltre il buffer (O*I nibble a * 4/byte). Qui i byte del peso devono corrispondere a un layout noto e i byte * della scala alla cardinalita' attesa (O per-row, O*ng per-gruppo) — altrimenti - * si termina invece di sforare. Ritorna fmt (1/2/3/4/5) e scrive *gs. */ + * si termina invece di sforare. Ritorna fmt (1/2/3/4/5/6/7) e scrive *gs. + * No container-level format stamp is consulted here: a self-describing + * __metadata__ stamp that could resolve a genuine byte-collision below + * (instead of refusing it unconditionally) is a follow-up proposal, not + * present in this build -- every ambiguous shape this function can see + * refuses, full stop. */ static int qt_resolve_fmt(const char *name, int O, int I, int64_t nb, int64_t ns, int *gs){ int64_t exp_i8=(int64_t)O*I, exp_i4=(int64_t)O*((I+1)/2), exp_i2=(int64_t)O*((I+3)/4); int64_t exp_i3=(int64_t)O*i3_rowbytes(I); /* int3-g64 (fmt=5): 24B per 64-input group */ /* fmt=6 (E8/IQ3, #452): scales live inside the 98B super-blocks, so the .qs * convention is kept with a single-float tag — ns==4 is the discriminator - * (every other format carries at least O floats of real scales). */ - if(ns==4 && nb==(int64_t)O*e8_rowbytes(I)){ *gs=0; return 6; } + * (every other format carries at least O floats of real scales). + * + * SECOND DESIGN LANDMINE: e8_rowbytes(I) = ceil(I/256)*98 is the constant 98 + * for every I in (0,256], so this check's nb==O*e8_rowbytes(I) collapses to + * nb==O*98 -- which is ALSO fp8-e4m3-b128's raw-byte weight count (O*I) at + * the ONE value of I where I itself equals 98 (solving 98k==I for I in + * ((k-1)*256,k*256] only admits k=1, I=98). At that same I=98 a handful of + * fp8 scale-array shapes ALSO carry exactly ns==4 bytes, same as this + * check's own tag, and each is a genuine collision candidate this block + * must catch before the unconditional `return 6` below: + * - a SINGLE-BLOCK fp8 tensor (O<=128, fp8_nblk(O)*fp8_nblk(I)==1) with + * f32 block scales: ns==1*4==4. + * - a FOUR-BLOCK fp8 tensor (fp8_nblk(O)*fp8_nblk(I)==4, e.g. O in + * (384,512] at this I) with UE8M0 (1 byte/block) scales: ns==4*1==4 -- + * the SAME arithmetic collision the fmt=1-vs-fmt=7 landmine below has + * with UE8M0, just against fmt=6's tag instead of fmt=1's per-row + * count; see that landmine's own comment for the general shape. + * - at O==1 specifically, fmt=1's own per-row tag (O*4) is also 4. + * Refuse unconditionally when any of these coincide with the E8/IQ3 tag -- + * this build has no stamp to break the tie (see the function's own header + * comment) -- rather than let dev's own unconditional `return 6` silently + * misread an fp8-e4m3-b128 (or, at O==1, plain int8) tensor as E8/IQ3- + * lattice-decoded garbage with no error. No real GLM tensor has these + * shapes; the discipline exists for untrusted containers. */ + if(ns==4 && nb==(int64_t)O*e8_rowbytes(I)){ + int fp8_blk_f32_also = (nb==(int64_t)O*I) && (fp8_nblk(O)*fp8_nblk(I)==1); + int fp8_blk_ue8m0_also = (nb==(int64_t)O*I) && (fp8_nblk(O)*fp8_nblk(I)==4); + int i8_row_also = (nb==(int64_t)O*I) && (O==1); /* ns==O*4==4 iff O==1 */ + if(fp8_blk_f32_also || fp8_blk_ue8m0_also || i8_row_also){ + fprintf(stderr,"%s: [%d,%d] byte layout (nb=%lld ns=%lld) matches E8/IQ3 " + "(fmt=6, 4-byte tag)%s%s%s; refusing rather than guessing (untrusted " + "container, fmt=6 collision at I=98)\n", + name,O,I,(long long)nb,(long long)ns, + fp8_blk_f32_also ? " AND per-128x128-block FP8 f32 scales (fmt=7, single block)" : "", + fp8_blk_ue8m0_also ? " AND per-128x128-block FP8 ue8m0 scales (fmt=7, 4 blocks, recognized-not-implemented)" : "", + i8_row_also ? " AND plain int8 per-row (fmt=1, O=1)" : ""); + exit(1); + } + *gs=0; return 6; + } /* Row formats take precedence: for tiny I the int3-g64 byte count can coincide with * a row layout (e.g. [O,48]: ceil(48/2)=24=1*24). For real tensor shapes the counts * are distinct, and the weight bytes — not the scale size — are the int3 tag, because * int3-g64 and grouped-int4-at-gs=64 carry the SAME scale cardinality O*ceil(I/64). */ int fmt = (nb==exp_i8)?1 : (nb==exp_i4)?2 : (nb==exp_i2)?3 : (nb==exp_i3)?5 : 0; if(!fmt){ - fprintf(stderr,"%s: quantized weight is %lld bytes — no int8/int4/int2/int3-g64 layout for [%d,%d], refusing (untrusted container)\n", + fprintf(stderr,"%s: quantized weight is %lld bytes — no int8/int4/int2/int3-g64/fp8 layout for [%d,%d], refusing (untrusted container)\n", name,(long long)nb,O,I); exit(1); } *gs=0; if(fmt==2){ int g=detect_group_size(O,I,ns); if(g>0){ fmt=4; *gs=g; } } + /* fmt=1 vs fmt=7 (native FP8-e4m3 passthrough): THE DESIGN LANDMINE. Weight + * bytes are IDENTICAL (O*I raw bytes, both matched exp_i8 above) -- fmt=1 + * and fmt=7 can only be told apart by the SCALE array's geometry: fmt=1 is + * per-row (ns==O*4 bytes); fmt=7 is per-128x128-BLOCK, ns==ceil(O/128)* + * ceil(I/128)*4 bytes for THIS build's implemented f32 scale encoding. For + * most real shapes these two byte counts are distinct and the match is + * unambiguous. But for small O (<=128) and/or small I, ceil(O/128)* + * ceil(I/128) can equal O exactly (e.g. O=1,I<=128 -> 1*1=1=O; O=2,I in + * [129,256] -> 1*2=2=O; O=256,I in (16256,16384] -> 2*128=256=O) -- a + * tensor whose scale array satisfies BOTH conventions at once. Guessing + * either way risks silently reading a genuine per-row-int8 tensor as + * block-scaled FP8 (or vice versa), which would corrupt every weight + * without any error. Refuse instead (same "untrusted container" + * discipline as the rest of this function). + * + * SCALE ENCODING IS A DECLARED PROPERTY, not a hardcoded constant: f32 (4 + * bytes/block, above) is what THIS build implements, but it is not the + * only encoding real fp8-e4m3-b128 weight geometry ships with. DeepSeek-V4 + * ships the SAME weight layout (FP8 E4M3, 128x128 blocks) with UE8M0 + * scales instead -- one byte per block, a power-of-two exponent, dtype + * F8_E8M0 (maintainer finding, colibri #524) -- so a fp8-e4m3-b128 + * container's scale sidecar can legitimately be ceil(O/128)*ceil(I/128)*1 + * bytes, not *4. That byte count is a REAL, distinct signature (never + * equal to the f32 block-scale count above for any O,I>=1, since one is + * exactly 4x the other and 4n==n only at n==0) -- recognized here rather + * than silently misread as a truncated or corrupt f32 scale array. This + * build does not implement UE8M0 decode: recognizing the signature and + * refusing BY NAME (rather than falling through to the generic + * "wrong byte count" refusal below, or worse, matching it against the + * wrong candidate) is the whole point -- a future decoder lands into this + * seam rather than a near-duplicate format. Checked for collision against + * every OTHER format's ns arithmetic reachable from this nb==O*I branch + * (fmt=1 and fmt=7-f32, both above): the byte counts are realistically + * distinct (nblkO*nblkI is orders of magnitude smaller than O*4 for any + * real GLM-sized matrix), but NOT categorically distinct -- the same + * small-O regime that makes the fmt=1-vs-fmt=7-f32 collision above + * possible also makes a fmt=1-vs-fmt=7-ue8m0 collision possible (e.g. + * O=1, I in (384,512]: nblkO=1, nblkI=4, ue8m0 ns=4=O*4=fmt=1's own + * per-row count) -- handled below by refusing whenever the ue8m0 + * signature matches, noting the row collision too when it also applies. + * No real GLM tensor has these shapes; see also the SECOND DESIGN + * LANDMINE above for the analogous ue8m0-vs-fmt=6 corner this same + * signature can hit. */ + if(fmt==1){ + int64_t nblkO=fp8_nblk(O), nblkI=fp8_nblk(I); + int64_t ns_row=(int64_t)O*4, ns_blk=nblkO*nblkI*4, ns_blk_ue8m0=nblkO*nblkI; + int is_row=(ns==ns_row), is_blk=(ns==ns_blk), is_blk_ue8m0=(ns==ns_blk_ue8m0); + if(is_row && is_blk){ + fprintf(stderr,"%s: [%d,%d] scale array is %lld bytes — matches BOTH per-row " + "int8 (fmt=1) and per-128x128-block FP8 (fmt=7) scale geometry; refusing " + "rather than guessing (untrusted container, THE DESIGN LANDMINE)\n", + name,O,I,(long long)ns); exit(1); + } + if(is_blk_ue8m0){ + fprintf(stderr,"%s: [%d,%d] fp8-e4m3-b128 with ue8m0 scales recognized but not " + "implemented; only f32 block scales are supported in this build (nb=%lld " + "bytes matches raw e4m3 weight bytes, ns=%lld bytes matches %lld blocks x " + "1 byte/block)%s\n", + name,O,I,(long long)nb,(long long)ns,(long long)(nblkO*nblkI), + is_row ? " -- scale array ALSO matches per-row int8 (fmt=1); refusing either way (untrusted container)" : ""); + exit(1); + } + if(is_blk && !is_row) fmt=7; + } int64_t exp_scale = (fmt==4)? (int64_t)O*((I+*gs-1)/(*gs)) - : (fmt==5)? (int64_t)O*i3_groups(I) : (int64_t)O; /* in FLOAT */ + : (fmt==5)? (int64_t)O*i3_groups(I) + : (fmt==7)? fp8_nblk(O)*fp8_nblk(I) + : (int64_t)O; /* in FLOAT */ if(ns != exp_scale*4){ fprintf(stderr,"%s: scale array is %lld bytes — expected %lld for [%d,%d] fmt=%d, refusing (untrusted container)\n", name,(long long)ns,(long long)(exp_scale*4),O,I,fmt); exit(1); } @@ -1072,16 +1297,25 @@ static void qt_from_disk(Model *m, const char *name, int O, int I, int bits, int else if(fmt==6){ /* E8/IQ3: everything in-block, .qs is the 4-byte tag */ if(t->fmt!=6||!t->q4){ t->fmt=6; t->O=O; t->I=I; t->gs=0; t->q4=qalloc(nb); t->s=qsalloc(1); } st_read_raw(&m->S,name,t->q4,drop); } + else if(fmt==7){ int64_t nblk=fp8_nblk(O)*fp8_nblk(I); /* fp8-e4m3: raw bytes (q8, like fmt=1) + * + one f32 scale per 128x128 block. BOTH weights and scale via qalloc, not falloc, + * from day one -- the GPU-visibility lesson fmt=4 and fmt=5 each had to learn + * separately (see review round 1 audit history in the report). */ + if(t->fmt!=7||!t->q8){ t->fmt=7; t->O=O; t->I=I; t->gs=0; t->q8=qalloc(nb); t->s=(float*)qalloc((size_t)nblk*sizeof(float)); } + st_read_raw(&m->S,name,t->q8,drop); } else { if(t->fmt!=fmt||!t->q4){ t->fmt=fmt; t->O=O; t->I=I; t->gs=0; t->q4=qalloc(nb); t->s=qsalloc(O); } st_read_raw(&m->S,name,t->q4,drop); } /* cap MUST match the scale cardinality qt_resolve_fmt already validated and * the falloc above actually reserved, per format: grouped-int4 (fmt=4) keeps - * O*ceil(I/gs) scales and int3-g64 (fmt=5) keeps O*i3_groups(I); everything - * else is per-row O. Using the per-row bound for a grouped format would - * reject a legitimate container (fmt=5 regressed exactly that way). */ + * O*ceil(I/gs) scales, int3-g64 (fmt=5) keeps O*i3_groups(I), E8/IQ3 (fmt=6) + * keeps the 1-float tag, and fp8-e4m3 (fmt=7) keeps ceil(O/128)*ceil(I/128) + * block scales; everything else is per-row O. Using the per-row bound for a + * grouped/blocked format would reject a legitimate container (fmt=5 regressed + * exactly that way). */ st_read_f32_cap(&m->S,sn,t->s, fmt==4 ? (int64_t)O*((I+gs-1)/gs) : fmt==5 ? (int64_t)O*i3_groups(I) : - fmt==6 ? (int64_t)1 : (int64_t)O, drop); + fmt==6 ? (int64_t)1 : + fmt==7 ? fp8_nblk(O)*fp8_nblk(I) : (int64_t)O, drop); } else { if(!t->qf && !t->q8 && !t->q4) qt_alloc(t,O,I,bits); if(t->fmt==0) st_read_f32_cap(&m->S,name,t->qf,(int64_t)O*I,drop); @@ -2389,10 +2623,19 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p * (step_decode_batch) passes per-row kvs[]/positions[] with pos_base=0, so the kernel * would rope every row at position 0 and attend over a 1-token window of the wrong * cache -> greedy decode hits EOS at token 2 (mux answers truncated to 1 token). - * Ragged rows take the CPU absorb path below, which reads kvs[s]/positions[s]. */ + * Ragged rows take the CPU absorb path below, which reads kvs[s]/positions[s]. + * fmt!=7 guards (q_a/q_b/kv_a/o): coli_metal_attn_decode's WP_() macro and the + * mm_gemv shader it dispatches through (bind_gemv) have no fmt=7 case in this + * build -- an fp8-passthrough tensor reaching here would be silently misread as + * f32 by the shader's default branch. Fail closed to the CPU path below instead + * (matmul_qt_ex there dispatches fmt=7 correctly via matmul_fp8). kv_b is already + * pinned to fmt==2 above for an unrelated reason (its absorb kernel is int4-only); + * these four checks are the same discipline extended to the tensors that flow + * through the generic per-fmt shader. */ if(g_metal_enabled && !kvs && S<=4 && (g_absorb==1||(g_absorb<0&&S<=4)) && m->kv_start[layer]==0 && D==6144 && H==64 && c->q_lora==2048 && c->kv_lora==512 && c->qk_nope==192 - && c->qk_rope==64 && vh==256 && l->kv_b.fmt==2){ + && c->qk_rope==64 && vh==256 && l->kv_b.fmt==2 + && l->q_a.fmt!=7 && l->q_b.fmt!=7 && l->kv_a.fmt!=7 && l->o.fmt!=7){ int sel_active = m->has_dsa && layern_layers && c->idx_type[layer] && (pos_base+S) > c->index_topk; if(!sel_active){ if(m->has_dsa && layern_layers && c->idx_type[layer]){ /* index keys for future selection */ @@ -4227,12 +4470,19 @@ static void layer_forward_rows(Model *m, Layer *l, int li, float *x, int S, int * in un solo submit GPU; la CPU legge il routing e fa solo resolve/disk/expert-CB. * Fallback: qualsiasi condizione mancante -> percorso CPU intero qui sotto. * !kvs: ragged mux rows (per-row KV/position) are not expressible in this kernel's - * single Lc/Rc + pos_base contract — see the matching guard in attention_rows. */ + * single Lc/Rc + pos_base contract — see the matching guard in attention_rows. + * fmt!=7 guards (q_a/q_b/kv_a/o/sh_gate/sh_up/sh_down): same fail-closed discipline + * as attention_rows, extended to the shared-expert MLP this fused kernel also + * covers -- none of these seven bind_gemv-routed tensors has fmt=7 support in this + * build's mm_gemv shader, so any of them carrying fmt=7 must skip the whole fused + * path (CPU below dispatches fmt=7 correctly via matmul_qt_ex/matmul_fp8). */ if(g_metal_enabled && !kvs && S<=4 && lin_layers && l->sparse && (g_absorb==1||(g_absorb<0&&S<=4)) && m->kv_start[li]==0 && D==6144 && c->n_heads==64 && c->q_lora==2048 && c->kv_lora==512 && c->qk_nope==192 && c->qk_rope==64 && c->v_head==256 && l->kv_b.fmt==2 - && c->n_experts==256 && c->topk==8 && c->n_shared==1 && c->moe_inter==2048){ + && c->n_experts==256 && c->topk==8 && c->n_shared==1 && c->moe_inter==2048 + && l->q_a.fmt!=7 && l->q_b.fmt!=7 && l->kv_a.fmt!=7 && l->o.fmt!=7 + && l->sh_gate.fmt!=7 && l->sh_up.fmt!=7 && l->sh_down.fmt!=7){ int sel_active = m->has_dsa && c->idx_type[li] && (pos_base+S) > c->index_topk; if(!sel_active){ static float *linrm,*lnrm,*lsh,*lw; static int *lidx,*lkeff; diff --git a/c/quant.h b/c/quant.h index 7ff019f6..3f3aea47 100644 --- a/c/quant.h +++ b/c/quant.h @@ -402,6 +402,111 @@ static void matmul_i3(float *y, const float *x, const uint8_t *q3, const float * } } +/* ---- fmt=7: native FP8-e4m3 passthrough (Z.ai GLM-5.2-FP8 read path) ------ + * PUBLIC ordinal, assigned by the maintainer on #524. This format was minted + * fmt=6 during original development, then re-tagged fmt=100 (PRIVATE ORDINAL + * BLOCK, see colibri.c's QT struct comment) after #465 (E8/IQ3, above) claimed + * ordinal 6 upstream out from under it; fmt=7 is the number it graduated to at + * merge, per the private-block convention's own "find-and-replace, zero + * on-disk impact" promise. See qt_resolve_fmt in colibri.c ("THE DESIGN + * LANDMINE") for the disambiguation this needs against BOTH fmt=1 (int8) and + * fmt=6 (E8/IQ3). + * + * fmt=7 weight bytes are O*I raw e4m3 bytes -- byte-identical to int8 (fmt=1). + * What makes this a DIFFERENT format is the scale: one f32 per 128x128 BLOCK of + * the [O,I] weight matrix (shape [ceil(O/128),ceil(I/128)]), not one f32 per + * output row. Dequant is w[o,i] = e4m3_decode(byte) * scale[o/128, i/128] + * (MULTIPLY, not divide) -- mirrors tools/convert_fp8_to_int4.py's dequant() + * exactly, which is the authoritative reference for how Z.ai's checkpoints read + * on the source side. This f32 encoding is a declared PROPERTY of the format, + * not the only one it can carry -- DeepSeek-V4 ships this identical weight + * geometry with UE8M0 (1-byte, power-of-two) block scales instead; qt_resolve_fmt + * recognizes that byte signature and refuses it BY NAME rather than + * misreading it as f32 (UE8M0 decode itself is not implemented in this build). + * + * 256-entry decode LUT, cross-checked byte-for-byte against PyTorch's + * torch.float8_e4m3fn (the exact dtype safetensors reports for these shards): + * sign(1) exp(4,bias=7) mant(3), subnormal at exp==0, and the OCP E4M3-FN + * convention that exp==0xF is NOT reserved for infinity -- only mant==0x7 at + * exp==0xF is NaN (max finite is exp=0xF,mant=0x6 -> 448). A static compile-time + * table (not a lazily-initialized one) is used deliberately: matmul_fp8 below + * runs inside #pragma omp parallel for, and a lazy-init global would race under + * concurrent first use. + * + * NaN POLICY (decided, documented per the build spec): a NaN byte (0x7F/0xFF) + * decodes to a real IEEE NaN and is left to propagate through the dot product, + * exactly like any other source of a NaN weight (fmt=0/f32 tensors are never + * scrubbed either). The engine already has a dedicated, TESTED safety net for + * NaN reaching the sampler (argmax_v/dist_build, see tests/test_logit_nan.c: + * "degrade + diagnose, never silently corrupt") -- fmt=7 relies on that existing + * downstream net rather than adding a second, redundant weight-level scrub. */ +static const float E4M3_LUT[256] = { + 0x0.0p+0f,0x1.0000000000000p-9f,0x1.0000000000000p-8f,0x1.8000000000000p-8f,0x1.0000000000000p-7f,0x1.4000000000000p-7f,0x1.8000000000000p-7f,0x1.c000000000000p-7f, + 0x1.0000000000000p-6f,0x1.2000000000000p-6f,0x1.4000000000000p-6f,0x1.6000000000000p-6f,0x1.8000000000000p-6f,0x1.a000000000000p-6f,0x1.c000000000000p-6f,0x1.e000000000000p-6f, + 0x1.0000000000000p-5f,0x1.2000000000000p-5f,0x1.4000000000000p-5f,0x1.6000000000000p-5f,0x1.8000000000000p-5f,0x1.a000000000000p-5f,0x1.c000000000000p-5f,0x1.e000000000000p-5f, + 0x1.0000000000000p-4f,0x1.2000000000000p-4f,0x1.4000000000000p-4f,0x1.6000000000000p-4f,0x1.8000000000000p-4f,0x1.a000000000000p-4f,0x1.c000000000000p-4f,0x1.e000000000000p-4f, + 0x1.0000000000000p-3f,0x1.2000000000000p-3f,0x1.4000000000000p-3f,0x1.6000000000000p-3f,0x1.8000000000000p-3f,0x1.a000000000000p-3f,0x1.c000000000000p-3f,0x1.e000000000000p-3f, + 0x1.0000000000000p-2f,0x1.2000000000000p-2f,0x1.4000000000000p-2f,0x1.6000000000000p-2f,0x1.8000000000000p-2f,0x1.a000000000000p-2f,0x1.c000000000000p-2f,0x1.e000000000000p-2f, + 0x1.0000000000000p-1f,0x1.2000000000000p-1f,0x1.4000000000000p-1f,0x1.6000000000000p-1f,0x1.8000000000000p-1f,0x1.a000000000000p-1f,0x1.c000000000000p-1f,0x1.e000000000000p-1f, + 0x1.0000000000000p+0f,0x1.2000000000000p+0f,0x1.4000000000000p+0f,0x1.6000000000000p+0f,0x1.8000000000000p+0f,0x1.a000000000000p+0f,0x1.c000000000000p+0f,0x1.e000000000000p+0f, + 0x1.0000000000000p+1f,0x1.2000000000000p+1f,0x1.4000000000000p+1f,0x1.6000000000000p+1f,0x1.8000000000000p+1f,0x1.a000000000000p+1f,0x1.c000000000000p+1f,0x1.e000000000000p+1f, + 0x1.0000000000000p+2f,0x1.2000000000000p+2f,0x1.4000000000000p+2f,0x1.6000000000000p+2f,0x1.8000000000000p+2f,0x1.a000000000000p+2f,0x1.c000000000000p+2f,0x1.e000000000000p+2f, + 0x1.0000000000000p+3f,0x1.2000000000000p+3f,0x1.4000000000000p+3f,0x1.6000000000000p+3f,0x1.8000000000000p+3f,0x1.a000000000000p+3f,0x1.c000000000000p+3f,0x1.e000000000000p+3f, + 0x1.0000000000000p+4f,0x1.2000000000000p+4f,0x1.4000000000000p+4f,0x1.6000000000000p+4f,0x1.8000000000000p+4f,0x1.a000000000000p+4f,0x1.c000000000000p+4f,0x1.e000000000000p+4f, + 0x1.0000000000000p+5f,0x1.2000000000000p+5f,0x1.4000000000000p+5f,0x1.6000000000000p+5f,0x1.8000000000000p+5f,0x1.a000000000000p+5f,0x1.c000000000000p+5f,0x1.e000000000000p+5f, + 0x1.0000000000000p+6f,0x1.2000000000000p+6f,0x1.4000000000000p+6f,0x1.6000000000000p+6f,0x1.8000000000000p+6f,0x1.a000000000000p+6f,0x1.c000000000000p+6f,0x1.e000000000000p+6f, + 0x1.0000000000000p+7f,0x1.2000000000000p+7f,0x1.4000000000000p+7f,0x1.6000000000000p+7f,0x1.8000000000000p+7f,0x1.a000000000000p+7f,0x1.c000000000000p+7f,0x1.e000000000000p+7f, + 0x1.0000000000000p+8f,0x1.2000000000000p+8f,0x1.4000000000000p+8f,0x1.6000000000000p+8f,0x1.8000000000000p+8f,0x1.a000000000000p+8f,0x1.c000000000000p+8f, NAN, + -0x0.0p+0f,-0x1.0000000000000p-9f,-0x1.0000000000000p-8f,-0x1.8000000000000p-8f,-0x1.0000000000000p-7f,-0x1.4000000000000p-7f,-0x1.8000000000000p-7f,-0x1.c000000000000p-7f, + -0x1.0000000000000p-6f,-0x1.2000000000000p-6f,-0x1.4000000000000p-6f,-0x1.6000000000000p-6f,-0x1.8000000000000p-6f,-0x1.a000000000000p-6f,-0x1.c000000000000p-6f,-0x1.e000000000000p-6f, + -0x1.0000000000000p-5f,-0x1.2000000000000p-5f,-0x1.4000000000000p-5f,-0x1.6000000000000p-5f,-0x1.8000000000000p-5f,-0x1.a000000000000p-5f,-0x1.c000000000000p-5f,-0x1.e000000000000p-5f, + -0x1.0000000000000p-4f,-0x1.2000000000000p-4f,-0x1.4000000000000p-4f,-0x1.6000000000000p-4f,-0x1.8000000000000p-4f,-0x1.a000000000000p-4f,-0x1.c000000000000p-4f,-0x1.e000000000000p-4f, + -0x1.0000000000000p-3f,-0x1.2000000000000p-3f,-0x1.4000000000000p-3f,-0x1.6000000000000p-3f,-0x1.8000000000000p-3f,-0x1.a000000000000p-3f,-0x1.c000000000000p-3f,-0x1.e000000000000p-3f, + -0x1.0000000000000p-2f,-0x1.2000000000000p-2f,-0x1.4000000000000p-2f,-0x1.6000000000000p-2f,-0x1.8000000000000p-2f,-0x1.a000000000000p-2f,-0x1.c000000000000p-2f,-0x1.e000000000000p-2f, + -0x1.0000000000000p-1f,-0x1.2000000000000p-1f,-0x1.4000000000000p-1f,-0x1.6000000000000p-1f,-0x1.8000000000000p-1f,-0x1.a000000000000p-1f,-0x1.c000000000000p-1f,-0x1.e000000000000p-1f, + -0x1.0000000000000p+0f,-0x1.2000000000000p+0f,-0x1.4000000000000p+0f,-0x1.6000000000000p+0f,-0x1.8000000000000p+0f,-0x1.a000000000000p+0f,-0x1.c000000000000p+0f,-0x1.e000000000000p+0f, + -0x1.0000000000000p+1f,-0x1.2000000000000p+1f,-0x1.4000000000000p+1f,-0x1.6000000000000p+1f,-0x1.8000000000000p+1f,-0x1.a000000000000p+1f,-0x1.c000000000000p+1f,-0x1.e000000000000p+1f, + -0x1.0000000000000p+2f,-0x1.2000000000000p+2f,-0x1.4000000000000p+2f,-0x1.6000000000000p+2f,-0x1.8000000000000p+2f,-0x1.a000000000000p+2f,-0x1.c000000000000p+2f,-0x1.e000000000000p+2f, + -0x1.0000000000000p+3f,-0x1.2000000000000p+3f,-0x1.4000000000000p+3f,-0x1.6000000000000p+3f,-0x1.8000000000000p+3f,-0x1.a000000000000p+3f,-0x1.c000000000000p+3f,-0x1.e000000000000p+3f, + -0x1.0000000000000p+4f,-0x1.2000000000000p+4f,-0x1.4000000000000p+4f,-0x1.6000000000000p+4f,-0x1.8000000000000p+4f,-0x1.a000000000000p+4f,-0x1.c000000000000p+4f,-0x1.e000000000000p+4f, + -0x1.0000000000000p+5f,-0x1.2000000000000p+5f,-0x1.4000000000000p+5f,-0x1.6000000000000p+5f,-0x1.8000000000000p+5f,-0x1.a000000000000p+5f,-0x1.c000000000000p+5f,-0x1.e000000000000p+5f, + -0x1.0000000000000p+6f,-0x1.2000000000000p+6f,-0x1.4000000000000p+6f,-0x1.6000000000000p+6f,-0x1.8000000000000p+6f,-0x1.a000000000000p+6f,-0x1.c000000000000p+6f,-0x1.e000000000000p+6f, + -0x1.0000000000000p+7f,-0x1.2000000000000p+7f,-0x1.4000000000000p+7f,-0x1.6000000000000p+7f,-0x1.8000000000000p+7f,-0x1.a000000000000p+7f,-0x1.c000000000000p+7f,-0x1.e000000000000p+7f, + -0x1.0000000000000p+8f,-0x1.2000000000000p+8f,-0x1.4000000000000p+8f,-0x1.6000000000000p+8f,-0x1.8000000000000p+8f,-0x1.a000000000000p+8f,-0x1.c000000000000p+8f, NAN, +}; +static inline float e4m3_decode(uint8_t b){ return E4M3_LUT[b]; } + +#define FP8_BLOCK 128 +static inline int64_t fp8_nblk(int n){ return ((int64_t)n + FP8_BLOCK - 1) / FP8_BLOCK; } + +/* y[S,O] = x[S,I] @ W^T, W raw e4m3 bytes (byte-identical layout to fmt=1) + + * per-128x128-BLOCK f32 scale [ceil(O/128),ceil(I/128)]. Scalar reference path + * (no SIMD in v1 -- BW-bound like the Metal kernel, and this format's hot path + * is the GPU one; a vectorized CPU kernel is future work if measured needed). + * Mirrors matmul_i3's double-accumulate-across-groups / float-within-group + * convention so cross-block cancellation doesn't cost precision unfairly. */ +static void matmul_fp8(float *y, const float *x, const uint8_t *q8, const float *bscale, + int S, int I, int O){ + int64_t nblkI = fp8_nblk(I); + #pragma omp parallel for schedule(static) + for(int o=0;oI) blen=I-base; + float sc=scl[bi]; float acc=0; + for(int i=base;i +#include +#include +#ifndef _WIN32 +#include +#include +#endif + +static int fails = 0; +#define CHECK(c) do{ if(!(c)){ printf("FAIL %s:%d: %s\n", __FILE__, __LINE__, #c); fails++; } }while(0) + +static uint64_t rng = 0xFEEDFACE0DDBA11ull; +static float rndf(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; + return ((int64_t)(rng & 0xFFFFF) - 0x80000) / (float)0x80000; } +static uint8_t rndbyte_nonan(void){ + for(;;){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; + uint8_t b = (uint8_t)(rng & 0xFF); + if(b != 0x7F && b != 0xFF) return b; } +} + +/* ---- Part A: qt_resolve_fmt disambiguation (in-process for non-refusing + * cases, fork+waitpid for refusing ones -- qt_resolve_fmt exit(1)s in place, + * it does not return an error code). ---- */ + +static int expect_fmt(int O, int I, int64_t nb, int64_t ns, int expect_fmt_val, const char *tag){ + int gs=0; + int fmt = qt_resolve_fmt(tag, O, I, nb, ns, &gs); + if(fmt != expect_fmt_val){ + printf("FAIL %s: got fmt=%d, expected fmt=%d (O=%d I=%d nb=%lld ns=%lld)\n", + tag, fmt, expect_fmt_val, O, I, (long long)nb, (long long)ns); + return 0; + } + return 1; +} + +static int expect_refuse(int O, int I, int64_t nb, int64_t ns, const char *tag){ +#ifndef _WIN32 + int pipefd[2]; if(pipe(pipefd)!=0) return 0; + pid_t pid = fork(); + if(pid < 0) return 0; + if(pid == 0){ + dup2(pipefd[1],2); close(pipefd[0]); close(pipefd[1]); + int gs=0; + qt_resolve_fmt(tag, O, I, nb, ns, &gs); /* must exit(1) inside; must NOT return */ + _exit(42); /* reaching here is the bug */ + } + close(pipefd[1]); + char err[1024]={0}; ssize_t n=read(pipefd[0],err,sizeof(err)-1); (void)n; + close(pipefd[0]); + int status=0; waitpid(pid,&status,0); + int ok = WIFEXITED(status) && WEXITSTATUS(status)==1; + if(!ok){ + printf("FAIL %s: expected exit(1) refusal, got status=%d, stderr=%.200s\n", tag, status, err); + return 0; + } + if(!strstr(err,"refus")){ + printf("FAIL %s: exited(1) but message lacked a refusal explanation: %.200s\n", tag, err); + return 0; + } + return 1; +#else + /* fork/pipe/waitpid are POSIX -- mirrors tests/test_st_pread.c's own + * Windows arm: skip the exit(1) subprocess check, print an explicit + * (never silent) skip line per case, count it as passing rather than + * failing. Every non-refusing case in this file (expect_fmt and its + * callers) still runs and asserts for real on Windows. */ + printf("skipped on Windows (no fork): %s\n", tag); + (void)O; (void)I; (void)nb; (void)ns; + return 1; +#endif +} + +static void test_disambiguation(void){ + /* --- non-degenerate golden paths (unambiguous either way) --- */ + CHECK(expect_fmt(4096,4096,(int64_t)4096*4096,4096*4,1,"plain int8 4096x4096")); + /* spec's own worked example: [2048,6144] expert -> block scale [16,48] */ + CHECK(expect_fmt(2048,6144,(int64_t)2048*6144,16LL*48*4,7,"fp8 2048x6144 (spec example)")); + CHECK(expect_fmt(384,6144,(int64_t)384*6144,3LL*48*4,7,"fp8 384x6144 non-square block grid")); + + /* --- degenerate shapes: O<=128 makes nblkO==1, so ns_blk==nblkI*4 can + * equal ns_row==O*4 whenever nblkI==O. Exhaustive-in-spirit sweep of the + * boundary the design landmine describes. --- */ + CHECK(expect_refuse(1,1, 1, 4, "degenerate O=1 I=1 (nblkI=1=O)")); + CHECK(expect_refuse(1,128, 128, 4, "degenerate O=1 I=128 (nblkI=1=O, I at block edge)")); + CHECK(expect_refuse(2,256, 2LL*256, 8, "degenerate O=2 I=256 (nblkI=2=O)")); + CHECK(expect_refuse(6,768, 6LL*768, 24, "degenerate O=6 I=768 (nblkI=6=O)")); + CHECK(expect_refuse(128,16384,128LL*16384, 512, "degenerate O=128 I=16384 (nblkO=1,nblkI=128=O)")); + /* O>128 degenerate case: nblkO=2, need nblkI=O/2 -- O=256,I=16384 -> nblkI=128, 2*128=256=O */ + CHECK(expect_refuse(256,16384,256LL*16384, 1024, "degenerate O=256 I=16384 (nblkO=2,nblkI=128, product=O)")); + /* k=3: the ambiguity isn't a one-off O=256 coincidence -- it's structural for ANY + * O that's a multiple of 128 (nblkO=k), since nblkI==128 (I in (16256,16384]) always + * makes nblkO*nblkI == k*128 == O. One more multiple (O=384=3*128) confirms the + * condition generalizes past the k=2 worked example, not just a re-derivation. */ + CHECK(expect_refuse(384,16384,384LL*16384, 1536, "degenerate O=384 I=16384 (nblkO=3,nblkI=128, k=3, product=O)")); + + /* --- boundary-ADJACENT non-degenerate cases: one step past each + * degenerate case above, both interpretations now legitimately resolve. --- */ + CHECK(expect_fmt(1,129, 129, 4, 1, "adjacent O=1 I=129 as fmt=1 (ns=row)")); + CHECK(expect_fmt(1,129, 129, 8, 7, "adjacent O=1 I=129 as fmt=7 (ns=block, nblkI=2)")); + CHECK(expect_fmt(2,257, 2LL*257, 8, 1, "adjacent O=2 I=257 as fmt=1 (ns=row)")); + CHECK(expect_fmt(2,257, 2LL*257, 12, 7, "adjacent O=2 I=257 as fmt=7 (ns=block, nblkI=3)")); + + /* --- neither interpretation matches: garbage .qs size, must still refuse + * (the pre-existing generic-mismatch path, exercised through the fmt=7-aware + * function to confirm the new code didn't disturb it). --- */ + CHECK(expect_refuse(10,10, 100, 999, "garbage ns matches neither row nor block layout")); +} + +/* ---- Part A2: fmt=6 (E8/IQ3, upstream #465, merged into dev) vs fmt=7 + * (this branch's fp8-e4m3-b128) collision -- SECOND DESIGN LANDMINE, see the + * derivation in qt_resolve_fmt's own comment. e8_rowbytes(I) is the constant + * 98 for every I in (0,256], so dev's fmt=6 tag check (ns==4 && + * nb==O*e8_rowbytes(I)) collapses to nb==O*98 -- which coincides with + * fp8-e4m3-b128's raw weight bytes (O*I) at the ONE value I==98, where a + * SINGLE-BLOCK (O<=128) fp8 tensor with f32 block scales, OR a FOUR-BLOCK + * fp8 tensor with ue8m0 (1 byte/block) scales, ALSO carries exactly ns==4, + * same as the E8 tag. An unstamped [I=98] fp8-e4m3-b128 tensor at either of + * those shapes is therefore byte-for-byte indistinguishable from a genuine + * fmt=6 tensor: nb AND ns both coincide, not just ns (contrast the fmt=1/ + * fmt=7 collision in Part A, where only ns ever coincides). O==1 stacks a + * THIRD candidate: fmt=1's per-row ns (O*4) is also 4 there. This build has + * no stamp to resolve any of these -- every one refuses. */ +static void test_fmt6_fp8_collision(void){ + /* O=64, I=98: nblkO=nblkI=1 (fmt=7, single block, f32 scales) -> ns=4; + * e8_blocks(98)=1 -> nb=O*98=6272 for BOTH interpretations, and fmt=6's + * .qs tag is always exactly one f32 -> ns=4 too. Must refuse. */ + int64_t nb64=(int64_t)64*98; + CHECK(expect_refuse(64,98, nb64, 4, "fmt=6/fmt=7(f32) collision O=64 I=98")); + /* boundary O=128 variant: still nblkO=1 for fmt=7 (128<=128), same collision. */ + int64_t nb128=(int64_t)128*98; + CHECK(expect_refuse(128,98, nb128, 4, "fmt=6/fmt=7(f32) collision O=128 I=98")); + + /* O=1, I=98: a THIRD candidate stacks on (fmt=1 plain int8 per-row, ns==O*4==4 + * too) -- a genuine three-way ambiguity. Must refuse. */ + int64_t nb1=(int64_t)1*98; + CHECK(expect_refuse(1,98, nb1, 4, "fmt=1/fmt=6/fmt=7(f32) THREE-way collision O=1 I=98")); + + /* O in (384,512], I=98: nblkO=4, nblkI=1, product=4 -- a fmt=7 tensor with + * UE8M0 (1 byte/block) scales also lands at ns==4*1==4 here, the SAME tag + * fmt=6 uses. Must refuse (and, since this build doesn't implement ue8m0 + * decode at all, would refuse on that basis alone even without the fmt=6 + * collision -- this case exercises the OUTER fmt=6 guard specifically). */ + int64_t nb400=(int64_t)400*98; + CHECK(expect_refuse(400,98, nb400, 4, "fmt=6/fmt=7(ue8m0, 4-block) collision O=400 I=98")); + + /* regression guard: a GENUINE (non-colliding) fmt=6 fixture -- I!=98, so + * e8_rowbytes(I)==98 does NOT equal O*I -- must keep resolving to fmt=6. + * Mirrors test_e8_kernel.c's own O=24,I=512 shape and a small + * single-super-block shape (I=256, inside the (0,256] range where + * e8_rowbytes(I)==98 but I!=98, so still non-colliding). */ + CHECK(expect_fmt(24,512, (int64_t)24*e8_rowbytes(512), 4, 6, + "genuine fmt=6 (non-colliding) O=24 I=512 -> still resolves to fmt=6")); + CHECK(expect_fmt(64,256, (int64_t)64*e8_rowbytes(256), 4, 6, + "genuine fmt=6 (non-colliding) O=64 I=256 (I!=98, no collision) -> fmt=6")); +} + +/* ---- Part A3: fmt=7's scale ENCODING is a declared property -- UE8M0 + * recognized, refused by name (not implemented in this build). ---- */ +static void test_ue8m0_scale_refusal(void){ + /* [2048,6144] (spec example shape, same as Part A's fmt=7/f32 golden + * path): nblkO=16, nblkI=48, product=768 blocks. A UE8M0 sidecar is + * exactly 1 byte/block -> ns=768, distinct from BOTH fmt=1's per-row + * count (O*4=8192) and this build's f32 block-scale count (768*4=3072). + * Clean, unambiguous UE8M0 signature -- must name-refuse, not silently + * treat it as a truncated/corrupt f32 array or match it to fmt=1. */ + int64_t nb=(int64_t)2048*6144; + CHECK(expect_refuse(2048,6144, nb, 768, + "fp8-e4m3-b128 with ue8m0 scales (spec-shaped, non-degenerate) -> recognized, refused by name")); + + /* O=1, I=400: nblkO=1, nblkI=ceil(400/128)=4, product=4 -- a UE8M0 sidecar + * here is ns=4*1=4, which ALSO equals fmt=1's per-row count (O*4=4): the + * same small-O regime that produces the f32-vs-fmt=1 collision in Part A + * produces a ue8m0-vs-fmt=1 collision too. Must still refuse (combined + * message), not silently pick fmt=1. */ + int64_t nb1=(int64_t)1*400; + CHECK(expect_refuse(1,400, nb1, 4, + "fp8-e4m3-b128 with ue8m0 scales, ALSO colliding with fmt=1 per-row (O=1) -> refused")); +} + +/* ---- Part B: qt_from_disk loader-seam (real safetensors file) ---- */ + +static void deq_fmt7(const QT *t, float *dq){ + int64_t nblkI = fp8_nblk(t->I); + for(int o=0;oO;o++){ + int64_t blkO = o/FP8_BLOCK; const float *scl = t->s + blkO*nblkI; + for(int i=0;iI;i++){ + int64_t bi = i/FP8_BLOCK; + dq[(int64_t)o*t->I+i] = e4m3_decode(t->q8[(int64_t)o*t->I+i]) * scl[bi]; + } + } +} +static void deq_fmt1(const QT *t, float *dq){ + for(int o=0;oO;o++){ float s=t->s[o]; + for(int i=0;iI;i++) dq[(int64_t)o*t->I+i]=(float)t->q8[(int64_t)o*t->I+i]*s; } +} + +#define CDIV(n,d) (((n)+(d)-1)/(d)) + +static void test_loader_seam(void){ + enum { O7=8, I7=256 }; /* nblkO=1, nblkI=2 -> 2 block scales total */ + enum { O1=5, I1=64 }; /* DIFFERENT shape from the fp8 tensor: O1*I1=320 + * bytes, ns=O1*4=20 -- neither collides with the + * fp8 tensor's own byte counts (kept deliberately + * distinct so this is a plain, non-degenerate + * negative control, not another landmine case). */ + enum { NBLK7 = CDIV(O7,128) * CDIV(I7,128) }; /* must be a compile-time constant expression for + * the static array below -- fp8_nblk() is a real + * function (runtime, not constexpr), so it can't + * size a `static` array even with literal inputs. */ + static uint8_t q7[O7*I7]; static float s7[NBLK7]; + for(int i=0;i real IEEE NaN -> flows through the dot product; see quant.h's + * e4m3_decode/E4M3_LUT comment for the documented policy + rationale). + * Pure quant.h test: no Model/QT/disk dependency (see quant.h's own "pure + * compute" header comment) -- the loader-seam (qt_resolve_fmt disambiguation, + * qt_from_disk) is covered separately in test_fp8_load.c. + * + * fmt=7, PUBLIC ordinal assigned by the maintainer on #524: this format was + * minted fmt=6 during original development of this branch, before dev's own + * #465 (E8/IQ3) claimed that ordinal upstream; re-tagged fmt=100 (PRIVATE + * ORDINAL BLOCK, see colibri.c's QT struct comment) from that point forward -- + * there was never a build in this branch's history where this format was + * reachable as fmt=6 -- and graduated to fmt=7 at merge. */ +#include "../quant.h" +#include +#include +#include + +static int fails = 0; +#define CHECK(c) do{ if(!(c)){ printf("FAIL %s:%d: %s\n", __FILE__, __LINE__, #c); fails++; } }while(0) + +static uint64_t rng = 0xC0FFEE1234567ull; +static float rndf(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; + return ((int64_t)(rng & 0xFFFFF) - 0x80000) / (float)0x80000; } +/* random byte EXCLUDING the two NaN codes (0x7F/0xFF) -- used by the accuracy/ + * magnitude tests below so a stray NaN term doesn't turn `rel>tol` trivially + * false (NaN compares false against everything) and mask a real kernel bug. + * NaN handling itself is checked separately, deliberately, below. */ +static uint8_t rndbyte_nonan(void){ + for(;;){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; + uint8_t b = (uint8_t)(rng & 0xFF); + if(b != 0x7F && b != 0xFF) return b; } +} + +/* independent reference decode (bit manipulation, not the LUT under test) -- + * cross-checked offline against torch.float8_e4m3fn byte-for-byte (256/256 + * match, RAN evidence, see the build report). Kept here so a transcription + * bug in E4M3_LUT's 256 hex-float literals is still caught by this test even + * though both were derived from the same formula. */ +static float ref_e4m3(uint8_t b){ + unsigned sign=(b>>7)&1, exp=(b>>3)&0xF, mant=b&0x7; + if(exp==0xF && mant==0x7) return NAN; + float val; + if(exp==0) val = (float)mant * (1.0f/8.0f) * powf(2.0f, 1.0f-7.0f); + else val = (1.0f + (float)mant*(1.0f/8.0f)) * powf(2.0f, (float)exp-7.0f); + return sign ? -val : val; +} + +static void test_lut(void){ + int nan_codes=0; + for(int b=0;b<256;b++){ + float got = e4m3_decode((uint8_t)b); + float ref = ref_e4m3((uint8_t)b); + if(isnan(ref)){ CHECK(isnan(got)); nan_codes++; continue; } + CHECK(got == ref); + if(got==0.f) CHECK(signbit(got)==signbit(ref)); /* +-0 must keep its sign bit */ + } + CHECK(nan_codes==2); /* exactly 0x7F and 0xFF */ + CHECK(isnan(e4m3_decode(0x7F)) && isnan(e4m3_decode(0xFF))); + CHECK(e4m3_decode(0x7E)==448.0f && e4m3_decode(0xFE)==-448.0f); /* max finite, OCP E4M3FN */ + CHECK(e4m3_decode(0x01)==0x1.0p-9f); /* min positive subnormal = 2^-9 */ + CHECK(e4m3_decode(0x00)==0.0f && !signbit(e4m3_decode(0x00))); + CHECK(e4m3_decode(0x80)==0.0f && signbit(e4m3_decode(0x80))); /* -0.0 */ +} + +/* double-precision block-scale reference: mirrors matmul_fp8's formula exactly + * (w[o,i] = e4m3_decode(byte) * scale[o/128,i/128]) but every intermediate is + * double so it can serve as ground truth for matmul_fp8's own float/double mix. */ +static void ref_matmul_fp8(double *y, double *mag, const float *x, const uint8_t *q8, + const float *bscale, int S, int I, int O){ + int64_t nblkI = fp8_nblk(I); + for(int o=0;o1e-30 ? d/mag[i] : d; + if(rel>worst) worst=rel; + if(rel>tol) bad++; + } + if(bad) printf(" %-40s worst_rel=%.3e FAIL (%d/%d over tol)\n", tag, worst, bad, n); + return bad==0; +} + +static void run_block_case(int O, int I, const char *tag){ + int64_t nblkO=fp8_nblk(O), nblkI=fp8_nblk(I), nblk=nblkO*nblkI; + uint8_t *q8 = malloc((size_t)O*I); + float *bscale = malloc((size_t)nblk*sizeof(float)); + float *x = malloc((size_t)3*I*sizeof(float)); + float *y = malloc((size_t)3*O*sizeof(float)); + double *yr = malloc((size_t)3*O*sizeof(double)); + double *mag = malloc((size_t)3*O*sizeof(double)); + for(int64_t i=0;i<(int64_t)O*I;i++) q8[i]=rndbyte_nonan(); + /* distinct, easily-distinguishable scale per block (b+1)*0.01 with alternating + * sign and a random jitter -- a swapped blkO/blkI stride or a wrong nblkI in the + * indexing picks up the WRONG block's scale, which this makes numerically loud + * rather than "close enough to pass by luck" (same idea as the GPU stride-audit + * case, applied here to the CPU kernel). */ + for(int64_t b=0;b that OUTPUT ROW is NaN (the + * documented policy: propagate, rely on the existing argmax_v/dist_build net + * downstream -- see quant.h's e4m3_decode comment). Rows that never touch the + * NaN byte stay finite: NaN contamination is per-dot-product, not global. + * O=3<=128 -> nblkO=1, so all three rows share the SAME nblkI=2-entry scale + * block-row (bscale[0..1]); q8 is the only thing that differs per row. */ + enum { O=3, I=200 }; + static uint8_t q8[O*I]; + static float bscale[2]; /* nblkO(1)*nblkI(2) */ + for(int i=0;i Date: Wed, 22 Jul 2026 15:47:07 -0400 Subject: [PATCH 02/12] fp8: Metal mm_gemv kernel for fmt=7 native FP8-e4m3 passthrough Ports fmt6/fp8-passthrough-r2@7baaada onto this branch's fmt=7 renumbering (previous commit): mm_gemv shader's fmt==7 branch (in-kernel e4m3 bit decode, OCP E4M3-FN NaN policy matching quant.h's e4m3_decode exactly; block-scale indexed scale[(o/128)*nblkI+i/128] and folded into acc), fmt_bytes()/fmt_scale_bytes() fmt==7 cases, coli_metal_matmul's explicit allow-list entry (fmt>3 && fmt!=7, since 7 is no longer adjacent to the public 0-3 range this file otherwise supports), all renumbered from the source commit's fmt=100. coli_metal_gemm/bind_gemv unchanged: fmt=7 was never in scope for either (matches matmul_qt_ex's CPU-dispatch exclusion from the previous commit). tests/test_backend_metal.mm: the fmt=7 GPU test group (own bit- manipulation CPU reference decode, magnitude-relative block-scale oracle, 7 shapes incl. block-edge/degenerate/non-square-stride-audit cases, an exhaustive 256-code LUT-exactness check run through the actual GPU kernel, and a gate-off check confirming coli_metal_gemm refuses fmt=7/returns 0), renumbered. Also carries forward run_fp8_moe_gate (originally landed alongside the source branch's metadata-stamp commit, but content-wise it is pure Metal-gate regression coverage with no stamp dependency: proves BY TEST, with deliberately-invalid 0xdeadbeef weight/ scale pointers that must never be dereferenced, that coli_metal_moe_block/moe_submit's existing `fmt != 1 && fmt != 2` gate already refuses fmt=7 before MB_BUILD's pointer-selection ternary could submit the wrong pointer -- in scope for this PR's "Metal kernel + all its tests", independent of the registry/stamp PR). RAN (this commit): make glm METAL=1 (clean rebuild) -- zero warnings. make metal-test under COLI_METAL_RESSET=0: full suite green, exit 0, including all 11 fp8 GPU lines (LUT 0/256 mismatches; 7 shape cases worst_rel in [3e-8, 7e-8]; GEMM gate rc=0; MB_BUILD/moe_submit gate rc=0) and every pre-existing case (int8/int4/int2/f32, moe_block, gemm S=64, fused attention, top-8 select) unaffected. make metal-test under COLI_METAL_RESSET=1: same, plus "[METAL] residency-set: on (macOS 15+, moe_submit skips per-buffer useResource:)" confirmed active. Co-Authored-By: Claude Fable 5 --- c/backend_metal.h | 11 ++- c/backend_metal.mm | 49 ++++++++++- c/tests/test_backend_metal.mm | 154 +++++++++++++++++++++++++++++++++- 3 files changed, 207 insertions(+), 7 deletions(-) diff --git a/c/backend_metal.h b/c/backend_metal.h index 35c4931c..bffb349b 100644 --- a/c/backend_metal.h +++ b/c/backend_metal.h @@ -27,9 +27,14 @@ void coli_metal_stats(size_t *tensor_count, size_t *tensor_bytes); int coli_metal_mem_info(size_t *used_bytes, size_t *total_bytes); /* - * y[S,O] = (x[S,I] @ W[O,I]^T) * scale[o]. - * fmt matches QT in glm.c: 0=f32, 1=int8, 2=int4(packed), 3=int2(packed). - * The first successful call wraps W and its row scales in GPU-visible buffers; + * y[S,O] = (x[S,I] @ W[O,I]^T) * scale[o]. fmt=7 (fp8 passthrough -- see + * colibri.c) instead folds a per-128x128-block scale into the accumulation + * -- see the shader comment in backend_metal.mm. + * fmt matches QT in colibri.c: 0=f32, 1=int8, 2=int4(packed), 3=int2(packed), + * 7=fp8-e4m3 (one raw byte/element, same layout as fmt=1) + per-128x128- + * block scale (scale array is [ceil(O/128),ceil(I/128)] floats; no group-size + * parameter needed -- the block is a fixed 128x128, not caller-configurable). + * The first successful call wraps W and its scales in GPU-visible buffers; * later calls reuse them (weights are assumed stable at the same address). * Returns 1 on success, 0 if Metal is unavailable or fmt is invalid. */ diff --git a/c/backend_metal.mm b/c/backend_metal.mm index b894eb06..d0f9c363 100644 --- a/c/backend_metal.mm +++ b/c/backend_metal.mm @@ -9,6 +9,16 @@ // ---- shader: general quantized GEMV, one threadgroup per output element (o,si) ---- // y[si,o] = (sum_i dequant(W[o,i]) * x[si,i]) * scale[o]. fmt: 0=f32 1=i8 2=i4 3=i2. +// fmt=7 (native FP8-e4m3 passthrough -- see colibri.c): raw byte +// layout identical to fmt=1 (one e4m3 byte per element, no packing), but the scale is +// per-128x128 BLOCK of [O,I] -- buffer(1) holds [ceil(O/128),ceil(I/128)] floats indexed +// scale[(o/128)*nblkI+i/128]. Folded into acc like a per-group scale would be, for the +// same reason: it is not constant across one output row. e4m3->float is bit manipulation +// in-kernel (sign/exp/mantissa, OCP E4M3-FN: exp==0 subnormal, exp==0xF&&mant==0x7 is the +// only NaN code) -- BW-bound kernel, decode ALU is free, no LUT texture. Must byte-match +// quant.h's E4M3_LUT/e4m3_decode (the CPU reference) including the NaN policy, which is +// why the metal-test suite runs all 256 byte codes through this exact kernel path (not +// just spot values) -- see run_fp8_lut(). static const char *SHADER = R"METAL( #include using namespace metal; @@ -53,6 +63,25 @@ kernel void mm_gemv(device const uchar* w [[buffer(0)]], // raw weight by for (int i = slane; i < I; i += 32) { uchar b = wr[i>>2]; int v = (b >> (2*(i&3))) & 0x3; acc += float(v-2) * xr[i]; } + } else if (fmt == 7) { // fp8 e4m3 passthrough: one raw byte per + // element (like fmt=1), scale per 128x128 + // block folded into acc (like a grouped fmt). + int nblkI = (I + 127) / 128; + device const uchar* wr = w + (long)o * I; + device const float* scl = scale + (long)(o/128) * nblkI; + for (int i = slane; i < I; i += 32) { + uchar b = wr[i]; + uint sign = b >> 7, exp = (b >> 3) & 0xF, mant = b & 0x7; + float wv; + if (exp == 0xF && mant == 0x7) { + wv = as_type(0x7fc00000u); // qNaN -- matches quant.h's e4m3_decode + } else { + float mag = (exp == 0) ? (float(mant) * 0.001953125f) // subnormal: mant*2^-9 + : (1.0f + float(mant)*0.125f) * exp2(float(int(exp) - 7)); + wv = sign ? -mag : mag; + } + acc += wv * xr[i] * scl[i/128]; + } } else { // f32 device const float* wr = (device const float*)(w) + (long)o * I; device const float4* w4 = (device const float4*)wr; @@ -60,7 +89,8 @@ kernel void mm_gemv(device const uchar* w [[buffer(0)]], // raw weight by for (int i = I8*8 + slane; i < I; i += 32) acc += wr[i] * xr[i]; } acc = simd_sum(acc); - if (slane == 0) y[row] = acc * scale[o]; + // fmt==7 already folds its per-block scale into acc above -- do not scale again. + if (slane == 0) y[row] = (fmt == 7) ? acc : acc * scale[o]; } // Batched bindless expert GEMV: each row gr belongs to expert erow[gr], whose weight and @@ -423,8 +453,19 @@ static size_t fmt_bytes(int fmt, int I, int O) { if (fmt == 1) return (size_t)O * I; if (fmt == 2) return (size_t)O * ((I+1)/2); if (fmt == 3) return (size_t)O * ((I+3)/4); + if (fmt == 7) return (size_t)O * I; // fp8 e4m3: one raw byte/element, same as fmt=1 return (size_t)O * I * sizeof(float); } +// fp8 (fmt=7) scale-array size: one f32 per 128x128 BLOCK -> ceil(O/128)*ceil(I/128) (2D, +// not per-row -- quant.h isn't included here, so the ceil-div is inlined rather than sharing +// colibri.c's qt_scale_bytes/quant.h's fp8_nblk). f32 is this build's implemented scale +// ENCODING for fmt=7 (see quant.h/colibri.c) -- this file has no reason to know that a +// UE8M0 encoding exists at all: qt_resolve_fmt refuses it on the CPU read path before any +// tensor in that encoding could ever reach this Metal-side sizing helper. +static size_t fmt_scale_bytes(int fmt, int I, int O) { + if (fmt == 7) return (size_t)((O + 127) / 128) * (size_t)((I + 127) / 128) * sizeof(float); + return (size_t)O * sizeof(float); +} // Wrap host memory zero-copy if page-aligned, else copy into a shared buffer. static id wrap(const void *p, size_t n) { @@ -590,14 +631,16 @@ static size_t fmt_bytes(int fmt, int I, int O) { extern "C" int coli_metal_matmul(ColiMetalTensor **tp, float *y, const float *x, const void *weights, const float *scales, int fmt, int S, int I, int O) { - if (!g_dev || fmt < 0 || fmt > 3) return 0; + /* fmt==7 (fp8 passthrough) is an explicit allow-list entry, not folded into the 0..3 + * contiguous range check below: it is not adjacent to it. */ + if (!g_dev || fmt < 0 || (fmt > 3 && fmt != 7)) return 0; @autoreleasepool { ColiMetalTensor *t = *tp; if (!t) { t = new ColiMetalTensor(); t->fmt = fmt; t->I = I; t->O = O; t->wbytes = fmt_bytes(fmt, I, O); t->w = wrap(weights, t->wbytes); - t->s = wrap(scales, (size_t)O * sizeof(float)); + t->s = wrap(scales, fmt_scale_bytes(fmt, I, O)); *tp = t; g_tensor_count++; g_tensor_bytes += t->wbytes; } diff --git a/c/tests/test_backend_metal.mm b/c/tests/test_backend_metal.mm index bff0444b..99b77f42 100644 --- a/c/tests/test_backend_metal.mm +++ b/c/tests/test_backend_metal.mm @@ -7,7 +7,7 @@ #include #include -enum { F32=0, I8=1, I4=2, I2=3 }; +enum { F32=0, I8=1, I4=2, I2=3, FP8=7 }; static void cpu_ref(int fmt, const void *W, const float *s, const float *x, float *y, int S, int I, int O) { @@ -52,6 +52,142 @@ static int run(int fmt, int O, int I, int S, const char *name) { return ok?0:1; } +// ---- fmt=7 (native FP8-e4m3 passthrough -- see colibri.c): own CPU reference + +// harness -------------------------------------------------------------------------- +// Independent reference decode (bit manipulation, NOT quant.h's E4M3_LUT -- this file +// stays self-contained like the rest of its CPU references) for OCP E4M3-FN: exp==0 is +// subnormal, exp==0xF&&mant==0x7 is the only NaN code (both signs), else normal with +// bias 7. Must match quant.h's e4m3_decode AND the mm_gemv fmt==7 branch's in-kernel +// bit manipulation exactly -- see run_fp8_lut() below for the exhaustive 256-code check +// that proves the GPU kernel agrees with this reference (and by transitivity with the +// CPU LUT, which was cross-checked byte-for-byte against torch.float8_e4m3fn offline). +static float ref_e4m3(uint8_t b) { + unsigned sign=(b>>7)&1, exp=(b>>3)&0xF, mant=b&0x7; + if (exp==0xF && mant==0x7) return NAN; + float val = (exp==0) ? (float)mant*(1.0f/8.0f)*powf(2.0f,1.0f-7.0f) + : (1.0f+(float)mant*(1.0f/8.0f))*powf(2.0f,(float)exp-7.0f); + return sign ? -val : val; +} +static int fp8_nblk(int n){ return (n+127)/128; } + +static void cpu_ref_fp8(const uint8_t *q8, const float *bscale, const float *x, + double *y, double *mag, int S, int I, int O) { + int nblkI = fp8_nblk(I); + for (int o=0;o W((size_t)O*I); + std::vector scale((size_t)nblk), x((size_t)S*I), yg((size_t)S*O); + std::vector yr((size_t)S*O), mag((size_t)S*O); + srand(5150 + I*7 + O*3 + S*13); + for (auto &b : W) { do { b = (uint8_t)(rand()&0xFF); } while (b==0x7F || b==0xFF); } // exclude NaN codes + // distinct, easily-distinguishable scale per block (stride-audit idiom, same as + // test_fp8_passthrough.c's CPU version): a swapped o/128 vs i/128 stride, or a wrong + // nblkI in the shader's indexing, lands on a DIFFERENT block's scale and shows up as + // a loud numeric mismatch rather than passing by luck. + for (int b=0;b 1e-30 ? d/mag[i] : d; + if (rel > worst) worst = rel; + if (rel > 1e-4) bad++; + } + int ok = (bad == 0); + printf(" %-42s worst_rel=%.2e (I=%d O=%d nblkO=%d nblkI=%d S=%d) %s\n", + name, worst, I, O, nblkO, nblkI, S, ok?"ok":"*** MISMATCH"); + coli_metal_tensor_free(t); + return ok?0:1; +} + +// Exhaustive LUT-exactness check run THROUGH the actual GPU kernel: O=256,I=1 means +// each of the 256 output rows has exactly ONE weight byte, set to that row's own index +// (0..255) -- so row o's dequant is exactly e4m3_decode(o). x=[1.0] and both blocks' +// scale=1.0 isolate the decode from the block-scale/accumulation logic entirely. This +// mirrors test_fp8_passthrough.c's CPU LUT-exactness test (test_lut), but exercised +// through mm_gemv's in-kernel bit manipulation instead of quant.h's table -- proving +// the two independently-written decoders agree on all 256 codes, including both NaN +// codes and the sign of zero. +static int run_fp8_lut(const char *name) { + enum { O=256, I=1 }; + std::vector W(O*I); for (int b=0;b scale(fp8_nblk(O)*fp8_nblk(I), 1.0f); // nblkO=2,nblkI=1 -> both blocks scale=1 + std::vector x(I, 1.0f), yg(O); + ColiMetalTensor *t=nullptr; + if (!coli_metal_matmul(&t, yg.data(), x.data(), W.data(), scale.data(), FP8, 1, I, O)) { + printf(" %-42s FAIL (matmul returned 0)\n", name); return 1; } + int bad=0; + for (int b=0;b>1]; int v=(i&1)?(b>>4):(b&0xF); return (float)(v-8); } static size_t roundpg(size_t n){ size_t p=16384; return ((n+p-1)/p)*p; } @@ -251,6 +387,22 @@ int main(void) { fail |= run(I8, 2048,6144,4, "int8 gate/up S=4"); fail |= run(I4, 2048,6144,7, "int4 gate/up S=7 (odd)"); fail |= run(I4, 2050,6146,3, "int4 non-mult-4 dims"); + printf("Metal fmt=7 native FP8-e4m3 passthrough tests:\n"); + fail |= run_fp8_lut("fp8 LUT exactness (256/256 codes via GPU kernel)"); + fail |= run_fp8(2048,6144,1, "fp8 gate/up-shaped O=2048 I=6144 (spec example) S=1"); + fail |= run_fp8(6144,2048,1, "fp8 down-shaped O=6144 I=2048 S=1"); + fail |= run_fp8(2048,6144,4, "fp8 gate/up-shaped O=2048 I=6144 S=4"); + // block edges: O,I not multiples of 128 (the partial-tile clamp) + fail |= run_fp8(130, 200, 2, "fp8 block edges: O,I both non-mult-128"); + fail |= run_fp8(129, 128, 1, "fp8 block edges: O just over 128, I exact"); + fail |= run_fp8(128, 129, 1, "fp8 block edges: O exact, I just over 128"); + fail |= run_fp8(1, 1, 1, "fp8 degenerate 1x1 (single sub-block)"); + // non-square block grid (nblkO=3 != nblkI=48) with a distinct scale per block: a + // swapped o/128 vs i/128 index, or a wrong nblkI stride, lands on the wrong block's + // scale and this shape/scale choice makes that numerically loud. + fail |= run_fp8(384, 6144, 3, "fp8 non-square block grid nblkO=3 nblkI=48 (stride audit)"); + fail |= run_fp8_gemm_gate("fp8 GEMM entry explicitly gated off (coli_metal_gemm refuses)"); + fail |= run_fp8_moe_gate("fp8 MB_BUILD/moe_submit entry gated off (shared-expert fmt=7 hazard)"); printf("Metal batched moe_block tests:\n"); fail |= run_moe({1,1,1,1,1,1,1,1}, "moe decode nb=8"); fail |= run_moe({3,1,4,2,1,5}, "moe ragged nb=6"); From 3d9aba8dc2e06aeb216c1b5c649c59fa57991e29 Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 15:49:15 -0400 Subject: [PATCH 03/12] fp8: repack tool (Z.ai FP8 shards -> byte-preserved fmt=7 container) Ports fmt6/fp8-passthrough-r2@b927911 (repack tool folded with its kv_b_proj self-review fix, as the source branch shipped it), renumbered fmt=100 -> fmt=7. This is the tool's PRE-metadata-stamp shape: it writes no __metadata__["colibri.fmt"] stamp, matching this PR's scope (stamp writer + reader + registry are a separate, follow-up PR per the maintainer's #524 reply). The module docstring says so explicitly now, and documents the scale-encoding property this tool implements (f32 block scales only -- Z.ai's GLM-5.2-FP8 checkpoints ship f32 weight_scale_inv, so there is no UE8M0 source data this tool would ever need to convert). tools/repack_fp8_passthrough.py: unlike convert_fp8_to_int4.py (dequant -> requant to a lossy format), copies fp8 weight bytes AS-IS and only renames/reshapes the _scale_inv sidecar to the engine's .qs convention (flat F32, ceil(O/128)*ceil(I/128) elements). Selection reuses convert_fp8_to_int4's classify()/check_or_record_params to target resident kinds only: shared expert, o_proj, other attention projections, dense-MLP first layers, and the generic resident fallback. Routed experts stay on the existing int4-g64 path; kv_b_proj is excluded because the CPU MLA-absorb path (qt_addrow/qt_matvec_rows) and the CUDA absorb kernels have no fmt=7 case and would silently misread it as int2-packed data. _check_geometry refuses (ValueError) rather than silently repack a shard whose _scale_inv shape doesn't match ceil(O/128)xceil(I/128) -- the write-side twin of qt_resolve_fmt's read-side "THE DESIGN LANDMINE" refusal. tests/test_fp8_repack.py: synthetic-fixture-only suite (glm_fp8_emit's exact real-checkpoint FP8 layout, no real Z.ai shard read or written) covering selection (resident kinds byte-preserved; routed experts/io/ f32/kv_b_proj excluded), byte-for-byte weight preservation + .qs rename, the write-side geometry-refusal path, --dry-run writing nothing, and the #383-class resume/params-guard idiom. Renumbered from the source's fmt=100. tools/README.md: one-line entry for the tool, renumbered. RAN: python3 -m unittest tests.test_fp8_repack -v -- 9/9 tests pass (this commit, torch + safetensors present in this environment). Co-Authored-By: Claude Fable 5 --- c/tests/test_fp8_repack.py | 210 +++++++++++++++++++++++ c/tools/README.md | 2 + c/tools/repack_fp8_passthrough.py | 272 ++++++++++++++++++++++++++++++ 3 files changed, 484 insertions(+) create mode 100644 c/tests/test_fp8_repack.py create mode 100644 c/tools/repack_fp8_passthrough.py diff --git a/c/tests/test_fp8_repack.py b/c/tests/test_fp8_repack.py new file mode 100644 index 00000000..283de225 --- /dev/null +++ b/c/tests/test_fp8_repack.py @@ -0,0 +1,210 @@ +"""tools/repack_fp8_passthrough.py: fmt=7 repack tool tests. + +fmt=7 is a PUBLIC ordinal, assigned by the maintainer on #524: see +repack_fp8_passthrough.py's module docstring for the fmt=6 -> fmt=100 -> +fmt=7 history (dev's own #465 merged a REAL fmt=6, E8/IQ3, so this tool's +format moved to the PRIVATE ORDINAL BLOCK as fmt=100 during development +before graduating to fmt=7 at merge). + +Synthetic fixtures ONLY (tools/glm_fp8_emit.py, the exact real-checkpoint FP8 +layout that convert_fp8_to_int4.py's dequant() reads) -- no real Z.ai shard is +read or written by this suite, per the build's hard constraint. Covers: +selection (resident kinds byte-preserved, routed experts / io / f32 excluded), +byte-for-byte preservation of the fp8 weight and the .qs scale rename, the +scale-geometry refusal path (THE DESIGN LANDMINE's write-side twin: a +malformed source shard must be refused, not silently repacked into a +container the engine's qt_resolve_fmt would then misread), --dry-run writing +nothing, and the #383-class resume/params-guard behavior. +""" +import glob, json, os, struct, sys, tempfile, unittest + +try: + import torch + from safetensors import safe_open + from safetensors.torch import save_file +except ImportError as e: + raise unittest.SkipTest(f"torch/safetensors not installed: {e}") + +sys.path.insert(0, os.path.join(os.path.dirname(__file__), "..", "tools")) +from glm_fp8_emit import save_fp8_safetensors +import repack_fp8_passthrough as rp + + +def _read_header(path): + with open(path, "rb") as f: + hlen = struct.unpack(" is_repack_target will select it + save_file({"model.layers.0.self_attn.o_proj.weight": w, + "model.layers.0.self_attn.o_proj.weight_scale_inv": bad_scale}, bad_path) + with self.assertRaises(ValueError): + rp.repack_shard(bad_path, n_layers=5) + with self.assertRaises(ValueError): + rp.shard_inventory(bad_path, n_layers=5) + + +class ResumeAndParamsGuardTest(unittest.TestCase): + def setUp(self): + self.tmp = tempfile.TemporaryDirectory() + self.indir = os.path.join(self.tmp.name, "fp8src") + os.makedirs(self.indir) + self.shard = os.path.join(self.indir, "model-00001-of-00001.safetensors") + _make_fixture(self.shard) + self.outdir = os.path.join(self.tmp.name, "out") + self.tool = os.path.join(os.path.dirname(__file__), "..", "tools", "repack_fp8_passthrough.py") + + def tearDown(self): + self.tmp.cleanup() + + def _run(self, n_layers): + return os.system(f'{sys.executable} "{self.tool}" --indir "{self.indir}" ' + f'--outdir "{self.outdir}" --n-layers {n_layers} >/dev/null 2>&1') + + def test_output_container_geometry(self): + self.assertEqual(self._run(5), 0) + outs = glob.glob(os.path.join(self.outdir, "out-fp8pass-*.safetensors")) + self.assertEqual(len(outs), 1) + hdr = _read_header(outs[0]) + # o_proj is [256,256] -> nblkO=nblkI=2 -> 4 block scales + qs = hdr["model.layers.0.self_attn.o_proj.weight.qs"] + self.assertEqual(qs["dtype"], "F32") + self.assertEqual(qs["shape"], [4]) + w = hdr["model.layers.0.self_attn.o_proj.weight"] + self.assertEqual(w["dtype"], "U8") + self.assertEqual(w["shape"], [256, 256]) + # experts / io / f32 / kv_b_proj must be absent from the output container entirely + self.assertNotIn("model.layers.0.mlp.experts.0.gate_proj.weight", hdr) + self.assertNotIn("model.embed_tokens.weight", hdr) + self.assertNotIn("model.layers.0.input_layernorm.weight", hdr) + self.assertNotIn("model.layers.0.self_attn.kv_b_proj.weight", hdr) + + def test_resume_skips_completed_shard(self): + self.assertEqual(self._run(5), 0) + mtime1 = os.path.getmtime(glob.glob(os.path.join(self.outdir, "out-fp8pass-*.safetensors"))[0]) + self.assertEqual(self._run(5), 0) # rerun, same params + outs = glob.glob(os.path.join(self.outdir, "out-fp8pass-*.safetensors")) + self.assertEqual(len(outs), 1, "resume must not duplicate output shards") + self.assertEqual(os.path.getmtime(outs[0]), mtime1, "resume must not rewrite a completed shard") + + def test_params_mismatch_refused(self): + self.assertEqual(self._run(5), 0) + rc = self._run(78) # different n_layers, same outdir + self.assertNotEqual(rc, 0, "a params mismatch on the same outdir must be refused") + # and the FIRST run's output must be untouched by the refused second run + outs = glob.glob(os.path.join(self.outdir, "out-fp8pass-*.safetensors")) + self.assertEqual(len(outs), 1) + + +if __name__ == "__main__": + unittest.main() diff --git a/c/tools/README.md b/c/tools/README.md index 20937885..b058556f 100644 --- a/c/tools/README.md +++ b/c/tools/README.md @@ -4,6 +4,8 @@ These scripts support model preparation and offline engineering work. They are not runtime dependencies of the C engine. - `convert_fp8_to_int4.py`, `download_glm52.py`: model preparation +- `repack_fp8_passthrough.py`: fmt=7 repack (byte-preserved FP8, resident kinds only; + see the module docstring -- synthetic-fixture-tested only, no real-shard runs yet) - `make_glm_oracle.py`, `make_glm_bench_model.py`: deterministic fixtures - `benchmark_cuda_fixture.py`, `eval_glm.py`, `fetch_benchmarks.py`: benchmarks - `gen_unicode.py`: tokenizer table generation diff --git a/c/tools/repack_fp8_passthrough.py b/c/tools/repack_fp8_passthrough.py new file mode 100644 index 00000000..6a1a3e0d --- /dev/null +++ b/c/tools/repack_fp8_passthrough.py @@ -0,0 +1,272 @@ +"""fmt=7 repack tool: Z.ai GLM-5.2-FP8 shards -> the engine's native-FP8-passthrough +container (byte-preserved, no dequant/requant). + +fmt=7 is a PUBLIC ordinal, assigned by the maintainer on #524. This tool's +output format was minted fmt=6 during this branch's original development, +before dev's own #465 (E8/IQ3) claimed that ordinal upstream and merged it +into dev as a REAL fmt=6 (colibri.c's qt_resolve_fmt, quant.h's E8 +constants); re-tagged fmt=100 (PRIVATE ORDINAL BLOCK, see colibri.c's QT +struct comment) from that point forward -- this tool has never emitted a +container claiming ordinal 6 -- and graduated to fmt=7 at merge. + +Unlike convert_fp8_to_int4.py (which DEQUANTIZES fp8 -> f32 -> REQUANTIZES to a +different, lossy format), this tool copies the fp8 weight bytes AS-IS and only +renames/reshapes the scale sidecar to the engine's `.qs` convention. Same 8 +bits/weight streaming cost as the source, zero additional quantization loss. + +THE DESIGN LANDMINE (see qt_resolve_fmt in colibri.c, c/quant.h's E4M3_LUT +comment): the engine tells fmt=7 (this tool's output) apart from fmt=1 (plain +int8) ONLY by scale-array geometry -- per-row for fmt=1, per-128x128-block for +fmt=7. This tool must therefore emit .qs sidecars whose byte count is EXACTLY +ceil(O/128)*ceil(I/128)*4, never anything that could coincide with O*4 by +accident for the shapes it targets (see _check_geometry below, which refuses +rather than silently emitting a container the engine might misread). At I=98 +specifically, a single-block fp8 tensor's raw weight-byte count ALSO coincides +with a genuine fmt=6 (E8/IQ3) tensor's -- see qt_resolve_fmt's "SECOND DESIGN +LANDMINE" comment; this tool does not target I=98 shapes in practice (Z.ai's +real projections don't land there), but the engine's read-side refusal covers +it regardless. + +SCALE ENCODING: this tool always emits f32 block scales (4 bytes/block, the +`.to(torch.float32)` below) -- the ENCODING this build implements. It never +reads or emits a UE8M0 (1 byte/block) sidecar: Z.ai's GLM-5.2-FP8 checkpoints +ship f32 `weight_scale_inv`, so there is no source data this tool would need +to convert from. A DeepSeek-V4-shaped source (same weight geometry, UE8M0 +scales) is out of scope for this tool as written; the read side (qt_resolve_fmt) +recognizes and refuses that byte signature by name rather than misread it, +should a container carrying it ever reach the engine some other way. + +SELECTED kinds only: "resident" tensors (shared expert, o_proj, other +attention projections, dense-MLP first layers, and the generic resident +fallback) -- routed experts (kind "x" in convert_fp8_to_int4.classify) are +EXCLUDED and stay on the existing int4-g64 streaming path; routers/norms/bias +(kind "f32") and embed/lm_head (kind "io", BF16 in source, never FP8) are +untouched by this tool -- carrying those into a mixed fmt=7+int4 loadout +directory is separate, deferred "loadout index-rewrite blending" work. + +kv_b_proj (kind "kvb") is ALSO deliberately excluded, despite being a resident +tensor classify() would otherwise select: colibri.c's MLA-absorption CPU path +(qt_addrow/qt_matvec_rows, called only on l->kv_b -- the always-available +fallback whenever the Metal fused decode kernel isn't used) has no fmt==7 +case and would silently misread fp8 bytes as int2-packed data; the CUDA +absorb kernels (coli_cuda_attention_absorb/_kvdev in backend_cuda.cu) are +similarly int4-specific. Repacking kv_b_proj to fmt=7 needs that CPU+CUDA +absorb-path work first -- out of scope for this build, so this tool refuses +to produce a container the engine cannot safely read. + +This tool does NOT write a container metadata stamp: __metadata__-based +self-description (writer + qt_verify_fmt_stamp reader + docs/FORMATS.md +registry) is a separate, follow-up PR, per the maintainer's #524 scoping. +This PR's collision/refusal logic (qt_resolve_fmt) therefore refuses +UNCONDITIONALLY at every ambiguous shape this tool could produce -- there is +no stamp to break a tie. + +HARD CONSTRAINT for this build: unit-tested on synthetic fixtures ONLY (see +tests/test_fp8_repack.py, built with tools/glm_fp8_emit.py's exact real- +checkpoint layout) plus --dry-run inventory mode against those same fixtures. +NO read of any real Z.ai shard, NO full or partial repack of a real checkpoint +-- that is later, user-GO'd work once this build's plumbing has been reviewed. + +Usage (synthetic fixtures / local testing only): + python3 tools/repack_fp8_passthrough.py --indir --outdir --dry-run + python3 tools/repack_fp8_passthrough.py --indir --outdir --n-layers 78 +""" +import argparse, glob, json, os, sys + +sys.path.insert(0, os.path.dirname(__file__)) +from convert_fp8_to_int4 import classify, check_or_record_params # reuse: same taxonomy, + # same #383-class params guard + +# Kinds classify() can return that this tool targets: resident weights, NOT routed +# experts ("x"), NOT the always-F32 set ("f32"), NOT embed/lm_head ("io" -- BF16 in +# source, never FP8), NOT sidecars/skips ("consumed"/"skip"). "kvb" (kv_b_proj) is +# ALSO excluded -- see the module docstring: the CPU absorb path (qt_addrow/ +# qt_matvec_rows) and the CUDA absorb kernels have no fmt=7 case yet. +RESIDENT_KINDS = frozenset({"sh", "o", "attn", "dmlp", "q"}) + +BLOCK = 128 + + +def _nblk(n): + return (n + BLOCK - 1) // BLOCK + + +def is_repack_target(name, dtype, keys, n_layers): + """True if `name` is a resident-kind FP8 tensor this tool should byte-preserve + into the fmt=7 container. `keys` is the full set of tensor names in the shard + (needed to confirm the `_scale_inv` sidecar is actually present -- classify() + alone can't tell FP8 tensors from any other dtype a resident-kind name might + carry in a non-FP8 checkpoint variant).""" + if name.endswith("_scale_inv"): + return False # sidecar, handled together with its weight + kind = classify(name, n_layers) + if kind not in RESIDENT_KINDS: + return False + if dtype not in ("F8_E4M3", "float8_e4m3fn"): + return False + return (name + "_scale_inv") in keys + + +def _check_geometry(name, O, I, nblkO, nblkI): + """Refuse (loud, ValueError) rather than silently emit a .qs whose geometry + doesn't match what qt_resolve_fmt expects for [O,I] -- the same "untrusted + container" discipline the C side applies on read, applied here on write so a + malformed source checkpoint is caught at repack time, not at load time three + steps later with a confusing engine-side refusal.""" + exp_nblkO, exp_nblkI = _nblk(O), _nblk(I) + if (nblkO, nblkI) != (exp_nblkO, exp_nblkI): + raise ValueError( + f"{name}: scale_inv block shape ({nblkO},{nblkI}) != expected " + f"({exp_nblkO},{exp_nblkI}) for [{O},{I}] -- refusing to repack a shape " + f"the engine's qt_resolve_fmt would either refuse or (worse) misread") + + +def shard_inventory(path, n_layers): + """--dry-run: scan one shard's header only (get_slice, no tensor data pulled) + and return the list of tensors that WOULD be repacked, without writing + anything. Cheap: safe to run over an entire real checkpoint's headers.""" + from safetensors import safe_open + inv = [] + with safe_open(path, framework="pt") as f: + keys = set(f.keys()) + for name in f.keys(): + sl = f.get_slice(name) + dt = sl.get_dtype() + if not is_repack_target(name, dt, keys, n_layers): + continue + O, I = sl.get_shape() + sc_slice = f.get_slice(name + "_scale_inv") + nblkO, nblkI = sc_slice.get_shape() + _check_geometry(name, O, I, nblkO, nblkI) + inv.append({"name": name, "kind": classify(name, n_layers), + "O": O, "I": I, "nblkO": nblkO, "nblkI": nblkI, + "weight_bytes": O * I, "scale_bytes": nblkO * nblkI * 4}) + return inv + + +def repack_shard(path, n_layers): + """Real mode: byte-preserve every selected tensor's weight bytes (raw e4m3, + no dequant/requant -- `.view(torch.uint8)` is a pure bit-reinterpret, verified + byte-identical against torch.float8_e4m3fn's own storage) and rename/flatten + the `_scale_inv` sidecar to the engine's `name.qs` convention (flat F32, + nblkO*nblkI elements, row-major [blkO,blkI] -- matching + scale[(o/128)*nblkI+i/128] on the read side). Returns (out_dict, inventory).""" + import torch + from safetensors import safe_open + out = {} + inv = [] + with safe_open(path, framework="pt") as f: + keys = set(f.keys()) + for name in f.keys(): + sl = f.get_slice(name) + dt = sl.get_dtype() + if not is_repack_target(name, dt, keys, n_layers): + continue + w = f.get_tensor(name) # torch.float8_e4m3fn [O,I] + sc = f.get_tensor(name + "_scale_inv") # torch.float32 [nblkO,nblkI] + O, I = w.shape + nblkO, nblkI = sc.shape + _check_geometry(name, O, I, nblkO, nblkI) + out[name] = w.view(torch.uint8).contiguous() # BYTE-PRESERVED + out[name + ".qs"] = sc.to(torch.float32).reshape(-1).contiguous() + inv.append({"name": name, "kind": classify(name, n_layers), + "O": O, "I": I, "nblkO": nblkO, "nblkI": nblkI, + "weight_bytes": O * I, "scale_bytes": nblkO * nblkI * 4}) + return out, inv + + +def _print_inventory_summary(all_inv, dry_run): + by_kind = {} + tot_w = tot_s = 0 + for it in all_inv: + by_kind.setdefault(it["kind"], []).append(it) + tot_w += it["weight_bytes"]; tot_s += it["scale_bytes"] + tag = "[DRY-RUN]" if dry_run else "[REPACK]" + print(f"{tag} {len(all_inv)} tensor(s) selected across all shards " + f"({tot_w/1e9:.3f} GB weights, {tot_s/1e6:.3f} MB scales)") + for kind in sorted(by_kind): + items = by_kind[kind] + w = sum(it["weight_bytes"] for it in items) + print(f"{tag} kind={kind:5s} {len(items):5d} tensor(s) {w/1e9:.3f} GB") + + +def main(): + ap = argparse.ArgumentParser(description=__doc__, + formatter_class=argparse.RawDescriptionHelpFormatter) + ap.add_argument("--indir", required=True, help="directory of Z.ai FP8 *.safetensors shards") + ap.add_argument("--outdir", required=True, help="destination for repacked fmt=7 container shards") + ap.add_argument("--n-layers", type=int, default=78) + ap.add_argument("--dry-run", action="store_true", + help="inventory only: scan headers, print selection counts, write nothing") + a = ap.parse_args() + + shards = sorted(glob.glob(os.path.join(a.indir, "*.safetensors"))) + if not shards: + print(f"ERROR: no *.safetensors files in {a.indir}"); sys.exit(1) + + if a.dry_run: + all_inv = [] + for sp in shards: + all_inv.extend(shard_inventory(sp, a.n_layers)) + _print_inventory_summary(all_inv, dry_run=True) + return + + os.makedirs(a.outdir, exist_ok=True) + # #383-class guard, reused (not reimplemented): refuses a resumed run whose + # params differ from what's already partially in this outdir (the #355 failure + # mode -- a second pass with different flags silently mixing containers). + # Owns its own sidecar (.fp8pass-params.json), separate from the shard-progress + # manifest below, exactly like the --repo mtp/indexer loops in + # convert_fp8_to_int4.py use it. + params = {"mode": "fp8-passthrough", "n_layers": a.n_layers} + if not check_or_record_params(a.outdir, "fp8pass-", params): + sys.exit(1) + + # RESUME (same #383-class manifest idiom as convert_fp8_to_int4.py's --indir path): + # a sidecar records input-shard -> output-shard-name-or-"" (shards with no + # resident-FP8 tensors emit no file), written atomically so an interrupted run + # never leaves a half-written manifest for the next invocation to trust. The + # params guard above already owns "don't mix conversions" -- this manifest is + # purely which shards are done. + prog_path = os.path.join(a.outdir, ".fp8pass-progress.json") + prog = {} + if os.path.exists(prog_path): + try: + prog = json.loads(open(prog_path).read()) + except (OSError, ValueError): + prog = {} + done = prog.setdefault("shards", {}) + + from safetensors.torch import save_file + all_inv = [] + n = fresh = skipped = 0 + for sp in shards: + key = os.path.basename(sp) + prev = done.get(key) + if prev is not None and (prev == "" or os.path.exists(os.path.join(a.outdir, prev))): + if prev: + n += 1 + skipped += 1 + continue + out, inv = repack_shard(sp, a.n_layers) + all_inv.extend(inv) + if not out: + done[key] = "" + else: + name = f"out-fp8pass-{n:05d}.safetensors" + save_file(out, os.path.join(a.outdir, name)) + done[key] = name + n += 1 + fresh += 1 + tmp = prog_path + ".tmp" + with open(tmp, "w") as fh: + json.dump(prog, fh, indent=1) + os.replace(tmp, prog_path) + if skipped: + print(f"[RESUME] {skipped} shard(s) already done in {a.outdir}, skipped") + print(f"[REPACK] {fresh} new output shard(s) written") + _print_inventory_summary(all_inv, dry_run=False) + + +if __name__ == "__main__": + main() From 1c82bdecec1a0a01544e00f2f376302f6a9af85f Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 15:52:32 -0400 Subject: [PATCH 04/12] fp8: end-to-end test, real repack tool -> real C loader Ports fmt6/fp8-passthrough-r2@b2aea6c, renumbered fmt=100 -> fmt=7 and adapted to drop the metadata-stamp check the source commit's version carried: this branch's repack tool (previous commit) writes no __metadata__ stamp, so there is nothing to sanity-check on that axis. Closes the last coverage gap: neither test_fp8_load.c (hand-authored wire-format fixtures -- proves the LOADER's own logic in isolation) nor test_fp8_repack.py (proves the WRITER's Python-side output shape, never invokes the C loader) exercises the real round trip. tests/test_fp8_e2e_loader.c: minimal harness (`#define main coli_glm_main_unused; #include "../colibri.c"`, same pattern as every other colibri.c-embedding test in this suite). Takes a container directory plus (name,O,I) triples on argv, calls the REAL qt_from_disk for each, and asserts fmt==7 resolved, q8/s both non-NULL, and every dequantized value finite. tests/test_fp8_e2e_repack_load.py: builds a synthetic FP8 checkpoint via tools/glm_fp8_emit.py, repacks it with the REAL tools/repack_fp8_passthrough.py via subprocess (the actual CLI entry point, not an imported-function reimplementation), sanity-checks the real tool's own selection by tensor presence before ever invoking the C side (so a failure localizes to tool vs. loader) -- and, since this PR's tool writes no container metadata stamp, additionally asserts __metadata__ is ABSENT from the repacked output (a regression guard that this PR's tool really stays out of the stamp PR's scope, not just an omitted assertion) -- compiles test_fp8_e2e_loader.c with production-mirroring flags (asserted zero warnings), and runs it against the real repacked output directory. Includes a routed-expert tensor and an f32 norm as negative controls, proving the real tool's selection logic held on an actual round trip, not just in test_fp8_repack.py's own synthetic-fixture suite. This closes out Branch A / PR 1's content: CPU read path + matmul_fp8, Metal kernel, repack tool (no stamp), collision/refusal logic (fmt=1-vs-fmt=7 THE DESIGN LANDMINE + fmt=6-vs-fmt=7 SECOND DESIGN LANDMINE, both refuse-unconditionally, no stamped_name parameter anywhere), the UE8M0 scale-encoding recognition seam, and now this round-trip test. The metadata stamp itself (writer, qt_verify_fmt_stamp, stamped_name plumbing, docs/FORMATS.md) is Branch B / PR 2, stacked on top of this branch. RAN (this commit): python3 -m unittest tests.test_fp8_e2e_repack_load -v -> 1/1 pass ("ok"). make clean && make glm METAL=1 -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0 (includes test_fp8_passthrough, test_fp8_load, test_int3/test_int3_load, and every other pre-existing case). make test-python -> 152 tests, 10 skipped, 0 failures, exit 0 (includes this branch's test_fp8_repack and test_fp8_e2e_repack_load, plus dev's own suite). make metal-test under COLI_METAL_RESSET=0 AND =1: both full suite green, exit 0, all fp8 GPU lines ok, RESSET=1 confirmed active via "[METAL] residency-set: on". Co-Authored-By: Claude Fable 5 --- c/tests/test_fp8_e2e_loader.c | 86 +++++++++++++++ c/tests/test_fp8_e2e_repack_load.py | 160 ++++++++++++++++++++++++++++ 2 files changed, 246 insertions(+) create mode 100644 c/tests/test_fp8_e2e_loader.c create mode 100644 c/tests/test_fp8_e2e_repack_load.py diff --git a/c/tests/test_fp8_e2e_loader.c b/c/tests/test_fp8_e2e_loader.c new file mode 100644 index 00000000..c2ee9884 --- /dev/null +++ b/c/tests/test_fp8_e2e_loader.c @@ -0,0 +1,86 @@ +/* fp8 passthrough END-TO-END loader harness. + * + * Neither test_fp8_load.c (hand-built wire-format C fixtures -- proves the + * LOADER's own disambiguation logic in isolation, but the container bytes + * are hand-authored, not tool-produced) nor test_fp8_repack.py (proves the + * WRITER's Python-side output shape -- never invokes the C loader at all) + * exercises the REAL round trip: a container the REAL + * tools/repack_fp8_passthrough.py actually produced, fed into the REAL + * st_init/qt_from_disk. This binary is the C-side half of that round trip; + * tests/test_fp8_e2e_repack_load.py drives it end to end (builds a synthetic + * FP8 checkpoint via tools/glm_fp8_emit.py, repacks it with the real tool + * via subprocess -- the actual CLI, not a reimplementation -- then invokes + * this binary against the real output directory). + * + * Usage: test_fp8_e2e_loader [name O I]... + * For every (name,O,I) triple, calls the REAL qt_from_disk (the identical + * function every model load uses) and asserts: + * (a) fmt==7 resolved (native fp8-e4m3-passthrough, via byte-arithmetic + * inference alone -- this PR's repack tool writes no container + * metadata stamp, see repack_fp8_passthrough.py's module docstring); + * (b) the weight (q8) and scale (s) buffers are non-NULL; + * (c) every dequantized value (e4m3_decode(byte) * block scale) is finite + * -- catches a decode-table or block-index bug a pure byte-count + * check wouldn't. */ +#define main coli_glm_main_unused +#include "../colibri.c" +#undef main + +#include +#include +#include +#include + +int main(int argc, char **argv){ + if(argc < 2){ fprintf(stderr,"usage: %s [name O I]...\n", argv[0]); return 2; } + if((argc-2) % 3 != 0){ + fprintf(stderr,"args after must come in (name,O,I) triples (got %d)\n", argc-2); + return 2; + } + const char *dir = argv[1]; + int ntensors = (argc-2)/3; + if(ntensors == 0){ fprintf(stderr,"no (name,O,I) triples given -- nothing to check\n"); return 2; } + + static Model gm; memset(&gm,0,sizeof gm); + st_init(&gm.S, dir); + + int fails = 0; + for(int i=0;i= 0){ + printf("FAIL %s: non-finite dequantized value at flat index %lld\n", name, (long long)bad); + fails++; continue; + } + printf("ok %s: fmt=7 O=%d I=%d, weights+scale loaded through the real loader, all-finite\n", + name, O, I); + } + if(fails){ printf("fp8 e2e repack->load: %d/%d tensor(s) FAILED\n", fails, ntensors); return 1; } + printf("fp8 e2e repack->load: ok (%d tensor(s))\n", ntensors); + return 0; +} diff --git a/c/tests/test_fp8_e2e_repack_load.py b/c/tests/test_fp8_e2e_repack_load.py new file mode 100644 index 00000000..4cd5e8e5 --- /dev/null +++ b/c/tests/test_fp8_e2e_repack_load.py @@ -0,0 +1,160 @@ +"""End-to-end: REAL tools/repack_fp8_passthrough.py output fed into the REAL +C loader (st_init/qt_from_disk in colibri.c) via a small C harness. + +Neither test_fp8_load.c (hand-built wire-format C fixtures -- proves the C +loader's OWN logic, but the container bytes are hand-authored, not +tool-produced) nor test_fp8_repack.py (proves the Python tool's OUTPUT +shape, but never invokes the C loader at all) covers this round trip. This +test does, end to end: + + 1. build a synthetic FP8 checkpoint with tools/glm_fp8_emit.py (the same + fixture-generation helper test_fp8_repack.py uses -- not reimplemented + here, imported directly); + 2. repack it with the REAL tools/repack_fp8_passthrough.py, invoked as a + subprocess exactly as a user would run it from the command line (not + imported and called as a library, so the CLI entry point is exercised + too, not just the importable functions); + 3. compile tests/test_fp8_e2e_loader.c -- production flags mirroring the + Makefile's own CFLAGS, asserted to produce zero warnings -- into a + temp binary; + 4. run that binary against the REAL repacked output directory, asserting + every selected tensor loads fmt=7 through the real qt_from_disk with + finite dequantized values (byte-arithmetic inference alone -- this PR's + repack tool writes no container metadata stamp). + +Hermetic: everything (checkpoint, repacked output, compiled harness binary) +lives under one TemporaryDirectory; nothing is left behind. +""" +import glob, json, os, shutil, struct, subprocess, sys, tempfile, unittest + +try: + import torch +except ImportError as e: + raise unittest.SkipTest(f"torch not installed: {e}") + +HERE = os.path.dirname(os.path.abspath(__file__)) +C_DIR = os.path.normpath(os.path.join(HERE, "..")) +sys.path.insert(0, os.path.join(C_DIR, "tools")) +from glm_fp8_emit import save_fp8_safetensors # reuse: real fp8 block-quantize, not reimplemented + + +def _cc_flags(): + """Mirror the Makefile's CFLAGS closely enough to compile colibri.c + cleanly: -O3 + the same warning flags, plus libomp on macOS if present + (falls back to single-threaded -- exactly like the Makefile's own + OMPDIR probe -- if it's not, rather than failing the build). Returns + (cc, cflags, ldflags) or (None, None, None) if no compiler is found.""" + cc = shutil.which("cc") or shutil.which("clang") or shutil.which("gcc") + if not cc: + return None, None, None + cflags = ["-O3", "-Wall", "-Wextra", "-Wno-unused-parameter", + "-Wno-misleading-indentation", "-Wno-unused-function"] + ldflags = ["-lm"] + if sys.platform == "darwin": + try: + prefix = subprocess.run(["brew", "--prefix", "libomp"], capture_output=True, + text=True, timeout=10).stdout.strip() + except (OSError, subprocess.TimeoutExpired, FileNotFoundError): + prefix = "" + inc, lib = os.path.join(prefix, "include"), os.path.join(prefix, "lib") + if prefix and os.path.exists(os.path.join(inc, "omp.h")): + cflags += ["-Xclang", "-fopenmp", "-I", inc] + ldflags += ["-L", lib, "-lomp"] + return cc, cflags, ldflags + + +class Fp8RepackLoadE2ETest(unittest.TestCase): + """The real tools/repack_fp8_passthrough.py -> the real C loader, once.""" + + def setUp(self): + self.tmp = tempfile.TemporaryDirectory() + self.indir = os.path.join(self.tmp.name, "fp8src") + os.makedirs(self.indir) + self.outdir = os.path.join(self.tmp.name, "out") + self.shard = os.path.join(self.indir, "model-00001-of-00001.safetensors") + + def tearDown(self): + self.tmp.cleanup() + + def test_real_repack_output_loads_through_real_c_loader(self): + D, I_ = 256, 384 # non-degenerate shapes: nblkO*nblkI never coincides with O (no + # THE DESIGN LANDMINE ambiguity here -- that boundary is + # test_fp8_load.c's job, this test's job is the round trip) + torch.manual_seed(7) + sd = { + "model.layers.0.self_attn.o_proj.weight": torch.randn(D, D) * 0.02, + "model.layers.0.self_attn.q_a_proj.weight": torch.randn(D, D) * 0.02, + "model.layers.0.mlp.shared_experts.gate_proj.weight": torch.randn(D, I_) * 0.02, + "model.layers.0.mlp.gate_proj.weight": torch.randn(D, I_) * 0.02, + # routed expert: must NOT be selected (stays int4-g64 path) -- included + # as a negative control so this test also proves the real tool's selection + # logic held on the real round trip, not just in test_fp8_repack.py's own suite. + "model.layers.0.mlp.experts.0.gate_proj.weight": torch.randn(I_, D) * 0.02, + "model.layers.0.input_layernorm.weight": torch.randn(D), # f32: must NOT be selected + } + # (O, I) exactly as the engine/repack tool see them -- known here because this + # test authored the checkpoint; the real C loader never gets told this by us, + # only by the container it reads (st_init parses shape from the file itself; + # O/I are qt_from_disk's own [O,I] contract, same as every real model load). + shapes = { + "model.layers.0.self_attn.o_proj.weight": (D, D), + "model.layers.0.self_attn.q_a_proj.weight": (D, D), + "model.layers.0.mlp.shared_experts.gate_proj.weight": (D, I_), + "model.layers.0.mlp.gate_proj.weight": (D, I_), + } + save_fp8_safetensors(sd, self.shard) + + tool = os.path.join(C_DIR, "tools", "repack_fp8_passthrough.py") + rc = subprocess.run([sys.executable, tool, "--indir", self.indir, + "--outdir", self.outdir, "--n-layers", "5"], + capture_output=True, text=True) + self.assertEqual(rc.returncode, 0, f"real repack tool failed:\nSTDOUT:\n{rc.stdout}\nSTDERR:\n{rc.stderr}") + + outs = glob.glob(os.path.join(self.outdir, "out-fp8pass-*.safetensors")) + self.assertEqual(len(outs), 1, "expected exactly one repacked output shard") + + # Sanity-check the real tool's own selection BEFORE asking the C harness to load + # it: if this fails, the bug is in the tool, not the loader, and running the + # harness anyway would only muddy which side broke. This PR's repack tool writes + # NO container metadata stamp (that's a separate, follow-up PR -- see the tool's + # own module docstring), so selection is checked by tensor presence instead of a + # stamp's tensor-name set, and the stamp's ABSENCE is itself asserted below: a + # regression guard that this PR's tool really stays out of that scope. + with open(outs[0], "rb") as f: + hlen = struct.unpack("load: ok", run.stdout) + + +if __name__ == "__main__": + unittest.main() From 6438195905eae0fa8af44e3bc41e6b9d071ebda9 Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 19:39:13 -0400 Subject: [PATCH 05/12] fp8: invert THE DESIGN LANDMINE ambiguous-shape collision to fmt=1 Maintainer review, #528 (the load-bearing blocker): self_attn.o_proj.weight loads as [D,H*v_head] = [6144,16384] on GLM-5.2 -- nblkO=ceil(6144/128)=48, nblkI=ceil(16384/128)=128, product=6144==O, so a genuine, valid, pre-existing int8-row o_proj tensor satisfies BOTH the per-row (fmt=1) and per-128x128-block (fmt=7) scale-byte-count conventions at once. qt_resolve_fmt's prior behavior was to refuse this shape unconditionally (exit(1)) -- a real regression on an ordinary, already-shipping model, strictly worse than the misread it guarded against. Confirmed at shard-header level on this repo's own glm52_v1_e8x4g64 container: o_proj weight U8 100,663,296 B (==6144*16384), .qs scale blob 24,576 B (==6144*4==48*128*4, satisfying both formulas at once). colibri.c: qt_resolve_fmt's is_row&&is_blk branch no longer exits -- it falls through to fmt=1 (no explicit assignment needed: fmt is already 1 entering this branch, and the is_blk&&!is_row check further down evaluates false whenever is_row is true, per the comment left in its place). Sound specifically because the WRITER now refuses to ever emit an fmt=7 container at a shape satisfying this same predicate (see repack_fp8_passthrough.py below): an unstamped ambiguous tensor reaching this function is never a genuine fmt=7 candidate, so resolving to the incumbent format isn't a guess. The SECOND DESIGN LANDMINE (fmt=6 vs fmt=7 collision at I=98) and the UE8M0 recognized-not-implemented refusals are unrelated collision predicates and are unchanged -- still unconditional refusals. tools/repack_fp8_passthrough.py: _check_geometry() now also refuses to repack any [O,I] where nblkO*nblkI==O -- the exact shape the reader's collision predicate describes -- making the module docstring's existing promise ("never anything that could coincide with O*4 by accident") true rather than merely asserted. tools/fp8_collision_census.py (new): enumerates every resident-role tensor shape reachable from a container's config.json + safetensors headers (JSON header parse only, no tensor data, no model load) and evaluates the collision predicate over all of them -- run against this repo's own two GLM-5.2 containers (glm52_v1_e8x4g64, glm52_i4), the complete family is self_attn.o_proj at [6144,16384] (79 instances per container, one per layer including the MTP head layer) and nothing else among gate/up/down, q_a/q_b/kv_a/kv_b_proj, or routed/shared experts. tests/test_fp8_load.c: the three shapes the review named explicitly ([128,16384],[256,16384],[384,16384]) plus every other degenerate case in the same sweep (they all hit the identical is_row&&is_blk branch) flip from expect_refuse to expect_fmt(...,1,...); new explicit regression for the real o_proj shape [6144,16384]. tests/test_fp8_repack.py: new writer-refusal test for a well-formed-but-ambiguous [2,256] fixture (the smallest instance of the same family, distinct from the pre-existing malformed-shape refusal test). RAN (this commit): make clean && make portable -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0. make test-python -> 153 tests, 10 skipped, 0 failures, exit 0. make glm METAL=1 -> clean rebuild, zero warnings. make metal-test under COLI_METAL_RESSET=0 AND =1: both full suite green, exit 0. --- c/colibri.c | 49 ++++++-- c/tests/test_fp8_load.c | 79 +++++++++--- c/tests/test_fp8_repack.py | 40 ++++++- c/tools/fp8_collision_census.py | 191 ++++++++++++++++++++++++++++++ c/tools/repack_fp8_passthrough.py | 45 +++++-- 5 files changed, 364 insertions(+), 40 deletions(-) create mode 100644 c/tools/fp8_collision_census.py diff --git a/c/colibri.c b/c/colibri.c index 0a7445c7..922c5bf3 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -1209,11 +1209,34 @@ static int qt_resolve_fmt(const char *name, int O, int I, int64_t nb, int64_t ns * unambiguous. But for small O (<=128) and/or small I, ceil(O/128)* * ceil(I/128) can equal O exactly (e.g. O=1,I<=128 -> 1*1=1=O; O=2,I in * [129,256] -> 1*2=2=O; O=256,I in (16256,16384] -> 2*128=256=O) -- a - * tensor whose scale array satisfies BOTH conventions at once. Guessing - * either way risks silently reading a genuine per-row-int8 tensor as - * block-scaled FP8 (or vice versa), which would corrupt every weight - * without any error. Refuse instead (same "untrusted container" - * discipline as the rest of this function). + * tensor whose scale array satisfies BOTH conventions at once. + * + * REVIEW FINDING (maintainer, #528): this collision is not a theoretical + * corner case. GLM-5.2's own self_attn.o_proj.weight loads as + * [D,H*v_head] = [6144,16384]: nblkO=ceil(6144/128)=48, + * nblkI=ceil(16384/128)=128, nblkO*nblkI=6144==O -- a real, pre-existing, + * VALID int8-row o_proj tensor hits this exact byte-count collision on + * every GLM-5.2 checkpoint this engine has ever loaded. The general + * family is any int8 tensor where O==ceil(O/128)*ceil(I/128) (see the + * CENSUS SCAN, tools/fp8_collision_census.py, for the complete + * enumerated set over this repo's own containers). An earlier revision + * of this function refused unconditionally here -- exit(1) at load time + * on an ordinary, already-shipping, valid model -- which is a strictly + * worse failure mode than the misread it was guarding against. + * + * INVERSION: an ambiguous shape resolves to fmt=1 -- the incumbent, + * already-on-disk, decodable format -- instead of refusing. This is sound + * because the WRITER side (tools/repack_fp8_passthrough.py's + * _check_geometry) now refuses to EMIT an fmt=7 container at any shape + * satisfying this same collision predicate: no genuine fmt=7 tensor this + * engine's own tooling can produce will ever reach this branch, so + * resolving to fmt=1 here is not a guess against a live fmt=7 candidate -- + * it is the only remaining candidate for an UNSTAMPED container. A + * self-describing __metadata__ stamp that lets a THIRD-PARTY fmt=7 + * container resolve at this same colliding shape is a separate mechanism + * (see qt_verify_fmt_stamp / the stamped build of this function) -- this + * (unstamped) function has no stamp to consult and always commits to the + * incumbent format. * * SCALE ENCODING IS A DECLARED PROPERTY, not a hardcoded constant: f32 (4 * bytes/block, above) is what THIS build implements, but it is not the @@ -1247,12 +1270,16 @@ static int qt_resolve_fmt(const char *name, int O, int I, int64_t nb, int64_t ns int64_t nblkO=fp8_nblk(O), nblkI=fp8_nblk(I); int64_t ns_row=(int64_t)O*4, ns_blk=nblkO*nblkI*4, ns_blk_ue8m0=nblkO*nblkI; int is_row=(ns==ns_row), is_blk=(ns==ns_blk), is_blk_ue8m0=(ns==ns_blk_ue8m0); - if(is_row && is_blk){ - fprintf(stderr,"%s: [%d,%d] scale array is %lld bytes — matches BOTH per-row " - "int8 (fmt=1) and per-128x128-block FP8 (fmt=7) scale geometry; refusing " - "rather than guessing (untrusted container, THE DESIGN LANDMINE)\n", - name,O,I,(long long)ns); exit(1); - } + /* THE DESIGN LANDMINE, INVERTED (maintainer review, #528): is_row&&is_blk + * used to exit(1) here unconditionally -- see the comment above this + * function's "REVIEW FINDING"/"INVERSION" paragraphs for why that was + * wrong. No explicit branch is needed to resolve it to fmt=1: `fmt` is + * already 1 at this point (this whole `if(fmt==1)` block only runs when + * nb==exp_i8, which is what set fmt=1 above), and the `is_blk && !is_row` + * check further down evaluates false whenever is_row is true -- so an + * is_row&&is_blk shape simply falls through unchanged to fmt=1. This + * comment exists so a future reader doesn't mistake the missing branch + * for an oversight. */ if(is_blk_ue8m0){ fprintf(stderr,"%s: [%d,%d] fp8-e4m3-b128 with ue8m0 scales recognized but not " "implemented; only f32 block scales are supported in this build (nb=%lld " diff --git a/c/tests/test_fp8_load.c b/c/tests/test_fp8_load.c index f04cca02..49ca2e67 100644 --- a/c/tests/test_fp8_load.c +++ b/c/tests/test_fp8_load.c @@ -14,18 +14,32 @@ * The two are told apart ONLY by the scale array's byte count (per-row O*4 for * fmt=1, per-128x128-block ceil(O/128)*ceil(I/128)*4 for fmt=7, THIS build's * implemented f32 scale encoding). For some shapes those two counts coincide - * exactly -- qt_resolve_fmt must REFUSE (exit(1)) rather than guess. Refusal - * is tested via fork()+waitpid(), mirroring tests/test_st_pread.c's - * established convention for exit(1)-terminated paths. + * exactly -- INVERSION (maintainer review, #528): qt_resolve_fmt used to + * REFUSE (exit(1)) this ambiguous case; it now resolves to fmt=1 (the + * incumbent, already-on-disk, decodable format) instead, because the + * collision is not hypothetical (GLM-5.2's own self_attn.o_proj.weight hits + * it, see qt_resolve_fmt's own "REVIEW FINDING"/"INVERSION" comment) and the + * writer side (repack_fp8_passthrough.py's _check_geometry) now refuses to + * ever EMIT an fmt=7 container at this same shape, so an unstamped ambiguous + * tensor reaching this function is never a genuine fmt=7 candidate. The + * former refusal-testing convention (fork()+waitpid(), mirroring + * tests/test_st_pread.c's exit(1)-path idiom) is kept for the OTHER + * refusing cases below (Part A2's fmt=6 collision, Part A3's UE8M0 + * recognized-not-implemented refusal, and the generic garbage-byte-count + * refusal) -- only the is_row&&is_blk collision in this Part flipped from + * expect_refuse to expect_fmt(...,1,...). * * Part A2: fmt=6 (E8/IQ3, upstream #465) vs fmt=7 collision at [O<=128 or - * O in a 128-block-count range, I=98] -- SECOND DESIGN LANDMINE. + * O in a 128-block-count range, I=98] -- SECOND DESIGN LANDMINE. Unchanged by + * the #528 inversion above (a different collision predicate, still refused). * * Part A3: fmt=7's scale ENCODING is a declared property, not a hardcoded * constant -- f32 (Part A/A2 above) is what this build implements. A UE8M0 * (1 byte/block) encoding is a REAL, distinct byte signature (the DeepSeek-V4 * checkpoint format for this identical weight geometry) this build recognizes - * and refuses BY NAME rather than misreading. + * and refuses BY NAME rather than misreading. Unchanged by the #528 + * inversion (a stamp confirms the WEIGHT format, never a decoder this build + * doesn't have -- see qt_resolve_fmt's own comment). * * Part B: qt_from_disk loader-seam -- writes a real single-shard .safetensors * file containing an fmt=7 tensor (U8 weight + per-block F32 .qs) next to an @@ -34,12 +48,18 @@ * fmt=1 correctly and the loaded weights dequantize identically to a reference. * Mirrors tests/test_int3_load.c's structure for fmt=5. * - * Part C: qt_bytes()/qt_scale_bytes() byte-accounting for fmt=7. + * Part C: qt_bytes()/qt_scale_bytes() byte-accounting for fmt=7, plus + * qt_wire_split() -- the shared weight/scale byte-range split qt_wire_mmap + * and qt_unwire_mmap both now call (maintainer review, #528: qt_scale_bytes() + * existed correctly but neither call site actually used it, a defect only + * -Wno-unused-function's suppression let compile clean). * * This build has NO container metadata stamp (see qt_resolve_fmt's own header - * comment) -- every ambiguous/unimplemented-encoding shape below refuses - * unconditionally; stamp-resolves-ambiguity behavior is a separate, follow-up - * PR (registry + metadata stamp), not exercised here. */ + * comment) -- every ambiguous/unimplemented-encoding shape below EXCEPT the + * is_row&&is_blk collision (now resolved to fmt=1, see Part A above) refuses + * unconditionally; stamp-resolves-ambiguity behavior for the OTHER + * collisions is a separate, follow-up PR (registry + metadata stamp), not + * exercised here. */ #define main coli_glm_main_unused #include "../colibri.c" #undef main @@ -123,21 +143,46 @@ static void test_disambiguation(void){ CHECK(expect_fmt(2048,6144,(int64_t)2048*6144,16LL*48*4,7,"fp8 2048x6144 (spec example)")); CHECK(expect_fmt(384,6144,(int64_t)384*6144,3LL*48*4,7,"fp8 384x6144 non-square block grid")); + /* --- REGRESSION (maintainer review, #528): GLM-5.2's own + * self_attn.o_proj.weight, [D,H*v_head]=[6144,16384]. nblkO=ceil(6144/128)=48, + * nblkI=ceil(16384/128)=128, product=6144==O -- this shape hits the + * is_row&&is_blk collision below on every GLM-5.2 checkpoint, and is a + * REAL, pre-existing, valid int8-row tensor, not a hypothetical. Confirmed + * against this repo's own v1 container header (o_proj weight U8 + * 100,663,296 B == 6144*16384; .qs scale blob 24,576 B == 6144*4 == + * 48*128*4, both at once) -- see the CENSUS SCAN + * (tools/fp8_collision_census.py) for the full enumerated family. An + * earlier revision of qt_resolve_fmt refused this shape unconditionally + * (exit(1) at load time on an ordinary model); the INVERSION resolves it + * to fmt=1 instead -- this is that non-refusal, asserted directly. */ + CHECK(expect_fmt(6144,16384,(int64_t)6144*16384,6144LL*4,1, + "GLM-5.2 o_proj shape [6144,16384]: valid int8-row, ambiguous-by-byte-count, resolves fmt=1 (NOT a refusal)")); + /* --- degenerate shapes: O<=128 makes nblkO==1, so ns_blk==nblkI*4 can * equal ns_row==O*4 whenever nblkI==O. Exhaustive-in-spirit sweep of the - * boundary the design landmine describes. --- */ - CHECK(expect_refuse(1,1, 1, 4, "degenerate O=1 I=1 (nblkI=1=O)")); - CHECK(expect_refuse(1,128, 128, 4, "degenerate O=1 I=128 (nblkI=1=O, I at block edge)")); - CHECK(expect_refuse(2,256, 2LL*256, 8, "degenerate O=2 I=256 (nblkI=2=O)")); - CHECK(expect_refuse(6,768, 6LL*768, 24, "degenerate O=6 I=768 (nblkI=6=O)")); - CHECK(expect_refuse(128,16384,128LL*16384, 512, "degenerate O=128 I=16384 (nblkO=1,nblkI=128=O)")); + * boundary the design landmine describes. INVERSION (maintainer review, + * #528): every one of these used to be an expect_refuse (exit(1)) -- the + * general family this collision predicate describes (any int8 tensor + * where O==ceil(O/128)*ceil(I/128)) includes real, valid, pre-existing + * tensors (see the o_proj case just above), so refusing was the wrong + * call across the board, not just for the three shapes the maintainer's + * review named explicitly ([128,16384],[256,16384],[384,16384] below) -- + * every case in this sweep hits the exact same is_row&&is_blk branch and + * is flipped here for the same reason. --- */ + CHECK(expect_fmt(1,1, 1, 4, 1, "degenerate O=1 I=1 (nblkI=1=O) -> fmt=1, was refuse")); + CHECK(expect_fmt(1,128, 128, 4, 1, "degenerate O=1 I=128 (nblkI=1=O, I at block edge) -> fmt=1, was refuse")); + CHECK(expect_fmt(2,256, 2LL*256, 8, 1, "degenerate O=2 I=256 (nblkI=2=O) -> fmt=1, was refuse")); + CHECK(expect_fmt(6,768, 6LL*768, 24, 1, "degenerate O=6 I=768 (nblkI=6=O) -> fmt=1, was refuse")); + /* the three shapes the maintainer's review named explicitly (his exact expect_refuse + * calls, O=128/256/384 at I=16384) -- FLIPPED to assert fmt=1. */ + CHECK(expect_fmt(128,16384,128LL*16384, 512, 1, "degenerate O=128 I=16384 (nblkO=1,nblkI=128=O) -> fmt=1, was refuse")); /* O>128 degenerate case: nblkO=2, need nblkI=O/2 -- O=256,I=16384 -> nblkI=128, 2*128=256=O */ - CHECK(expect_refuse(256,16384,256LL*16384, 1024, "degenerate O=256 I=16384 (nblkO=2,nblkI=128, product=O)")); + CHECK(expect_fmt(256,16384,256LL*16384, 1024, 1, "degenerate O=256 I=16384 (nblkO=2,nblkI=128, product=O) -> fmt=1, was refuse")); /* k=3: the ambiguity isn't a one-off O=256 coincidence -- it's structural for ANY * O that's a multiple of 128 (nblkO=k), since nblkI==128 (I in (16256,16384]) always * makes nblkO*nblkI == k*128 == O. One more multiple (O=384=3*128) confirms the * condition generalizes past the k=2 worked example, not just a re-derivation. */ - CHECK(expect_refuse(384,16384,384LL*16384, 1536, "degenerate O=384 I=16384 (nblkO=3,nblkI=128, k=3, product=O)")); + CHECK(expect_fmt(384,16384,384LL*16384, 1536, 1, "degenerate O=384 I=16384 (nblkO=3,nblkI=128, k=3, product=O) -> fmt=1, was refuse")); /* --- boundary-ADJACENT non-degenerate cases: one step past each * degenerate case above, both interpretations now legitimately resolve. --- */ diff --git a/c/tests/test_fp8_repack.py b/c/tests/test_fp8_repack.py index 283de225..64bb3dc3 100644 --- a/c/tests/test_fp8_repack.py +++ b/c/tests/test_fp8_repack.py @@ -11,10 +11,14 @@ read or written by this suite, per the build's hard constraint. Covers: selection (resident kinds byte-preserved, routed experts / io / f32 excluded), byte-for-byte preservation of the fp8 weight and the .qs scale rename, the -scale-geometry refusal path (THE DESIGN LANDMINE's write-side twin: a -malformed source shard must be refused, not silently repacked into a -container the engine's qt_resolve_fmt would then misread), --dry-run writing -nothing, and the #383-class resume/params-guard behavior. +scale-geometry refusal path in TWO forms (THE DESIGN LANDMINE's write-side +twin) -- a malformed source shard (.qs shape doesn't match ceil(O/128)x +ceil(I/128) for its weight) must be refused, AND a well-formed-but-AMBIGUOUS +shape (nblkO*nblkI==O, where this tensor's block-scale byte count would +coincide with a per-row int8 scale byte count) must ALSO be refused +(maintainer review, #528: this is not hypothetical, GLM-5.2's own +self_attn.o_proj.weight is a real instance) -- --dry-run writing nothing, and +the #383-class resume/params-guard behavior. """ import glob, json, os, struct, sys, tempfile, unittest @@ -153,6 +157,34 @@ def test_geometry_refusal_on_malformed_scale(self): with self.assertRaises(ValueError): rp.shard_inventory(bad_path, n_layers=5) + def test_geometry_refusal_on_ambiguous_shape(self): + """WRITER-SIDE REFUSAL (maintainer review, #528): a WELL-FORMED .qs sidecar + (correct nblkOxnblkI shape for [O,I], unlike the malformed case above) must + still be refused when nblkO*nblkI==O -- the exact shape where this tensor's + block-scale byte count (nblkO*nblkI*4) would coincide with a per-row int8 + scale byte count (O*4). Not hypothetical: GLM-5.2's own + self_attn.o_proj.weight ([6144,16384], nblkO=48 nblkI=128 product=6144==O) + is a real instance of this exact family; O=2,I=256 here is the smallest + realistic analog (nblkO=1, nblkI=2, product=2==O), the same degenerate + shape family c/tests/test_fp8_load.c's test_disambiguation() sweeps. The + engine's qt_resolve_fmt reader resolves an unstamped collision at this + shape to fmt=1 (int8-row) rather than refusing (see colibri.c's + INVERSION) -- so this tool refusing to ever EMIT an fmt=7 container here + is what keeps that reader-side resolution correct: an fmt=7 tensor this + tool produced at this shape would be silently read back as plain int8.""" + amb_path = os.path.join(self.tmp.name, "ambiguous.safetensors") + O, I_ = 2, 256 + w = torch.randn(O, I_).to(torch.float8_e4m3fn) + nblkO, nblkI = (O + 127) // 128, (I_ + 127) // 128 + self.assertEqual(nblkO * nblkI, O, "fixture must actually hit the collision") + scale = torch.ones(nblkO, nblkI, dtype=torch.float32) # WELL-FORMED for [O,I] -- passes the first check + save_file({"model.layers.0.self_attn.o_proj.weight": w, + "model.layers.0.self_attn.o_proj.weight_scale_inv": scale}, amb_path) + with self.assertRaises(ValueError): + rp.repack_shard(amb_path, n_layers=5) + with self.assertRaises(ValueError): + rp.shard_inventory(amb_path, n_layers=5) + class ResumeAndParamsGuardTest(unittest.TestCase): def setUp(self): diff --git a/c/tools/fp8_collision_census.py b/c/tools/fp8_collision_census.py new file mode 100644 index 00000000..99ea37f6 --- /dev/null +++ b/c/tools/fp8_collision_census.py @@ -0,0 +1,191 @@ +#!/usr/bin/env python3 +"""fp8_collision_census.py: enumerate every RESIDENT-role tensor shape in a +safetensors container (JSON header parse ONLY -- no tensor data read, no +torch/safetensors import, no model load) and evaluate THE DESIGN LANDMINE's +collision predicate (qt_resolve_fmt in colibri.c): O == ceil(O/128)*ceil(I/128), +the shape family where an int8-row tensor's per-row scale byte count (O*4) +exactly equals a per-128x128-block fp8 scale byte count +(ceil(O/128)*ceil(I/128)*4) for the SAME [O,I]. + +Maintainer review, #528: this is not a hypothetical corner case. GLM-5.2's own +self_attn.o_proj.weight ([D,H*v_head] = [6144,16384]) hits it on every +checkpoint this engine has loaded. This script is the "enumerated, not +sampled" evidence for that finding, run against this repo's own containers. + +WHY config.json, not just the safetensors header: a container's on-disk +weight tensor is stored FLATTENED (one dimension: O*I raw bytes for int8-row, +O*ceil(I/2) for int4-packed, etc. -- see colibri.c's qt_bytes()) -- the header +alone cannot recover O and I separately for [O,I]. This script instead +computes each tensor's [O,I] from the model's config.json (hidden_size, +q_lora_rank, kv_lora_rank, num_attention_heads, qk_rope_head_dim, +qk_nope_head_dim, v_head_dim, intermediate_size, moe_intermediate_size -- the +same dimensions colibri.c's own Cfg struct reads at load time and passes into +qt_resolve_fmt as O,I), for every tensor NAME the header actually contains +that this script's role table recognizes, then cross-checks the computed +byte count against the header's own raw byte count (int8-row: O*I; int4- +packed: O*ceil(I/2)) as a sanity check -- reported, not silently trusted. + +Scope: RESIDENT-kind tensor roles only (o_proj, q_a/q_b_proj, +kv_a_proj_with_mqa, kv_b_proj, dense-MLP gate/up/down, shared-expert +gate/up/down) plus routed-expert gate/up/down for completeness -- these are +the only roles tools/repack_fp8_passthrough.py's RESIDENT_KINDS classify() +taxonomy could ever select for fmt=7 repacking (routed experts are excluded +by that tool's own design and stay on int4-g64/E8, but are included here +anyway since the collision predicate is a pure function of [O,I], independent +of current on-disk format). f32/norm/router/embed/lm_head tensors are not +quantized (no U8 weight) and are skipped. + +Usage: + python3 tools/fp8_collision_census.py [ ...] +""" +import glob, json, os, struct, sys + +BLOCK = 128 + + +def _nblk(n): + return (n + BLOCK - 1) // BLOCK + + +def is_ambiguous(O, I): + """THE DESIGN LANDMINE's collision predicate (colibri.c qt_resolve_fmt): + an int8-row tensor at this [O,I] has per-row scale bytes O*4 that exactly + equal per-128x128-block fp8 scale bytes ceil(O/128)*ceil(I/128)*4.""" + return _nblk(O) * _nblk(I) == O + + +def read_header(path): + """Header-only read: 8-byte little-endian length prefix + that many bytes + of JSON. Never reads a single byte of actual tensor data.""" + with open(path, "rb") as f: + hlen = struct.unpack(" header entry, across every + shard in the container. Header (JSON) only, see read_header above.""" + entries = {} + for path in sorted(glob.glob(os.path.join(container_dir, "*.safetensors"))): + hdr = read_header(path) + for k, v in hdr.items(): + if k == "__metadata__": + continue + entries[k] = v + return entries + + +def classify_shape(name, cfg): + """Returns (kind, O, I) for a tensor name this census can size from + config.json alone, or None if `name` isn't a role this script recognizes + (norms/router/embed/lm_head/MTP-head-only tensors -- none of these carry + a resident [O,I] weight this predicate applies to).""" + D = cfg["hidden_size"] + H = cfg["num_attention_heads"] + q_lora = cfg["q_lora_rank"] + kv_lora = cfg["kv_lora_rank"] + qk_rope = cfg["qk_rope_head_dim"] + qk_nope = cfg["qk_nope_head_dim"] + v_head = cfg["v_head_dim"] + inter = cfg["intermediate_size"] + moe_inter = cfg["moe_intermediate_size"] + qk_head = qk_nope + qk_rope + + if name.endswith("self_attn.o_proj.weight"): + return ("self_attn.o_proj", D, H * v_head) + if name.endswith("self_attn.q_a_proj.weight"): + return ("self_attn.q_a_proj", q_lora, D) + if name.endswith("self_attn.q_b_proj.weight"): + return ("self_attn.q_b_proj", H * qk_head, q_lora) + if name.endswith("self_attn.kv_a_proj_with_mqa.weight"): + return ("self_attn.kv_a_proj_with_mqa", kv_lora + qk_rope, D) + if name.endswith("self_attn.kv_b_proj.weight"): + return ("self_attn.kv_b_proj (kvb -- excluded from fmt=7 repack, see tool docstring)", + H * (qk_nope + v_head), kv_lora) + if ".mlp.experts." in name and name.endswith("gate_proj.weight"): + return ("mlp.experts.*.gate_proj (routed -- excluded from fmt=7 repack)", moe_inter, D) + if ".mlp.experts." in name and name.endswith("up_proj.weight"): + return ("mlp.experts.*.up_proj (routed -- excluded from fmt=7 repack)", moe_inter, D) + if ".mlp.experts." in name and name.endswith("down_proj.weight"): + return ("mlp.experts.*.down_proj (routed -- excluded from fmt=7 repack)", D, moe_inter) + if "shared_experts" in name and name.endswith("gate_proj.weight"): + return ("mlp.shared_experts.gate_proj", moe_inter, D) + if "shared_experts" in name and name.endswith("up_proj.weight"): + return ("mlp.shared_experts.up_proj", moe_inter, D) + if "shared_experts" in name and name.endswith("down_proj.weight"): + return ("mlp.shared_experts.down_proj", D, moe_inter) + if name.endswith("mlp.gate_proj.weight") and ".experts." not in name: + return ("mlp.gate_proj (dense)", inter, D) + if name.endswith("mlp.up_proj.weight") and ".experts." not in name: + return ("mlp.up_proj (dense)", inter, D) + if name.endswith("mlp.down_proj.weight") and ".experts." not in name: + return ("mlp.down_proj (dense)", D, inter) + return None + + +def _byte_check(nb_actual, O, I): + if nb_actual == O * I: + return "int8-row (nb=O*I)" + if nb_actual == O * ((I + 1) // 2): + return "int4-packed (nb=O*ceil(I/2))" + return f"UNVERIFIED (nb={nb_actual}, neither O*I={O*I} nor O*ceil(I/2)={O*((I+1)//2)} -- " \ + f"likely int3-g64/E8-lattice, expected for routed experts in some containers)" + + +def census(container_dir): + cfg_path = os.path.join(container_dir, "config.json") + with open(cfg_path) as f: + cfg = json.load(f) + entries = scan_tensor_headers(container_dir) + by_pattern = {} + unresolved = [] + for name, entry in entries.items(): + if name.endswith(".qs") or name.endswith("_scale_inv"): + continue + if entry.get("dtype") != "U8": + continue + cls = classify_shape(name, cfg) + if cls is None: + unresolved.append(name) + continue + kind, O, I = cls + shp = entry.get("shape", []) + nb_actual = 1 + for d in shp: + nb_actual *= d + key = (kind, O, I) + rec = by_pattern.setdefault(key, {"kind": kind, "O": O, "I": I, "count": 0, + "ambiguous": is_ambiguous(O, I), + "byte_encodings": set()}) + rec["count"] += 1 + rec["byte_encodings"].add(_byte_check(nb_actual, O, I)) + return cfg, by_pattern, unresolved + + +def main(): + dirs = sys.argv[1:] + if not dirs: + print("usage: fp8_collision_census.py [ ...]", file=sys.stderr) + sys.exit(1) + any_ambig = False + for d in dirs: + base = os.path.basename(os.path.normpath(d)) + cfg, by_pattern, unresolved = census(d) + print(f"=== {base} ({d}) ===") + print(f"{'tensor name pattern':70s} {'O':>8s} {'I':>8s} {'nblkO*nblkI':>12s} {'count':>6s} ambiguous? on-disk encoding(s)") + for (kind, O, I), rec in sorted(by_pattern.items()): + nblk_prod = _nblk(O) * _nblk(I) + flag = "AMBIGUOUS <---" if rec["ambiguous"] else "-" + if rec["ambiguous"]: + any_ambig = True + enc = ", ".join(sorted(rec["byte_encodings"])) + print(f"{kind:70s} {O:8d} {I:8d} {nblk_prod:12d} {rec['count']:6d} {flag:15s} {enc}") + if unresolved: + print(f" ({len(unresolved)} U8 tensor(s) in {base} not classified by this script's role " + f"table -- inspect if unexpected: {sorted(unresolved)[:5]}{'...' if len(unresolved)>5 else ''})") + print() + print("RESULT:", "AMBIGUOUS family found (see table above)" if any_ambig else "no ambiguous family found") + + +if __name__ == "__main__": + main() diff --git a/c/tools/repack_fp8_passthrough.py b/c/tools/repack_fp8_passthrough.py index 6a1a3e0d..704703af 100644 --- a/c/tools/repack_fp8_passthrough.py +++ b/c/tools/repack_fp8_passthrough.py @@ -19,13 +19,23 @@ int8) ONLY by scale-array geometry -- per-row for fmt=1, per-128x128-block for fmt=7. This tool must therefore emit .qs sidecars whose byte count is EXACTLY ceil(O/128)*ceil(I/128)*4, never anything that could coincide with O*4 by -accident for the shapes it targets (see _check_geometry below, which refuses -rather than silently emitting a container the engine might misread). At I=98 -specifically, a single-block fp8 tensor's raw weight-byte count ALSO coincides -with a genuine fmt=6 (E8/IQ3) tensor's -- see qt_resolve_fmt's "SECOND DESIGN -LANDMINE" comment; this tool does not target I=98 shapes in practice (Z.ai's -real projections don't land there), but the engine's read-side refusal covers -it regardless. +accident for the shapes it targets -- ENFORCED, not just asserted, by +_check_geometry below (maintainer review, #528): it refuses to repack any +[O,I] where nblkO*nblkI==O, the exact shape where this tool's own block-scale +byte count would coincide with a per-row int8 scale byte count. This is not +hypothetical: GLM-5.2's own self_attn.o_proj.weight ([6144,16384]) is a real +instance (nblkO=48, nblkI=128, product=6144==O). It matters more now than it +used to, because the engine's UNSTAMPED read-side collision policy resolves an +ambiguous shape to fmt=1 rather than refusing (see qt_resolve_fmt's +INVERSION) -- so an fmt=7 container this tool emitted at such a shape would be +silently read back as plain int8, not caught by a read-side refusal. Avoiding +the collision at write time is therefore load-bearing, not a redundant +belt-and-braces check. At I=98 specifically, a single-block fp8 tensor's raw +weight-byte count ALSO coincides with a genuine fmt=6 (E8/IQ3) tensor's -- see +qt_resolve_fmt's "SECOND DESIGN LANDMINE" comment; this tool does not target +I=98 shapes in practice (Z.ai's real projections don't land there), and that +narrower collision is unchanged by this PR (still an unconditional read-side +refusal), so _check_geometry does not additionally guard against it here. SCALE ENCODING: this tool always emits f32 block scales (4 bytes/block, the `.to(torch.float32)` below) -- the ENCODING this build implements. It never @@ -112,13 +122,32 @@ def _check_geometry(name, O, I, nblkO, nblkI): doesn't match what qt_resolve_fmt expects for [O,I] -- the same "untrusted container" discipline the C side applies on read, applied here on write so a malformed source checkpoint is caught at repack time, not at load time three - steps later with a confusing engine-side refusal.""" + steps later with a confusing engine-side refusal. + + WRITER-SIDE REFUSAL (maintainer review, #528): also refuses at the AMBIGUOUS + shape THE DESIGN LANDMINE describes (qt_resolve_fmt in colibri.c) -- where + O == ceil(O/128)*ceil(I/128), i.e. this tensor's block-scale byte count + (nblkO*nblkI*4) exactly equals a per-row int8 scale byte count (O*4). GLM-5.2's + own self_attn.o_proj.weight ([6144,16384]) is a REAL, non-hypothetical instance + of this shape. The engine's reader now resolves an unstamped collision to fmt=1 + (the incumbent, decodable format) rather than refusing -- which makes it load + bearing that THIS tool never emits an fmt=7 container at such a shape: this + docstring's own opening promise ("never anything that could coincide with O*4 + by accident") is enforced here, not just asserted.""" exp_nblkO, exp_nblkI = _nblk(O), _nblk(I) if (nblkO, nblkI) != (exp_nblkO, exp_nblkI): raise ValueError( f"{name}: scale_inv block shape ({nblkO},{nblkI}) != expected " f"({exp_nblkO},{exp_nblkI}) for [{O},{I}] -- refusing to repack a shape " f"the engine's qt_resolve_fmt would either refuse or (worse) misread") + if nblkO * nblkI == O: + raise ValueError( + f"{name}: [{O},{I}] is an AMBIGUOUS fp8-e4m3-b128 shape -- its " + f"block-scale byte count (nblkO*nblkI*4={nblkO*nblkI*4}) exactly " + f"coincides with a per-row int8 scale byte count (O*4={O*4}) for the " + f"same [O,I]; refusing to emit an fmt=7 container the engine's " + f"qt_resolve_fmt collision policy would resolve to fmt=1 (int8-row), " + f"not fmt=7, on an unstamped read (THE DESIGN LANDMINE, colibri.c)") def shard_inventory(path, n_layers): From de005418423356d6c1b80a84ca674e5658a73166 Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 19:39:50 -0400 Subject: [PATCH 06/12] fp8: wire qt_scale_bytes() into qt_wire_mmap/qt_unwire_mmap Maintainer review, #528: qt_scale_bytes() already existed and had the right fmt=7 branch (per-128x128-block scale count), but neither qt_wire_mmap nor qt_unwire_mmap actually called it -- both independently hardcoded scale_b=(int64_t)t->O*4 (per-row) inline, so the function compiled but was never referenced by its intended call sites. It only compiled clean because -Wno-unused-function is in CFLAGS: the claimed fmt=4/fmt=5/fmt=7 expert-pinning fix was never applied, and an fmt=7 tensor reaching qt_wire_mmap would get weight_b overshooting the allocation whenever nblk>O (mlocking/munlocking the wrong byte ranges on both halves). colibri.c: new qt_wire_split(t,&weight_b,&scale_b) -- the ONE place both call sites now get their weight/scale byte split from, calling qt_scale_bytes() internally. qt_wire_mmap and qt_unwire_mmap both replace their inline scale_b/weight_b computation with a call to it. qt_scale_bytes() is consequently no longer an orphan: it is exercised both directly (existing qt_bytes/qt_scale_bytes arithmetic tests) and through the real call sites via qt_wire_split. tests/test_fp8_load.c: check_wire_split() calls qt_wire_split() directly (no mlock syscall -- mem_wire's actual RLIMIT_MEMLOCK behavior is environment-dependent and not what changed) across fmt=1 (unaffected, per-row), fmt=4/fmt=5 (grouped formats qt_scale_bytes' own comment already named as previously broken too), and fmt=7 at two nblk>O shapes -- each asserting the split matches qt_scale_bytes()/qt_bytes() AND that the fmt=7 cases have actually moved off the old O*4 constant, not just coincidentally matched it. RAN (this commit): make clean && make portable -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0. make test-python -> 153 tests, 10 skipped, 0 failures, exit 0. make glm METAL=1 -> clean rebuild, zero warnings. make metal-test under COLI_METAL_RESSET=0 AND =1: both full suite green, exit 0. One-shot build with -Wno-unused-function removed from CFLAGS: qt_scale_bytes and qt_wire_split are both silent (no longer orphans); 8 unrelated, pre-existing orphan functions remain (st_init, st_prefetch, st_read_slice_f32, tier_pick_swap, stops_arm, cmp_fdesc, attention, and qt_unwire_mmap itself -- the last one only because this translation unit was built without -DCOLI_CUDA, its sole caller's #ifdef guard), all out of this PR's scope -- see the report. --- c/colibri.c | 45 +++++++++++++++++++++++++++-------------- c/tests/test_fp8_load.c | 39 +++++++++++++++++++++++++++++++++++ 2 files changed, 69 insertions(+), 15 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index 922c5bf3..83544db4 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -199,9 +199,9 @@ static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ * qt_wire_mmap/qt_unwire_mmap mlock the weight and scale ranges as TWO SEPARATE * regions (separate allocations, not one contiguous buffer: qt_from_disk always * qalloc's t->q8/t->q4 and t->s independently) and need the scale byte count on - * its own to split qt_bytes(t) into a weight half and a scale half. The two call - * sites used to hardcode scale_b=O*4 -- i.e. assume PER-ROW scale for every - * format -- which is right for fmt 0/1/2/3 but wrong for fmt=4 (grouped, + * its own to split qt_bytes(t) into a weight half and a scale half. Hardcoding + * scale_b=O*4 at those two call sites -- i.e. assuming PER-ROW scale for every + * format -- is right for fmt 0/1/2/3 but wrong for fmt=4 (grouped, * O*ceil(I/gs) scales), fmt=5 (int3-g64, O*ceil(I/64) scales), and fmt=7 * (per-128x128-block, ceil(O/128)*ceil(I/128) scales): any tensor in one of * those three formats reaching qt_wire_mmap/qt_unwire_mmap would mlock/munlock @@ -211,14 +211,21 @@ static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ * qt_bytes' `+4` literal above), and t->fmt==6 tensors never reach * qt_wire_mmap/qt_unwire_mmap's mlock path in this build (int3-g64/E8 stay * CPU-side, see qt_cuda_upload -- MMAP wiring is a CUDA/Metal-adjacent memory - * concern this build doesn't extend to them). fixed generically since the same - * bug shape applies to fmt=4/5, not just the format that surfaced it -- fmt=4 - * is the routed-expert "int4-g64" format (qt_resolve_fmt assigns it to ESlot - * g/u/d the same as any other tensor, see expert_load_impl), so this is not - * purely a dormant fmt=7-only fix: any COLI_MMAP + mem_should_wire() session - * pinning fmt=4 experts was already mlocking the wrong ranges before this - * branch existed. UNVERIFIED at runtime (no model run performed -- static/ - * arithmetic fix only, see the report). */ + * concern this build doesn't extend to them). This is not purely a dormant + * fmt=7-only fix: fmt=4 is the routed-expert "int4-g64" format (qt_resolve_fmt + * assigns it to ESlot g/u/d the same as any other tensor, see + * expert_load_impl), so any COLI_MMAP + mem_should_wire() session pinning + * fmt=4 experts mlocks the wrong ranges without this. UNVERIFIED at runtime + * (no model run performed -- static/arithmetic fix only, see the report). + * + * REVIEW FINDING (maintainer, #528): this function existed correctly (the + * fmt=7 branch below is right) but neither qt_wire_mmap nor qt_unwire_mmap + * actually called it -- both independently hardcoded scale_b=(int64_t)t->O*4 + * inline, so the "fix" this comment describes was never applied at either + * call site, and the resulting -Wunused-function warning on this then-dead + * function was suppressed by -Wno-unused-function in CFLAGS rather than + * caught. qt_wire_split() below is now the ONE place both call sites get + * weight_b/scale_b from -- see it and its call sites for the actual fix. */ static int64_t qt_scale_bytes(const QT *t){ if(t->fmt==4){ int ng=(t->I+t->gs-1)/t->gs; return (int64_t)t->O*ng*4; } if(t->fmt==5){ int64_t ng=((int64_t)t->I+63)/64; return (int64_t)t->O*ng*4; } @@ -227,6 +234,16 @@ static int64_t qt_scale_bytes(const QT *t){ * fmt=6, the 4-byte tag -- callers of this function never see fmt=6, * see the comment above; the fallback is harmless dead code for it). */ } +/* qt_wire_mmap/qt_unwire_mmap's shared weight/scale byte-range split -- the ONE + * place that computes it, so the two sites can never independently drift back + * to a hardcoded per-row-only formula (that drift -- qt_scale_bytes() existing + * but unused at both sites -- was the maintainer's #528 finding, see the + * comment above qt_scale_bytes()). Directly unit-tested (test_fp8_load.c) so a + * regression here, not just at a call site, fails the suite. */ +static void qt_wire_split(const QT *t, int64_t *weight_b, int64_t *scale_b){ + *scale_b = qt_scale_bytes(t); + *weight_b = qt_bytes(t) - *scale_b; +} typedef struct { float *in_ln, *post_ln; @@ -6367,8 +6384,7 @@ static int mem_wire(void *addr, size_t len){ static void qt_unwire_mmap(QT *t){ if(!g_mmap || !mem_should_wire()) return; if(!t->q8 && !t->q4) return; - int64_t scale_b=(int64_t)t->O*4; - int64_t weight_b=qt_bytes(t)-scale_b; + int64_t weight_b, scale_b; qt_wire_split(t,&weight_b,&scale_b); void *wp=t->q8?(void*)t->q8:(void*)t->q4; #if defined(__APPLE__) || defined(__linux__) || defined(__FreeBSD__) if(weight_b>0 && !munlock(wp,(size_t)weight_b)) g_mmap_wired-=weight_b; @@ -6381,8 +6397,7 @@ static void qt_unwire_mmap(QT *t){ static void qt_wire_mmap(QT *t, int64_t *wired, long *failed){ if(!t->q8 && !t->q4) return; if(t->cuda_eligible) return; /* resident in VRAM; host range is dead weight */ - int64_t scale_b=(int64_t)t->O*4; - int64_t weight_b=qt_bytes(t)-scale_b; + int64_t weight_b, scale_b; qt_wire_split(t,&weight_b,&scale_b); void *wp=t->q8?(void*)t->q8:(void*)t->q4; if(weight_b>0){ if(mem_wire(wp,(size_t)weight_b)==0) *wired+=weight_b; else (*failed)++; } if(t->s && scale_b>0){ if(mem_wire(t->s,(size_t)scale_b)==0) *wired+=scale_b; else (*failed)++; } diff --git a/c/tests/test_fp8_load.c b/c/tests/test_fp8_load.c index 49ca2e67..10d84074 100644 --- a/c/tests/test_fp8_load.c +++ b/c/tests/test_fp8_load.c @@ -399,6 +399,40 @@ static void check_fp8_bytes(int O, int I, const char *tag){ if(!(O==1 && I==1)) CHECK(got_total != old_wrong); } +/* qt_wire_mmap()/qt_unwire_mmap() (colibri.c) both compute weight_b/scale_b via + * qt_wire_split(t,&weight_b,&scale_b) -- this calls that EXACT shared function + * directly (no mlock syscall: same reasoning as check_fp8_bytes above, mem_wire's + * actual RLIMIT_MEMLOCK behavior is environment-dependent and not what changed) + * so a regression at qt_wire_split() itself -- e.g. reverting to the + * scale_b=(int64_t)t->O*4 hardcode the maintainer's #528 review found dead-coded + * behind an unused qt_scale_bytes() (compiling clean only because + * -Wno-unused-function suppressed the warning that should have caught it) -- + * fails here. Covers every format qt_wire_split's callers can see in practice + * (fmt=1 per-row unaffected; fmt=4/5 grouped-scale formats qt_scale_bytes' + * own comment names as previously-broken too; fmt=7 nblk>O, the shape this + * review round is about). */ +static void check_wire_split(int fmt, int O, int I, int gs, const char *tag){ + QT t; memset(&t,0,sizeof t); t.fmt=fmt; t.O=O; t.I=I; t.gs=gs; + int64_t want_scale = qt_scale_bytes(&t); + int64_t want_weight = qt_bytes(&t) - want_scale; + int64_t got_weight=-1, got_scale=-1; + qt_wire_split(&t,&got_weight,&got_scale); + if(got_scale != want_scale) + printf("FAIL %s: qt_wire_split scale_b=%lld want=%lld\n", tag, (long long)got_scale, (long long)want_scale); + CHECK(got_scale == want_scale); + if(got_weight != want_weight) + printf("FAIL %s: qt_wire_split weight_b=%lld want=%lld\n", tag, (long long)got_weight, (long long)want_weight); + CHECK(got_weight == want_weight); + /* the regression this test exists to catch: a per-row-only scale_b==O*4 + * must NOT be what qt_wire_split returns for a format whose real scale + * cardinality differs from O (fmt=4/5/7 here) -- confirm the value has + * actually moved off that old constant, not just matched qt_scale_bytes() + * by coincidence at a degenerate shape. */ + int64_t old_wrong_scale = (int64_t)O*4; + if((fmt==4 || fmt==5 || fmt==7) && want_scale != old_wrong_scale) + CHECK(got_scale != old_wrong_scale); +} + int main(void){ test_disambiguation(); test_fmt6_fp8_collision(); @@ -408,6 +442,11 @@ int main(void){ check_fp8_bytes(6144,2048, "qt_bytes fmt=7 down-shaped O=6144 I=2048"); check_fp8_bytes(130,200, "qt_bytes fmt=7 block edges O,I both non-mult-128"); check_fp8_bytes(1,1, "qt_bytes fmt=7 degenerate 1x1"); + check_wire_split(1, 4096,4096, 0, "qt_wire_split fmt=1 plain int8 (per-row scale, unaffected by the fix)"); + check_wire_split(4, 2048,6144, 64, "qt_wire_split fmt=4 grouped int4 (O*ceil(I/gs) scale, not O*4)"); + check_wire_split(5, 2048,6144, 0, "qt_wire_split fmt=5 int3-g64 (O*ceil(I/64) scale, not O*4)"); + check_wire_split(7, 2,16384, 0, "qt_wire_split fmt=7 nblk(128) >> O(2): scale=512B, NOT O*4=8B"); + check_wire_split(7, 2048,6144, 0, "qt_wire_split fmt=7 spec example: scale=3072B, NOT O*4=8192B"); if(fails){ printf("fp8 loader-seam tests: %d FAILED\n", fails); return 1; } printf("fp8 loader-seam tests: ok\n"); return 0; From fec7b09b562c0881795925bdf93a916c89301219 Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 19:40:04 -0400 Subject: [PATCH 07/12] fp8: correct stale Metal fmt=7 guard comments Self-review while responding to the #528 maintainer round: attention_rows' and layer_forward_rows' fmt!=7 guard comments claimed "coli_metal_attn_decode's WP_() macro and the mm_gemv shader it dispatches through (bind_gemv) have no fmt=7 case in this build" -- stale since this PR's own Metal kernel commit (96baf22, "Metal mm_gemv kernel for fmt=7 native FP8-e4m3 passthrough") added exactly that: mm_gemv's shader body has a real, working fmt==7 branch (fp8-e4m3 decode + per-128x128-block scale lookup, backend_metal.mm). The shader is not the gap. The actual reason these two fused-kernel paths still fail closed for fmt=7 is the WP_() macro immediately below each guard: `(q).fmt==1?(const void*)(q).q8:(const void*)(q).q4` picks q8 only for fmt==1 and q4 otherwise -- and q4 is NULL/unallocated for fmt=7 (same convention as fmt=1, see the QT struct comment), so an fmt=7 tensor reaching either fused path through WP_() would hand the kernel a NULL weight pointer, not a shader that misreads bytes as f32. Both guard comments now name the real blocker and note that fixing WP_() and threading fmt=7 through bind_gemv is a deferred follow-up (this PR's own GPU-path note), not attempted in this round. No code changes: the guards (fmt!=7 exclusions) and their behavior are unchanged, only the comments explaining WHY they exist. RAN (this commit): make clean && make portable -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0. make glm METAL=1 -> clean rebuild, zero warnings. --- c/colibri.c | 33 ++++++++++++++++++++++----------- 1 file changed, 22 insertions(+), 11 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index 83544db4..aef24891 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2668,14 +2668,20 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p * would rope every row at position 0 and attend over a 1-token window of the wrong * cache -> greedy decode hits EOS at token 2 (mux answers truncated to 1 token). * Ragged rows take the CPU absorb path below, which reads kvs[s]/positions[s]. - * fmt!=7 guards (q_a/q_b/kv_a/o): coli_metal_attn_decode's WP_() macro and the - * mm_gemv shader it dispatches through (bind_gemv) have no fmt=7 case in this - * build -- an fp8-passthrough tensor reaching here would be silently misread as - * f32 by the shader's default branch. Fail closed to the CPU path below instead - * (matmul_qt_ex there dispatches fmt=7 correctly via matmul_fp8). kv_b is already - * pinned to fmt==2 above for an unrelated reason (its absorb kernel is int4-only); - * these four checks are the same discipline extended to the tensors that flow - * through the generic per-fmt shader. */ + * fmt!=7 guards (q_a/q_b/kv_a/o): STALE-COMMENT FIX -- this PR's own mm_gemv + * Metal shader (backend_metal.mm) DOES have a real fmt==7 branch (fp8-e4m3 + * decode + per-128x128-block scale); the shader is not the gap. The actual + * blocker is the WP_() macro two lines below: it picks q8 only for fmt==1, + * else q4 -- and q4 is NULL/unallocated for fmt=7 (same convention as fmt=1, + * see the QT struct comment), so an fmt=7 tensor reaching + * coli_metal_attn_decode would hand the kernel a NULL weight pointer, not a + * misread-as-f32 shader. Fail closed to the CPU path below instead + * (matmul_qt_ex there dispatches fmt=7 correctly via matmul_fp8); fixing + * WP_() and actually wiring fmt=7 through bind_gemv/coli_metal_attn_decode is + * a deferred follow-up (see this PR's GPU-path note), not done in this + * round. kv_b is already pinned to fmt==2 above for an unrelated reason (its + * absorb kernel is int4-only); these four checks are the same discipline + * extended to the tensors that flow through the shared per-fmt shader. */ if(g_metal_enabled && !kvs && S<=4 && (g_absorb==1||(g_absorb<0&&S<=4)) && m->kv_start[layer]==0 && D==6144 && H==64 && c->q_lora==2048 && c->kv_lora==512 && c->qk_nope==192 && c->qk_rope==64 && vh==256 && l->kv_b.fmt==2 @@ -4517,9 +4523,14 @@ static void layer_forward_rows(Model *m, Layer *l, int li, float *x, int S, int * single Lc/Rc + pos_base contract — see the matching guard in attention_rows. * fmt!=7 guards (q_a/q_b/kv_a/o/sh_gate/sh_up/sh_down): same fail-closed discipline * as attention_rows, extended to the shared-expert MLP this fused kernel also - * covers -- none of these seven bind_gemv-routed tensors has fmt=7 support in this - * build's mm_gemv shader, so any of them carrying fmt=7 must skip the whole fused - * path (CPU below dispatches fmt=7 correctly via matmul_qt_ex/matmul_fp8). */ + * covers. STALE-COMMENT FIX -- the mm_gemv shader these bind_gemv calls dispatch + * through DOES have a real fmt==7 branch (this PR's own Metal kernel commit); the + * actual gap is the WP_() macro below, which picks q8 only for fmt==1 and q4 + * (NULL/unallocated for fmt=7) otherwise, so none of these seven bind_gemv-routed + * tensors can safely carry fmt=7 through this fused path yet -- CPU below + * dispatches fmt=7 correctly via matmul_qt_ex/matmul_fp8. Fixing WP_() and wiring + * fmt=7 through bind_gemv is the same deferred follow-up noted in attention_rows, + * not done in this round. */ if(g_metal_enabled && !kvs && S<=4 && lin_layers && l->sparse && (g_absorb==1||(g_absorb<0&&S<=4)) && m->kv_start[li]==0 && D==6144 && c->n_heads==64 && c->q_lora==2048 && c->kv_lora==512 From f9774dcd9f84d0987a79bd09dce02ee27e49a372 Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 21:15:49 -0400 Subject: [PATCH 08/12] fp8: site-level wire regression test for qt_wire_mmap/qt_unwire_mmap FIX ROUND, validator finding (mutation-proven gap): check_wire_split() exercises qt_wire_split() directly, so it does NOT notice a mutation that reverts ONLY qt_wire_mmap's and qt_unwire_mmap's call sites back to the old inline scale_b=(int64_t)t->O*4 hardcode, leaving qt_wire_split() itself untouched -- the two sites would simply stop calling the now-orphaned-again helper, and nothing existing would catch it. tests/test_fp8_load.c: new observer seam (#define mlock/munlock to per-file names before the colibri.c #include, then define those names as self-contained recorders after it -- both are external POSIX library functions with no body in this translation unit, so there is nothing to accidentally shadow-out-of-existence, unlike an attempted mem_wire()-level shadow which was tried first and found NOT to work: mem_wire is a real, internally-defined function, so renaming it renames its call sites too, and qt_wire_mmap ends up calling the renamed-but-still-real function directly, bypassing anything defined under the freed-up old name -- caught by inspecting the preprocessed output before it could hide a broken test, see the comment left in its place). test_wire_site_regression() calls the real qt_wire_mmap()/qt_unwire_mmap() and asserts, through that seam, that the (addr,len) each site actually passes down to the platform lock call matches qt_scale_bytes()/qt_bytes() -- observing the call sites' own behavior, not a reimplementation of it. POSIX-only (mlock/munlock are only called on the __APPLE__/__linux__/__FreeBSD__ arm; Windows uses compat_mlock/compat_munlock, untouched here, matching this file's existing POSIX-only test-seam convention). PROVEN TO BITE (RAN): applied the exact mutation this test exists to catch -- reverted only the two call sites to `int64_t scale_b=(int64_t) t->O*4; int64_t weight_b=qt_bytes(t)-scale_b;`, leaving qt_wire_split() itself byte-for-byte untouched (diffed against the pre-mutation file to confirm) -- rebuilt, and got: FAIL wire-site regression: qt_wire_mmap's SCALE mlock() call got len=8 want=512 (this is exactly what a reverted scale_b=O*4 hardcode breaks) FAIL wire-site regression: qt_unwire_mmap's SCALE munlock() call got len=8 want=512 (this is exactly what a reverted scale_b=O*4 hardcode breaks) fp8 loader-seam tests: 4 FAILED (4 distinct CHECK() failures -- weight+scale on each of wire/unwire; each line appears repeated in the raw run due to a pre-existing, unrelated stdio-buffering artifact where this file's LATER fork()-based stamp-refusal tests each flush a copy of the still-buffered output when their child calls exit(1) -- confirmed by counting: exactly as many repeats as later fork+exit(1) children, same total whether piped or redirected to a file, and the "N FAILED" summary counts the real 4 distinct failures once each). Reverted the mutation (diffed back to byte-identical with the pre-mutation file) and reran: "fp8 loader-seam tests: ok", exit 0. RAN (this commit): make clean && make portable -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0. make test-python -> 155 tests, 10 skipped, 0 failures, exit 0. make glm METAL=1 -> clean rebuild, zero warnings. make metal-test under COLI_METAL_RESSET=0 AND =1: both full suite green, exit 0. --- c/tests/test_fp8_load.c | 158 +++++++++++++++++++++++++++++++++++++++- 1 file changed, 157 insertions(+), 1 deletion(-) diff --git a/c/tests/test_fp8_load.c b/c/tests/test_fp8_load.c index 10d84074..9a22a69a 100644 --- a/c/tests/test_fp8_load.c +++ b/c/tests/test_fp8_load.c @@ -52,7 +52,15 @@ * qt_wire_split() -- the shared weight/scale byte-range split qt_wire_mmap * and qt_unwire_mmap both now call (maintainer review, #528: qt_scale_bytes() * existed correctly but neither call site actually used it, a defect only - * -Wno-unused-function's suppression let compile clean). + * -Wno-unused-function's suppression let compile clean). check_wire_split() + * exercises qt_wire_split() itself directly; test_wire_site_regression() + * (FIX ROUND, validator finding) additionally exercises the real + * qt_wire_mmap()/qt_unwire_mmap() call sites through a mem_wire()/munlock() + * observer seam (defined right below the #include below) -- a mutation that + * reverts ONLY those two call sites back to the old scale_b=(int64_t)t->O*4 + * hardcode, leaving qt_wire_split() itself untouched, is invisible to + * check_wire_split() but fails test_wire_site_regression() (proven by + * actually running that exact mutation -- see the report). * * This build has NO container metadata stamp (see qt_resolve_fmt's own header * comment) -- every ambiguous/unimplemented-encoding shape below EXCEPT the @@ -60,9 +68,41 @@ * unconditionally; stamp-resolves-ambiguity behavior for the OTHER * collisions is a separate, follow-up PR (registry + metadata stamp), not * exercised here. */ +/* WIRE-SITE REGRESSION SEAM (FIX ROUND, validator finding). First attempt + * (superseded, kept as a note): renaming mem_wire() itself via macro does + * NOT work as an observer seam -- mem_wire is a real, internally-defined + * static function, so the SAME rename that frees up the name "mem_wire" + * for a shadow ALSO renames every CALL SITE (qt_wire_mmap's) to the new + * name, meaning qt_wire_mmap ends up calling the renamed-but-still-real + * function directly, bypassing any shadow defined under the old name + * entirely (confirmed by inspecting the preprocessed output -- caught + * before it could hide a broken test). The seam that actually works is one + * level lower: mlock()/munlock() themselves are EXTERNAL POSIX library + * functions with no body anywhere in this translation unit (only a + * declaration, via , plus mem_wire's/qt_unwire_mmap's own call + * sites) -- renaming them redirects those call sites to a name THIS FILE + * provides its own (self-contained) definition for, with no real + * implementation being shadowed out of existence. This observes the EXACT + * (addr,len) qt_wire_mmap (via mem_wire) and qt_unwire_mmap actually pass + * down to the platform lock/unlock call -- not a reimplementation of what + * they SHOULD pass, and immune to a future mem_wire refactor since the + * seam sits at the syscall boundary, not the wrapper. #ifndef _WIN32: + * mlock/munlock are only called on this `#if defined(__APPLE__) || + * defined(__linux__) || defined(__FreeBSD__)` arm; Windows uses + * compat_mlock/compat_munlock instead (untouched here, matching this + * file's existing POSIX-only test-seam convention -- the fork/pipe/waitpid + * refusal tests below skip analogously on Windows). */ +#ifndef _WIN32 +#define mlock test_mlock_seam +#define munlock test_munlock_seam +#endif #define main coli_glm_main_unused #include "../colibri.c" #undef main +#ifndef _WIN32 +#undef mlock +#undef munlock +#endif #include #include @@ -72,6 +112,38 @@ #include #endif +/* Shadow definitions for the seam above -- must come after the #include so + * mem_wire's (unmodified, real) call to mlock() and qt_unwire_mmap's + * (unmodified, real) call to munlock() have already been renamed to these + * names by the #define above. Neither mlock() nor munlock() has a body + * anywhere in this translation unit (both are declared only, via + * ) -- these are the ONLY definitions the renamed call sites + * can resolve to, both self-contained: returning 0 (success) without + * actually locking/unlocking anything is fine for this test, since the + * buffer under test is never really mlocked in the first place (mirrors + * check_fp8_bytes'/check_wire_split's own stated reasoning that the real + * RLIMIT_MEMLOCK-gated syscall's success is environment-dependent and not + * what any of these tests need to prove). Non-static: both must match the + * extern linkage of the (renamed) declarations already left in + * this translation unit -- a static definition here would conflict with + * that non-static declaration ("static declaration follows non-static + * declaration"). g_seam_wire_* observes mem_wire's (hence qt_wire_mmap's) + * calls; g_seam_unwire_* observes qt_unwire_mmap's direct calls. */ +#ifndef _WIN32 +static void *g_seam_wire_addr[4]; static size_t g_seam_wire_len[4]; static int g_seam_wire_n; +int test_mlock_seam(const void *addr, size_t len){ + if(g_seam_wire_n < 4){ g_seam_wire_addr[g_seam_wire_n]=(void*)addr; g_seam_wire_len[g_seam_wire_n]=len; } + g_seam_wire_n++; + return 0; +} +static void *g_seam_unwire_addr[4]; static size_t g_seam_unwire_len[4]; static int g_seam_unwire_n; +int test_munlock_seam(const void *addr, size_t len){ + if(g_seam_unwire_n < 4){ g_seam_unwire_addr[g_seam_unwire_n]=(void*)addr; g_seam_unwire_len[g_seam_unwire_n]=len; } + g_seam_unwire_n++; + return 0; +} +#endif + static int fails = 0; #define CHECK(c) do{ if(!(c)){ printf("FAIL %s:%d: %s\n", __FILE__, __LINE__, #c); fails++; } }while(0) @@ -433,6 +505,89 @@ static void check_wire_split(int fmt, int O, int I, int gs, const char *tag){ CHECK(got_scale != old_wrong_scale); } +/* ---- Site-level wire regression (FIX ROUND, validator finding: mutation- + * proven gap). check_wire_split() above calls qt_wire_split() directly -- + * it would NOT notice a mutation that reverts ONLY qt_wire_mmap's and + * qt_unwire_mmap's call sites back to the old inline + * `scale_b=(int64_t)t->O*4` hardcode, leaving qt_wire_split() itself + * intact (the two sites would simply stop CALLING the now-orphaned-again + * helper). This test calls the real qt_wire_mmap()/qt_unwire_mmap() + * functions and asserts, through the mlock()/munlock() observer seam + * defined above (right after the #include), that the (addr,len) each site + * ACTUALLY passed down to the platform lock call matches + * qt_scale_bytes()/qt_bytes() -- i.e. it observes the call sites' own + * behavior, not a reimplementation of it. Shape chosen so a reverted site + * is unmistakably wrong, not coincidentally right: O=2, I=16384 -> + * nblkO=1, nblkI=128, nblk=128, scale_b=512B; the old hardcode would + * compute O*4=8B instead, a 64x difference. Proven to bite: see the + * report's mutation output (the exact single-line revert, applied and + * reverted, with the resulting FAIL line pasted in). POSIX-only (like this + * file's fork-based refusal tests): the observer seam only intercepts the + * `#if defined(__APPLE__) || defined(__linux__) || defined(__FreeBSD__)` + * arm both qt_wire_mmap (via mem_wire) and qt_unwire_mmap take -- Windows + * takes a different call (compat_mlock/compat_munlock) not intercepted + * here. */ +static void test_wire_site_regression(void){ +#ifndef _WIN32 + g_seam_wire_n = 0; g_seam_unwire_n = 0; + enum { O=2, I=16384 }; + enum { NBLKO = CDIV(O,128), NBLKI = CDIV(I,128), NBLK = NBLKO*NBLKI }; + static uint8_t q7[O*I]; static float s7[NBLK]; + for(int i=0;i= 1 && (int64_t)g_seam_wire_len[0] != want_weight) + printf("FAIL wire-site regression: qt_wire_mmap's WEIGHT mlock() call got len=%lld want=%lld\n", + (long long)g_seam_wire_len[0], (long long)want_weight); + CHECK(g_seam_wire_n>=1 && (int64_t)g_seam_wire_len[0]==want_weight); + if(g_seam_wire_n >= 2 && (int64_t)g_seam_wire_len[1] != want_scale) + printf("FAIL wire-site regression: qt_wire_mmap's SCALE mlock() call got len=%lld want=%lld " + "(this is exactly what a reverted scale_b=O*4 hardcode breaks)\n", + (long long)g_seam_wire_len[1], (long long)want_scale); + CHECK(g_seam_wire_n>=2 && (int64_t)g_seam_wire_len[1]==want_scale); + + /* qt_unwire_mmap early-returns unless g_mmap && mem_should_wire() are + * both true -- force both on for this call, then restore, so this test + * doesn't change global state for any test that runs after it. */ + int saved_g_mmap = g_mmap, saved_g_mlock = g_mlock; + g_mmap = 1; g_mlock = 1; + qt_unwire_mmap(&t); + g_mmap = saved_g_mmap; g_mlock = saved_g_mlock; + + if(g_seam_unwire_n != 2) + printf("FAIL wire-site regression: qt_unwire_mmap called munlock() %d times, expected 2\n", g_seam_unwire_n); + CHECK(g_seam_unwire_n == 2); + if(g_seam_unwire_n >= 1 && (int64_t)g_seam_unwire_len[0] != want_weight) + printf("FAIL wire-site regression: qt_unwire_mmap's WEIGHT munlock() call got len=%lld want=%lld\n", + (long long)g_seam_unwire_len[0], (long long)want_weight); + CHECK(g_seam_unwire_n>=1 && (int64_t)g_seam_unwire_len[0]==want_weight); + if(g_seam_unwire_n >= 2 && (int64_t)g_seam_unwire_len[1] != want_scale) + printf("FAIL wire-site regression: qt_unwire_mmap's SCALE munlock() call got len=%lld want=%lld " + "(this is exactly what a reverted scale_b=O*4 hardcode breaks)\n", + (long long)g_seam_unwire_len[1], (long long)want_scale); + CHECK(g_seam_unwire_n>=2 && (int64_t)g_seam_unwire_len[1]==want_scale); +#else + printf("skipped on Windows (no mlock/munlock observer seam): wire-site regression\n"); +#endif +} + int main(void){ test_disambiguation(); test_fmt6_fp8_collision(); @@ -447,6 +602,7 @@ int main(void){ check_wire_split(5, 2048,6144, 0, "qt_wire_split fmt=5 int3-g64 (O*ceil(I/64) scale, not O*4)"); check_wire_split(7, 2,16384, 0, "qt_wire_split fmt=7 nblk(128) >> O(2): scale=512B, NOT O*4=8B"); check_wire_split(7, 2048,6144, 0, "qt_wire_split fmt=7 spec example: scale=3072B, NOT O*4=8192B"); + test_wire_site_regression(); if(fails){ printf("fp8 loader-seam tests: %d FAILED\n", fails); return 1; } printf("fp8 loader-seam tests: ok\n"); return 0; From 06a8631ea3cc37ad29213c0374d81bdb325b76e6 Mon Sep 17 00:00:00 2001 From: monotophic Date: Wed, 22 Jul 2026 21:16:29 -0400 Subject: [PATCH 09/12] fp8: fmt=6 (E8/IQ3) branch in qt_scale_bytes() FIX ROUND, audit finding (SHOULD-FIX): qt_scale_bytes() had no explicit fmt=6 branch and fell back to the per-row O*4 formula. Verified at source before implementing (per instructions, not trusting the audit's own derivation): qt_from_disk's fmt==6 branch allocates t->s via qsalloc(1) -- ONE float, 4 bytes total, independent of O -- and its st_read_f32_cap call right below validates the SAME cardinality (`fmt==6 ? (int64_t)1 : ...`). qt_bytes()'s own fmt==6 branch already treats the scale as a fixed `+4` literal, never O*4. All three sites agree: fmt=6's scale allocation is a fixed 4 bytes, confirming the audit's claim rather than contradicting it. This function's own header comment used to ALSO claim fmt=6 tensors "never reach qt_wire_mmap/qt_unwire_mmap's mlock path in this build" because they "stay CPU-side" -- that conflated two unrelated mechanisms: qt_cuda_upload decides GPU-VRAM upload eligibility (fmt=5/6 genuinely excluded there, no CUDA kernel for either), while qt_wire_mmap/qt_unwire_mmap mlock HOST RAM pages under COLI_MMAP, an entirely CPU-side concern -- staying CPU-side is exactly what makes a tensor a wiring CANDIDATE, not exempt from it. expert_load_impl assigns fmt=6 to ESlot g/u/d exactly like fmt=4/5 (same qt_resolve_fmt call site), and pin_wire calls qt_wire_mmap/qt_unwire_mmap on every pinned ESlot's g/u/d unconditionally, with no format filter -- so a pinned E8/IQ3 expert under COLI_MMAP + mem_should_wire() does reach here. Before this fix, that path's O*4 fallback would return a scale_b wildly larger than the tensor's real 4-byte allocation, then mlock/munlock that many bytes starting at t->s -- past the end of a 4-byte allocation. This is a strict correction, not scope creep introduced by wiring the helper: pre-existing behavior at the old hardcoded call sites (scale_b=(int64_t)t->O*4, inline, before qt_wire_split existed) was identically wrong for fmt=6 already -- this closes the same gap the #528 fix closed for fmt=4/5/7, for the one format that fix's own review round didn't reach. colibri.c: qt_scale_bytes() gains `if(t->fmt==6) return 4;`, and the header comment's stale reachability claim is corrected. tests/test_fp8_load.c: check_wire_split(6, ...) at a realistic expert shape (verifying the value moved off the old O*4 fallback) and at the O=1 degenerate shape (where O*4 coincidentally also equals 4, so only the new direct `want_scale==4` assertion -- not the generic "moved off O*4" guard -- catches a regression there; both are asserted). PROVEN TO BITE (RAN): removed the `if(t->fmt==6) return 4;` line only, rebuilt, ran the suite: FAIL tests/test_fp8_load.c:574: want_scale == 4 fp8 loader-seam tests: N FAILED Reverted (diffed back to byte-identical with the pre-mutation file) and reran clean: "fp8 loader-seam tests: ok", exit 0. RAN (this commit): make clean && make portable -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0. make test-python -> 155 tests, 10 skipped, 0 failures, exit 0. make glm METAL=1 -> clean rebuild, zero warnings. make metal-test under COLI_METAL_RESSET=0 AND =1: both full suite green, exit 0. One-shot -Wno-unused-function-stripped build: same 8 pre-existing, unrelated orphans as before this round, no new ones. --- c/colibri.c | 60 ++++++++++++++++++++++++++++------------- c/tests/test_fp8_load.c | 17 +++++++++--- 2 files changed, 54 insertions(+), 23 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index aef24891..874a6fca 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -202,21 +202,17 @@ static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ * its own to split qt_bytes(t) into a weight half and a scale half. Hardcoding * scale_b=O*4 at those two call sites -- i.e. assuming PER-ROW scale for every * format -- is right for fmt 0/1/2/3 but wrong for fmt=4 (grouped, - * O*ceil(I/gs) scales), fmt=5 (int3-g64, O*ceil(I/64) scales), and fmt=7 - * (per-128x128-block, ceil(O/128)*ceil(I/128) scales): any tensor in one of - * those three formats reaching qt_wire_mmap/qt_unwire_mmap would mlock/munlock - * the wrong byte ranges on BOTH halves (weight_b = qt_bytes(t)-scale_b shifts - * too). fmt=6 (E8) is deliberately NOT covered here: its .qs is a fixed 4-byte - * tag with no wire-mmap-relevant weight/scale split of its own (matches - * qt_bytes' `+4` literal above), and t->fmt==6 tensors never reach - * qt_wire_mmap/qt_unwire_mmap's mlock path in this build (int3-g64/E8 stay - * CPU-side, see qt_cuda_upload -- MMAP wiring is a CUDA/Metal-adjacent memory - * concern this build doesn't extend to them). This is not purely a dormant - * fmt=7-only fix: fmt=4 is the routed-expert "int4-g64" format (qt_resolve_fmt - * assigns it to ESlot g/u/d the same as any other tensor, see - * expert_load_impl), so any COLI_MMAP + mem_should_wire() session pinning - * fmt=4 experts mlocks the wrong ranges without this. UNVERIFIED at runtime - * (no model run performed -- static/arithmetic fix only, see the report). + * O*ceil(I/gs) scales), fmt=5 (int3-g64, O*ceil(I/64) scales), fmt=6 (E8/IQ3, + * a FIXED 4-byte tag, see below), and fmt=7 (per-128x128-block, + * ceil(O/128)*ceil(I/128) scales): any tensor in one of those four formats + * reaching qt_wire_mmap/qt_unwire_mmap would mlock/munlock the wrong byte + * ranges on BOTH halves (weight_b = qt_bytes(t)-scale_b shifts too). This is + * not purely a dormant fmt=7-only fix: fmt=4 is the routed-expert + * "int4-g64" format (qt_resolve_fmt assigns it to ESlot g/u/d the same as + * any other tensor, see expert_load_impl), so any COLI_MMAP + + * mem_should_wire() session pinning fmt=4 experts mlocks the wrong ranges + * without this. UNVERIFIED at runtime (no model run performed -- static/ + * arithmetic fix only, see the report). * * REVIEW FINDING (maintainer, #528): this function existed correctly (the * fmt=7 branch below is right) but neither qt_wire_mmap nor qt_unwire_mmap @@ -225,14 +221,40 @@ static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ * call site, and the resulting -Wunused-function warning on this then-dead * function was suppressed by -Wno-unused-function in CFLAGS rather than * caught. qt_wire_split() below is now the ONE place both call sites get - * weight_b/scale_b from -- see it and its call sites for the actual fix. */ + * weight_b/scale_b from -- see it and its call sites for the actual fix. + * + * FIX ROUND (audit finding, SHOULD-FIX): this comment used to ALSO claim + * fmt=6 (E8/IQ3) was "deliberately NOT covered" because "t->fmt==6 tensors + * never reach qt_wire_mmap/qt_unwire_mmap's mlock path in this build + * (int3-g64/E8 stay CPU-side, see qt_cuda_upload -- MMAP wiring is a + * CUDA/Metal-adjacent memory concern this build doesn't extend to them)" -- + * that conflated two UNRELATED mechanisms: qt_cuda_upload decides GPU-VRAM + * upload eligibility (fmt=5/6 genuinely excluded there, no CUDA kernel for + * either -- see qt_cuda_upload's own `fmt==5||fmt==6` guard), while + * qt_wire_mmap/qt_unwire_mmap mlock HOST RAM pages under COLI_MMAP, an + * entirely CPU-side concern. Staying CPU-side (no CUDA kernel) is exactly + * what makes a tensor a CANDIDATE for RAM wiring, not exempt from it. + * expert_load_impl assigns fmt=6 to ESlot g/u/d exactly like fmt=4/5 (same + * qt_resolve_fmt call site, three lines above the fmt=4 comment), and + * pin_wire calls qt_wire_mmap/qt_unwire_mmap on every pinned ESlot's g/u/d + * unconditionally, with no format filter -- so a pinned E8/IQ3 expert under + * COLI_MMAP + mem_should_wire() DOES reach here. Before the fmt=6 branch + * below, the O*4 fallback would have returned a scale_b wildly larger than + * the tensor's real scale allocation (a FIXED 4 bytes, qsalloc(1) in + * qt_from_disk's fmt==6 branch -- confirmed against qt_bytes' own `+4` + * literal and qt_from_disk's st_read_f32_cap cardinality switch, both of + * which already treat fmt=6's scale as exactly 1 float, not O floats), then + * mlock/munlock'd that many bytes starting at t->s -- past the end of a + * 4-byte allocation. UNVERIFIED at runtime (same static/arithmetic-only + * caveat as the fmt=4/5/7 fix above -- no model run performed). */ static int64_t qt_scale_bytes(const QT *t){ if(t->fmt==4){ int ng=(t->I+t->gs-1)/t->gs; return (int64_t)t->O*ng*4; } if(t->fmt==5){ int64_t ng=((int64_t)t->I+63)/64; return (int64_t)t->O*ng*4; } + if(t->fmt==6) return 4; /* E8/IQ3: FIXED 4-byte tag (qsalloc(1)), not O*4 -- see this + * function's own header comment for why this is reachable and + * load-bearing, not dead code. */ if(t->fmt==7){ int64_t nblkO=((int64_t)t->O+127)/128, nblkI=((int64_t)t->I+127)/128; return nblkO*nblkI*4; } - return (int64_t)t->O*4; /* fmt 0 (t->s NULL, guarded by callers)/1/2/3/6: per-row (or, for - * fmt=6, the 4-byte tag -- callers of this function never see fmt=6, - * see the comment above; the fallback is harmless dead code for it). */ + return (int64_t)t->O*4; /* fmt 0 (t->s NULL, guarded by callers)/1/2/3: per-row. */ } /* qt_wire_mmap/qt_unwire_mmap's shared weight/scale byte-range split -- the ONE * place that computes it, so the two sites can never independently drift back diff --git a/c/tests/test_fp8_load.c b/c/tests/test_fp8_load.c index 9a22a69a..9c01b24a 100644 --- a/c/tests/test_fp8_load.c +++ b/c/tests/test_fp8_load.c @@ -481,8 +481,10 @@ static void check_fp8_bytes(int O, int I, const char *tag){ * -Wno-unused-function suppressed the warning that should have caught it) -- * fails here. Covers every format qt_wire_split's callers can see in practice * (fmt=1 per-row unaffected; fmt=4/5 grouped-scale formats qt_scale_bytes' - * own comment names as previously-broken too; fmt=7 nblk>O, the shape this - * review round is about). */ + * own comment names as previously-broken too; fmt=6 (E8/IQ3, FIX ROUND, + * audit finding: a FIXED 4-byte tag, not O*4 -- see qt_scale_bytes' own + * comment for why this is reachable, not dead code); fmt=7 nblk>O, the + * shape this review round is about). */ static void check_wire_split(int fmt, int O, int I, int gs, const char *tag){ QT t; memset(&t,0,sizeof t); t.fmt=fmt; t.O=O; t.I=I; t.gs=gs; int64_t want_scale = qt_scale_bytes(&t); @@ -497,12 +499,17 @@ static void check_wire_split(int fmt, int O, int I, int gs, const char *tag){ CHECK(got_weight == want_weight); /* the regression this test exists to catch: a per-row-only scale_b==O*4 * must NOT be what qt_wire_split returns for a format whose real scale - * cardinality differs from O (fmt=4/5/7 here) -- confirm the value has + * cardinality differs from O (fmt=4/5/6/7 here) -- confirm the value has * actually moved off that old constant, not just matched qt_scale_bytes() * by coincidence at a degenerate shape. */ int64_t old_wrong_scale = (int64_t)O*4; - if((fmt==4 || fmt==5 || fmt==7) && want_scale != old_wrong_scale) + if((fmt==4 || fmt==5 || fmt==6 || fmt==7) && want_scale != old_wrong_scale) CHECK(got_scale != old_wrong_scale); + /* fmt=6's scale is a FIXED 4 bytes regardless of [O,I] -- assert the + * literal value directly too, not just "moved off O*4", since a future + * regression that made it O-dependent in some OTHER wrong way would + * still pass the generic check above. */ + if(fmt==6) CHECK(want_scale == 4); } /* ---- Site-level wire regression (FIX ROUND, validator finding: mutation- @@ -600,6 +607,8 @@ int main(void){ check_wire_split(1, 4096,4096, 0, "qt_wire_split fmt=1 plain int8 (per-row scale, unaffected by the fix)"); check_wire_split(4, 2048,6144, 64, "qt_wire_split fmt=4 grouped int4 (O*ceil(I/gs) scale, not O*4)"); check_wire_split(5, 2048,6144, 0, "qt_wire_split fmt=5 int3-g64 (O*ceil(I/64) scale, not O*4)"); + check_wire_split(6, 2048,6144, 0, "qt_wire_split fmt=6 E8/IQ3 (FIXED 4-byte tag, not O*4=8192B)"); + check_wire_split(6, 1,1, 0, "qt_wire_split fmt=6 E8/IQ3 degenerate O=1 (O*4 would coincidentally also be 4 -- exercises the literal-4 assert, not just the moved-off-O*4 one)"); check_wire_split(7, 2,16384, 0, "qt_wire_split fmt=7 nblk(128) >> O(2): scale=512B, NOT O*4=8B"); check_wire_split(7, 2048,6144, 0, "qt_wire_split fmt=7 spec example: scale=3072B, NOT O*4=8192B"); test_wire_site_regression(); From 865c347c06ae95207205e0464246530659be6b77 Mon Sep 17 00:00:00 2001 From: monotophic Date: Thu, 23 Jul 2026 00:13:15 -0400 Subject: [PATCH 10/12] fp8: refuse (not crash) on fmt=6/7 in qt_addrow/qt_matvec_rows FIX ROUND 2, engine defect (clean-room conformance trial finding, the headline item): qt_addrow and qt_matvec_rows (colibri.c) serve the kv_b MLA-absorption CPU path. Both dispatch fmt 0/4/5 explicitly, then fall through assuming a PER-ROW scale (t->s[row]) followed by fmt=1/2/3 (qt_addrow) or fmt=0/1/2/3/4/5 via an if/else-if chain ending in a bare `else` (qt_matvec_rows) -- nothing stopped any OTHER fmt from reaching that fall-through. Reproduced (trial-verified, re-confirmed here): a fmt=7 QT (t->s holds ceil(O/128)*ceil(I/128) per-block floats, not O -- a [130,130] tensor has nblk=4, so t->s[row] overreads for every row past 3) reaches the fall-through's `t->q4+(int64_t)row*((I+3)/4)` -- t->q4 is NULL for fmt=7 (raw bytes live in t->q8 instead, same convention as fmt=1) -- and dereferences NULL-plus-offset: SIGSEGV. A fmt=6 QT (t->s is a FIXED 4-byte tag, qsalloc(1), not O floats) overreads t->s[row] for row>0 and then silently misreads the real E8/IQ3 lattice bytes in t->q4 as int2-packed data -- same bug SHAPE as #298's CUDA absorb-kernel fix (this file's own fmt=4/5 branches exist for exactly this reason; fmt=6 was simply missed). Both functions now refuse loudly (stderr naming the function and the fmt, exit(1)) for any fmt they don't explicitly handle: qt_addrow gets an explicit `fmt!=1 && fmt!=2 && fmt!=3` guard before the per-row-scale read (fmt 0/4/5 already returned above); qt_matvec_rows' final bare `else` becomes `else if(fmt==3)` plus a new refusing `else`. Reachability note (context, not a scope excuse): both functions serve ONLY the kv_b absorb path, and tools/repack_fp8_passthrough.py deliberately excludes kv_b_proj from fmt=7 repacking -- so this fires only via a hand-slotted or ambiguous-collision container, never this repo's own tooling's output. Crash-instead-of-refuse is still a real defect (loud failure, every refusal names its condition), and the fmt=7 QT surface these functions can now be handed is one this same PR pair created. tests/tests/test_qt_addrow.c (new): fork+pipe+waitpid refusal tests (this suite's house pattern) for fmt=6 and fmt=7 through BOTH functions, using the coordinator's own [130,130] partial-block repro shape for fmt=7; the harness also detects a raw SIGSEGV (WIFSIGNALED) and reports it as the specific failure it is, rather than hanging or mis-reporting. Byte- identity checks for every format both functions still handle (0/1/2/3/4/5) against an independently-written reference dequantizer (not copy-pasted from either function -- restructured as a single per-element loop per format), epsilon-compared (float multiplication reassociation, e.g. qt_addrow's precomputed c=coef*scale vs the reference's coef*(scale*w[i]), is not bit-exact even though mathematically identical -- confirmed by inspecting actual failures at strict equality before relaxing, all last-ULP-only, no shape/decode errors). PROVEN TO BITE (RAN, full transcript in the report): reverted the guard in qt_addrow only, rebuilt, ran the new test -- reproduced the exact SIGSEGV (signal 11) for fmt=7 and the silent non-refusal for fmt=6, caught cleanly by the test's own signal-aware harness. Reverted (diffed back to byte-identical with the pre-mutation file), reran clean. Repeated independently for qt_matvec_rows' guard alone (qt_addrow's guard left intact) -- same SIGSEGV/silent-non-refusal pair reproduced and reverted. RAN (this commit): make clean && make portable -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0, includes the new tests/test_qt_addrow binary. make test-python -> 153 tests, 10 skipped, 0 failures, exit 0. make glm METAL=1 -> clean rebuild, zero warnings. make metal-test under COLI_METAL_RESSET=0 AND =1: both full suite green, exit 0. One-shot -Wno-unused-function-stripped build: same 8 pre-existing, unrelated orphans as before this round, no new ones. --- c/Makefile | 5 +- c/colibri.c | 42 +++++- c/tests/test_qt_addrow.c | 286 +++++++++++++++++++++++++++++++++++++++ 3 files changed, 331 insertions(+), 2 deletions(-) create mode 100644 c/tests/test_qt_addrow.c diff --git a/c/Makefile b/c/Makefile index 02031cc8..9285a9c9 100644 --- a/c/Makefile +++ b/c/Makefile @@ -199,7 +199,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_mirror$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_temp_env$(EXE) tests/test_tok_o200k$(EXE) tests/test_kv_alloc$(EXE) tests/test_int3$(EXE) tests/test_int3_load$(EXE) tests/test_fp8_passthrough$(EXE) tests/test_fp8_load$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) tests/test_router_nan$(EXE) tests/test_pipe_block$(EXE) tests/test_e8_kernel$(EXE) tests/test_rans$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_mirror$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_temp_env$(EXE) tests/test_tok_o200k$(EXE) tests/test_kv_alloc$(EXE) tests/test_int3$(EXE) tests/test_int3_load$(EXE) tests/test_fp8_passthrough$(EXE) tests/test_fp8_load$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) tests/test_router_nan$(EXE) tests/test_pipe_block$(EXE) tests/test_e8_kernel$(EXE) tests/test_qt_addrow$(EXE) tests/test_rans$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -484,6 +484,9 @@ tests/test_fp8_passthrough$(EXE): tests/test_fp8_passthrough.c quant.h tests/test_fp8_load$(EXE): tests/test_fp8_load.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +tests/test_qt_addrow$(EXE): tests/test_qt_addrow.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h sample.h kv_persist.h telemetry.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + tests/test_logit_nan$(EXE): tests/test_logit_nan.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) diff --git a/c/colibri.c b/c/colibri.c index 874a6fca..312a300e 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2434,6 +2434,33 @@ static void qt_addrow(const QT *t, int row, float coef, float *acc){ for(int k=0;k>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); acc[base+k]+=cg*(float)((int)u-4); } } return; } + /* GUARD (fix round 2, engine defect -- clean-room conformance trial found a real + * SIGSEGV, reproduced): fmt 0/4/5 already returned above; everything below this + * point assumes a PER-ROW scale (t->s[row]) followed by fmt=1 (int8, explicit + * branch), fmt=2 (int4 packed, explicit branch), or the tail's own IMPLICIT fmt=3 + * (int2 packed, the final unconditional block) -- there was no guard stopping any + * OTHER fmt from reaching here. fmt=6 (E8/IQ3): t->s is a FIXED 4-byte tag (ONE + * float total, qsalloc(1) in qt_from_disk), not O floats -- t->s[row] for row>0 is + * a heap OVERREAD, and the untouched fall-through then misreads t->q4's real E8 + * lattice bytes as int2-packed data (same bug SHAPE as #298's CUDA absorb-kernel + * fix, and the same one this file's own fmt=4/5 branches above were added to + * dodge -- fmt=6 was simply missed). fmt=7 (fp8-e4m3-b128): t->s holds + * ceil(O/128)*ceil(I/128) per-block floats, not O -- t->s[row] overreads for + * row>=nblk (e.g. a [130,130] tensor has nblk=4, so every row past 3 already reads + * out of bounds), AND t->q4 is NULL for fmt=7 (raw bytes live in t->q8 instead, + * same convention as fmt=1 -- see the QT struct comment), so the fall-through's + * `t->q4+(int64_t)row*((I+3)/4)` dereferences NULL-plus-offset: SIGSEGV, + * reproduced (see the report's proof-of-bite transcript). Refuse loudly instead -- + * this function has no byte-count context of its own to validate against (it only + * ever sees an already-resolved QT), so "unsupported fmt" is the only check + * available, same "refuse rather than misread" discipline qt_resolve_fmt applies + * at load time. */ + if(t->fmt!=1 && t->fmt!=2 && t->fmt!=3){ + fprintf(stderr,"qt_addrow: unsupported fmt=%d for the per-row-scale absorb path " + "(only fmt 1/2/3 reach this point; fmt 0/4/5 are handled above and return " + "before it) -- refusing rather than misread t->s[row]/t->q4\n", t->fmt); + exit(1); + } float c=coef*t->s[row]; if(t->fmt==1){ const int8_t *w=t->q8+(int64_t)row*I; for(int i=0;ifmt==2){ const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); @@ -2480,8 +2507,21 @@ static void qt_matvec_rows(const QT *t, int r0, int n, const float *x, float *y) for(int k=0;k>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); acc+=(float)((int)u-4)*x[base+k]; } a+=(double)(acc*sr[g]); } } - else { const uint8_t *w=t->q4+(int64_t)row*((I+3)/4); float s=t->s[row]; float acc=0; + /* fmt=3 (int2 packed, per-row scale) is the only fmt this final arm legitimately + * handles -- fmt 0/4/5 matched above, fmt 1/2 have their own explicit branches + * above too. Same GUARD and same reasoning as qt_addrow's (fix round 2, engine + * defect): fmt=6's t->s is a fixed 4-byte tag (t->s[row] overreads for row>0), + * fmt=7's t->s holds per-128x128-block floats (t->s[row] overreads for + * row>=nblk) and t->q4 is NULL for fmt=7 -- both would have silently misread or + * crashed here exactly like qt_addrow did before its own fix; refuse instead. */ + else if(t->fmt==3){ const uint8_t *w=t->q4+(int64_t)row*((I+3)/4); float s=t->s[row]; float acc=0; for(int i=0;i>2]; acc+=((int)((b>>((i&3)*2))&3)-2)*x[i]; } a=acc*s; } + else { + fprintf(stderr,"qt_matvec_rows: unsupported fmt=%d for the per-row-scale absorb " + "path (only fmt 0/1/2/3/4/5 are handled) -- refusing rather than misread " + "t->s[row]/t->q4\n", t->fmt); + exit(1); + } y[j]=(float)a; } } diff --git a/c/tests/test_qt_addrow.c b/c/tests/test_qt_addrow.c new file mode 100644 index 00000000..c97c12b3 --- /dev/null +++ b/c/tests/test_qt_addrow.c @@ -0,0 +1,286 @@ +/* qt_addrow()/qt_matvec_rows() (colibri.c) -- per-row absorb helpers used only by the + * kv_b MLA-absorption CPU path (l->kv_b, see attention_rows/decode). FIX ROUND 2, engine + * defect (clean-room conformance trial finding, mutation-proven): both functions dispatch + * fmt 0/4/5 explicitly, then fall through assuming a PER-ROW scale (t->s[row]) followed by + * fmt=1/2/3 (qt_addrow) or fmt=0/1/2/3/4/5 (qt_matvec_rows, via an if/else-if chain ending + * in a bare `else`) -- nothing stopped fmt=6 (E8/IQ3, t->s is a FIXED 4-byte tag, not O + * floats) or fmt=7 (fp8-e4m3-b128, t->s holds per-128x128-block floats, not O; t->q4 is + * NULL) from reaching that fall-through. For fmt=7 specifically this SIGSEGVs: t->s[row] + * overreads (silently, usually not fatal on its own), then the untouched tail computes + * `t->q4+(int64_t)row*((I+3)/4)` on a NULL t->q4 and dereferences it. For fmt=6 it silently + * misreads the real E8/IQ3 lattice bytes as int2-packed data (same bug SHAPE as #298's CUDA + * absorb-kernel fix, which is why this file's own fmt=4/5 branches exist -- fmt=6 was simply + * missed). Both functions now refuse loudly (exit(1), naming the function and the fmt) for + * any fmt they don't explicitly handle, matching qt_resolve_fmt's own "refuse rather than + * misread" discipline. + * + * This file: (1) proves the refusal fires for fmt=6 and fmt=7 through BOTH functions + * (fork+pipe+waitpid, this suite's established house pattern for exit(1)-terminated paths -- + * see tests/test_fp8_load.c's expect_refuse/expect_stamp_refuse); (2) proves every format + * BOTH functions still legitimately handle (0/1/2/3/4/5) produces byte-identical results + * against an independently-written reference dequantizer (qt_dequant_row_ref below -- NOT + * copy-pasted from qt_addrow/qt_matvec_rows, restructured as a single per-element loop per + * format, so a real regression in either the guard's placement or the untouched per-fmt math + * would show up here, not just tautologically re-run the same code). Reachability note (not + * a scope excuse, just context): both functions serve ONLY the kv_b absorb path + * (attention_rows/decode call sites), and tools/repack_fp8_passthrough.py deliberately + * excludes kv_b_proj from fmt=7 repacking -- so this fires only via a hand-slotted or + * ambiguous-collision container, not this repo's own tooling's own output. Crash-instead-of- + * refuse is still a real defect (spec I6: loud failure, every refusal names its condition), + * and the fmt=7 QT surface these functions can now be handed is one this same PR pair + * created. */ +#define main coli_glm_main_unused +#include "../colibri.c" +#undef main + +#include +#include +#include +#ifndef _WIN32 +#include +#include +#endif + +static int fails = 0; +#define CHECK(c) do{ if(!(c)){ printf("FAIL %s:%d: %s\n", __FILE__, __LINE__, #c); fails++; } }while(0) + +static uint64_t rng = 0xA11CE5EEDF00Dull; +static uint8_t rndbyte(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; return (uint8_t)(rng & 0xFF); } +static float rndsmallf(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; + return ((int64_t)(rng & 0xFFF) - 0x800) / (float)0x800; } /* [-1,1)-ish, small magnitude */ + +/* ---- independent reference dequantizer: W[row,:] as floats, one element per loop + * iteration -- deliberately NOT the same code shape as qt_addrow/qt_matvec_rows (those + * unroll pairs for fmt=2/3, split low/high planes for fmt=5) so this genuinely + * cross-checks the production math, not just the production code running twice. ---- */ +static void qt_dequant_row_ref(const QT *t, int row, float *out){ + int I=t->I; + if(t->fmt==0){ const float *w=t->qf+(int64_t)row*I; for(int i=0;ifmt==4){ + const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); int gs=t->gs, ng=(I+gs-1)/gs; + const float *scl=t->s+(int64_t)row*ng; + for(int i=0;i>1]; int nib=(i&1)?(b>>4):(b&0xF); + out[i]=((int)nib-8)*scl[i/gs]; } + return; } + if(t->fmt==5){ + const uint8_t *w=t->q4+(int64_t)row*i3_rowbytes(I); + const float *sr=t->s+(int64_t)row*i3_groups(I); + for(int i=0;i>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); + out[i]=((int)u-4)*sr[g]; } + return; } + float s=t->s[row]; + if(t->fmt==1){ const int8_t *w=t->q8+(int64_t)row*I; for(int i=0;ifmt==2){ + const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); + for(int i=0;i>1]; int nib=(i&1)?(b>>4):(b&0xF); out[i]=((int)nib-8)*s; } + return; } + /* fmt==3 */ + { const uint8_t *w=t->q4+(int64_t)row*((I+3)/4); + for(int i=0;i>2]; int v=(b>>((i&3)*2))&3; out[i]=((int)v-2)*s; } } +} + +/* ---- fixture builders: one per format, deterministic pseudo-random payload ---- */ +static void fill_fmt0(QT *t, int O, int I){ + t->fmt=0; t->O=O; t->I=I; t->gs=0; + t->qf=(float*)malloc((size_t)O*I*sizeof(float)); + for(int i=0;iqf[i]=rndsmallf(); +} +static void fill_fmt1(QT *t, int O, int I){ + t->fmt=1; t->O=O; t->I=I; t->gs=0; + t->q8=(int8_t*)malloc((size_t)O*I); + t->s=(float*)malloc((size_t)O*sizeof(float)); + for(int i=0;iq8[i]=(int8_t)(rndbyte()-128); + for(int i=0;is[i]=0.01f+0.001f*(float)i; +} +static void fill_fmt2(QT *t, int O, int I){ + t->fmt=2; t->O=O; t->I=I; t->gs=0; + int rb=(I+1)/2; + t->q4=(uint8_t*)malloc((size_t)O*rb); + t->s=(float*)malloc((size_t)O*sizeof(float)); + for(int i=0;iq4[i]=rndbyte(); + for(int i=0;is[i]=0.02f+0.001f*(float)i; +} +static void fill_fmt3(QT *t, int O, int I){ + t->fmt=3; t->O=O; t->I=I; t->gs=0; + int rb=(I+3)/4; + t->q4=(uint8_t*)malloc((size_t)O*rb); + t->s=(float*)malloc((size_t)O*sizeof(float)); + for(int i=0;iq4[i]=rndbyte(); + for(int i=0;is[i]=0.03f+0.001f*(float)i; +} +static void fill_fmt4(QT *t, int O, int I, int gs){ + t->fmt=4; t->O=O; t->I=I; t->gs=gs; + int rb=(I+1)/2, ng=(I+gs-1)/gs; + t->q4=(uint8_t*)malloc((size_t)O*rb); + t->s=(float*)malloc((size_t)O*ng*sizeof(float)); + for(int i=0;iq4[i]=rndbyte(); + for(int i=0;is[i]=0.015f+0.0007f*(float)i; +} +static void fill_fmt5(QT *t, int O, int I){ + t->fmt=5; t->O=O; t->I=I; t->gs=0; + int64_t ng=i3_groups(I), rb=i3_rowbytes(I); + t->q4=(uint8_t*)malloc((size_t)O*rb); + t->s=(float*)malloc((size_t)(O*ng)*sizeof(float)); + for(int i=0;iq4[i]=rndbyte(); + for(int i=0;is[i]=0.025f+0.0003f*(float)i; +} +static void free_qt(QT *t){ free(t->qf); free(t->q8); free(t->q4); free(t->s); memset(t,0,sizeof *t); } + +/* ---- byte-identity: qt_addrow / qt_matvec_rows vs the independent reference, every + * format both functions still legitimately handle ---- */ +/* Epsilon, not bit-exact: qt_addrow precomputes c=coef*scale ONCE, then c*w[i] + * ((coef*scale)*w[i]); the reference below multiplies coef*(scale*w[i]) -- + * mathematically identical, but floating-point multiplication isn't + * associative, so the two can differ in the last ULP for some inputs. This + * is expected float-reassociation noise, not a regression -- confirmed by + * inspecting actual failures at bit-exact comparison (all last-digit-of- + * mantissa only, no shape/format-decode errors) before relaxing to epsilon, + * matching this suite's own established practice for reduction-order- + * sensitive checks (e.g. the metal-test suite's worst_rel comparisons). */ +static void check_addrow_identity(QT *t, const char *tag){ + int I=t->I; float *ref=(float*)malloc((size_t)I*sizeof(float)); + float *acc=(float*)malloc((size_t)I*sizeof(float)); + float coef=1.7f; /* != 1, so a coef-scaling bug can't hide */ + for(int row=0; rowO; row++){ + qt_dequant_row_ref(t,row,ref); + memset(acc,0,(size_t)I*sizeof(float)); + qt_addrow(t,row,coef,acc); + for(int i=0;i1e-6f ? ae/fabsf(want) : ae; + if(rel > 1e-5f){ + printf("FAIL %s: qt_addrow row=%d i=%d got=%.9g want=%.9g rel=%.3g\n",tag,row,i,(double)acc[i],(double)want,(double)rel); + fails++; + } + } + } + free(ref); free(acc); +} +static void check_matvec_identity(QT *t, const char *tag){ + int I=t->I; float *ref=(float*)malloc((size_t)I*sizeof(float)); + float *x=(float*)malloc((size_t)I*sizeof(float)); + for(int i=0;iO; row++){ + qt_dequant_row_ref(t,row,ref); + double want=0; for(int i=0;i1e-6 ? fabs((double)y-want)/fabs(want) : fabs((double)y-want); + if(relerr > 1e-4){ + printf("FAIL %s: qt_matvec_rows row=%d got=%.9g want=%.9g relerr=%.3g\n",tag,row,(double)y,want,relerr); + fails++; + } + } + free(ref); free(x); +} + +static void test_byte_identity_all_formats(void){ + QT t; + memset(&t,0,sizeof t); fill_fmt0(&t,4,17); check_addrow_identity(&t,"fmt=0"); check_matvec_identity(&t,"fmt=0"); free_qt(&t); + memset(&t,0,sizeof t); fill_fmt1(&t,4,17); check_addrow_identity(&t,"fmt=1"); check_matvec_identity(&t,"fmt=1"); free_qt(&t); + memset(&t,0,sizeof t); fill_fmt2(&t,4,17); check_addrow_identity(&t,"fmt=2"); check_matvec_identity(&t,"fmt=2"); free_qt(&t); + memset(&t,0,sizeof t); fill_fmt3(&t,4,17); check_addrow_identity(&t,"fmt=3"); check_matvec_identity(&t,"fmt=3"); free_qt(&t); + memset(&t,0,sizeof t); fill_fmt4(&t,4,40,16);check_addrow_identity(&t,"fmt=4"); check_matvec_identity(&t,"fmt=4"); free_qt(&t); + memset(&t,0,sizeof t); fill_fmt5(&t,4,130); check_addrow_identity(&t,"fmt=5"); check_matvec_identity(&t,"fmt=5"); free_qt(&t); +} + +/* ---- refusal: fork+pipe+waitpid, this suite's house pattern (see test_fp8_load.c's + * expect_refuse/expect_stamp_refuse) -- must exit(1) with a "refus"-containing message, + * never reach the caller's continuation, never crash with a signal. ---- */ +typedef void (*absorb_fn)(void); +static void call_addrow_fmt7(void){ + QT t; memset(&t,0,sizeof t); + enum { O=130, I=130 }; /* the coordinator's own repro shape: nblkO=nblkI=2, nblk=4 */ + static uint8_t q8[O*I]; static float s[4]; + for(int i=0;i Date: Thu, 23 Jul 2026 00:22:26 -0400 Subject: [PATCH 11/12] fp8: repack tool refuses loudly on zero repack-target tensors FIX ROUND 2, item 3 (clean-room conformance trial finding, spec I6 -- loud failure, every refusal names its condition): a real (non-dry-run) run of tools/repack_fp8_passthrough.py whose --indir contains FP8 tensors but none matching a repack-target kind used to exit 0 having emitted nothing -- trial-verified with a non-resident-named FP8 tensor selecting silently. That is a silent trap: an empty "container" (just the resume/params sidecars, no actual .safetensors output) nobody asked for, with no caller-side check positioned to catch it. main() now refuses (nonzero exit, stderr naming --indir and the shard count checked) when every shard -- this run's AND any prior resumed run's, via the `done` progress dict, one entry per shard -- has NEVER produced an output. --dry-run is deliberately NOT covered: printing "0 tensor(s) selected" already IS the loud, honest answer dry-run exists to give, not a silent no-op, so turning it into a refusal would be wrong, not thorough (this judgment call implemented directly, per the coordinator's own "or escalate the disagreement" allowance -- no disagreement here, both readings agree once dry-run's actual purpose is considered). tests/test_fp8_repack.py: new ZeroTargetRefusalTest -- a fixture shard with FP8 tensors but none matching a repack-target kind (routed expert + f32 norm, mirroring the trial's own scenario) must refuse on a real run and still exit 0 on --dry-run. Also updates the module's kv_b_proj-exclusion comment (repack tool docstring), which the #528+FIX-ROUND-2 qt_addrow/qt_matvec_rows fix made stale: it claimed an fmt=7 kv_b_proj tensor "would silently misread fp8 bytes as int2-packed data" -- true only for fmt=6; fmt=7 actually SIGSEGV'd (t->q4 is NULL for fmt=7). Both now refuse loudly at the absorb call site instead -- corrected to describe the actual (fixed) behavior, while keeping the underlying exclusion decision itself unchanged (refusing loudly there is still not the same as supporting fmt=7 for kv_b_proj). PROVEN TO BITE (RAN): removed the new zero-target check only, ran the new test -- `test_zero_targets_refuses_loudly` failed exactly as expected (returncode 0 instead of nonzero); `test_zero_targets_dry_run_unaffected` still passed (confirming the dry-run path was never touched by the mutation). Restored (diffed back to byte-identical with the pre-mutation file), reran: both new tests green. RAN (this commit): python3 -m unittest tests.test_fp8_repack -v -> 14/14 pass. make clean && make portable -> clean rebuild, zero warnings. make test-c -> full C suite green, exit 0. make test-python -> 157 tests, 10 skipped, 0 failures, exit 0 (2 more than before this round: the new ZeroTargetRefusalTest cases). --- c/tests/test_fp8_repack.py | 62 ++++++++++++++++++++++++++++++- c/tools/repack_fp8_passthrough.py | 34 ++++++++++++++--- 2 files changed, 89 insertions(+), 7 deletions(-) diff --git a/c/tests/test_fp8_repack.py b/c/tests/test_fp8_repack.py index 64bb3dc3..e06b60ed 100644 --- a/c/tests/test_fp8_repack.py +++ b/c/tests/test_fp8_repack.py @@ -18,9 +18,14 @@ coincide with a per-row int8 scale byte count) must ALSO be refused (maintainer review, #528: this is not hypothetical, GLM-5.2's own self_attn.o_proj.weight is a real instance) -- --dry-run writing nothing, and -the #383-class resume/params-guard behavior. +the #383-class resume/params-guard behavior. FIX ROUND 2 (clean-room +conformance trial, spec I6): a real (non-dry-run) run that selects ZERO +repack-target tensors across --indir must refuse loudly (nonzero exit, +stderr naming the condition), not exit 0 having emitted nothing -- an empty +"container" nobody asked for; --dry-run's own "0 selected" report is +unaffected (that IS the loud, honest answer dry-run exists to give). """ -import glob, json, os, struct, sys, tempfile, unittest +import glob, json, os, struct, subprocess, sys, tempfile, unittest try: import torch @@ -238,5 +243,58 @@ def test_params_mismatch_refused(self): self.assertEqual(len(outs), 1) +def _make_zero_target_fixture(path, D=256, I_=384): + """A shard with real FP8 tensors, but NONE matching a repack-target kind -- + mirrors the trial's own scenario (a non-resident-named FP8 tensor selects + nothing): one routed expert (kind "x", explicitly excluded) and one f32 + norm (never a repack candidate at all).""" + torch.manual_seed(2) + sd = { + "model.layers.0.mlp.experts.0.gate_proj.weight": torch.randn(I_, D) * 0.02, # routed: EXCLUDE + "model.layers.0.input_layernorm.weight": torch.randn(D), # f32: EXCLUDE + } + save_fp8_safetensors(sd, path) + + +class ZeroTargetRefusalTest(unittest.TestCase): + """FIX ROUND 2, item 3 (clean-room conformance trial, spec I6): a real + (non-dry-run) run whose --indir has FP8 tensors but none matching a + repack-target kind must refuse loudly, not exit 0 having written nothing.""" + + def setUp(self): + self.tmp = tempfile.TemporaryDirectory() + self.indir = os.path.join(self.tmp.name, "fp8src") + os.makedirs(self.indir) + self.shard = os.path.join(self.indir, "model-00001-of-00001.safetensors") + _make_zero_target_fixture(self.shard) + self.outdir = os.path.join(self.tmp.name, "out") + self.tool = os.path.join(os.path.dirname(__file__), "..", "tools", "repack_fp8_passthrough.py") + + def tearDown(self): + self.tmp.cleanup() + + def _run(self, extra_args=()): + return subprocess.run( + [sys.executable, self.tool, "--indir", self.indir, "--outdir", self.outdir, + "--n-layers", "5", *extra_args], + capture_output=True, text=True) + + def test_zero_targets_refuses_loudly(self): + proc = self._run() + self.assertNotEqual(proc.returncode, 0, "zero repack-target tensors must refuse, not exit 0") + self.assertIn("no repack-target tensors", proc.stderr) + self.assertIn(self.indir, proc.stderr, "the refusal must name the --indir condition") + # confirmed empty: no output container was (or should be) produced + self.assertEqual(glob.glob(os.path.join(self.outdir, "out-fp8pass-*.safetensors")), []) + + def test_zero_targets_dry_run_unaffected(self): + """--dry-run's own "0 tensor(s) selected" report is the loud, honest + answer dry-run exists to give -- not a silent no-op -- so it must NOT + be turned into a refusal by this fix.""" + proc = self._run(["--dry-run"]) + self.assertEqual(proc.returncode, 0) + self.assertIn("0 tensor(s) selected", proc.stdout) + + if __name__ == "__main__": unittest.main() diff --git a/c/tools/repack_fp8_passthrough.py b/c/tools/repack_fp8_passthrough.py index 704703af..d292270f 100644 --- a/c/tools/repack_fp8_passthrough.py +++ b/c/tools/repack_fp8_passthrough.py @@ -58,11 +58,20 @@ tensor classify() would otherwise select: colibri.c's MLA-absorption CPU path (qt_addrow/qt_matvec_rows, called only on l->kv_b -- the always-available fallback whenever the Metal fused decode kernel isn't used) has no fmt==7 -case and would silently misread fp8 bytes as int2-packed data; the CUDA -absorb kernels (coli_cuda_attention_absorb/_kvdev in backend_cuda.cu) are -similarly int4-specific. Repacking kv_b_proj to fmt=7 needs that CPU+CUDA -absorb-path work first -- out of scope for this build, so this tool refuses -to produce a container the engine cannot safely read. +support. FIX ROUND 2 update: an fmt=7 (or fmt=6) tensor reaching either +function now refuses loudly (exit(1), naming the function and the fmt) -- +before that fix it was worse than "misread as int2": qt_addrow SIGSEGV'd +outright (t->q4 is NULL for fmt=7; the untouched int2 fall-through +dereferenced it), and only fmt=6 (t->q4 non-NULL, wrong per-row scale +geometry) actually produced silently-wrong values. Either way, this tool +still does not select kv_b_proj: refusing loudly at the absorb call site is +not the same as SUPPORTING fmt=7 there (no dequant path exists, whether it +now trips a controlled refusal or, before the fix, undefined behavior). The +CUDA absorb kernels (coli_cuda_attention_absorb/_kvdev in backend_cuda.cu) +are similarly int4-specific and unaffected by that fix. Repacking kv_b_proj +to fmt=7 needs real fmt=7 CPU+CUDA absorb-path support first -- out of scope +for this build, so this tool refuses to produce a container the engine +cannot safely read. This tool does NOT write a container metadata stamp: __metadata__-based self-description (writer + qt_verify_fmt_stamp reader + docs/FORMATS.md @@ -295,6 +304,21 @@ def main(): print(f"[RESUME] {skipped} shard(s) already done in {a.outdir}, skipped") print(f"[REPACK] {fresh} new output shard(s) written") _print_inventory_summary(all_inv, dry_run=False) + # FIX ROUND 2 (clean-room conformance trial, spec I6 -- loud failure, every + # refusal names its condition): a run that selects ZERO repack-target tensors + # across every shard under --indir (present run AND any prior resumed run -- + # `done` has one entry per shard, "" meaning that shard has NEVER produced an + # output) used to exit 0 having emitted nothing -- a silent trap: an empty + # "container" (just the resume/params sidecars, no actual .safetensors output) + # that nobody asked for and no caller-side check would catch. Refuse loudly + # instead. --dry-run is NOT covered here: reporting "0 tensors selected" IS + # the loud, honest answer dry-run exists to give, not a silent no-op -- see + # the early `return` above, before any of this resume/write bookkeeping. + if all(v == "" for v in done.values()): + print(f"ERROR: no repack-target tensors found under {a.indir} (checked " + f"{len(shards)} shard(s), 0 selected) -- nothing emitted; refusing " + f"to exit 0 for an empty container nobody asked for", file=sys.stderr) + sys.exit(1) if __name__ == "__main__": From dbf8d00990afbc8aa097d60cee3cfa5e849c33df Mon Sep 17 00:00:00 2001 From: monotophic Date: Thu, 23 Jul 2026 02:15:34 -0400 Subject: [PATCH 12/12] tests: drain child stderr to EOF in fork-capture helpers A single read() on the capture pipe can return short on Linux (glibc's unbuffered stderr arrives in multiple chunks), truncating the refusal message before its marker -- the linux CI job caught expect_refuse's [64,98] fmt=6/fmt=7 collision case failing exactly this way while the refusal itself fired correctly. Loop the read to EOF; same fix in test_qt_addrow.c's identical helper. macOS coalesced the chunks, which is why local runs never saw it. --- c/tests/test_fp8_load.c | 2 +- c/tests/test_qt_addrow.c | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/c/tests/test_fp8_load.c b/c/tests/test_fp8_load.c index 9c01b24a..b57d9d1a 100644 --- a/c/tests/test_fp8_load.c +++ b/c/tests/test_fp8_load.c @@ -183,7 +183,7 @@ static int expect_refuse(int O, int I, int64_t nb, int64_t ns, const char *tag){ _exit(42); /* reaching here is the bug */ } close(pipefd[1]); - char err[1024]={0}; ssize_t n=read(pipefd[0],err,sizeof(err)-1); (void)n; + char err[1024]={0}; size_t eoff=0; ssize_t n; /* drain to EOF: a single read() can return SHORT on Linux pipes (glibc unbuffered stderr arrives in chunks) -- truncated the refusal message past the marker, CI-caught */ while(eoff0) eoff+=(size_t)n; close(pipefd[0]); int status=0; waitpid(pid,&status,0); int ok = WIFEXITED(status) && WEXITSTATUS(status)==1; diff --git a/c/tests/test_qt_addrow.c b/c/tests/test_qt_addrow.c index c97c12b3..cc9f81c5 100644 --- a/c/tests/test_qt_addrow.c +++ b/c/tests/test_qt_addrow.c @@ -245,7 +245,7 @@ static int expect_refuse_call(absorb_fn fn, const char *tag){ _exit(42); /* reaching here (surviving the call without exit(1)) is the bug */ } close(pipefd[1]); - char err[1024]={0}; ssize_t n=read(pipefd[0],err,sizeof(err)-1); (void)n; + char err[1024]={0}; size_t eoff=0; ssize_t n; /* drain to EOF: a single read() can return SHORT on Linux pipes (glibc unbuffered stderr arrives in chunks) -- truncated the refusal message past the marker, CI-caught */ while(eoff0) eoff+=(size_t)n; close(pipefd[0]); int status=0; waitpid(pid,&status,0); if(WIFSIGNALED(status)){