diff --git a/c/Makefile b/c/Makefile index 9247ab81..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_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 @@ -478,6 +478,15 @@ 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_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/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/colibri.c b/c/colibri.c index 7783036e..312a300e 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,91 @@ 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. 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), 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 + * 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. + * + * 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: 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 + * 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; @@ -300,6 +433,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 +713,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 +756,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 +1178,162 @@ 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. + * + * 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 + * 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); + /* 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 " + "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 +1363,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); @@ -2134,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); @@ -2180,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; } } @@ -2389,10 +2729,25 @@ 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): 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){ + && 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 +4582,24 @@ 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. 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 && 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; @@ -6090,8 +6457,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; @@ -6104,8 +6470,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/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 -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"); 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() diff --git a/c/tests/test_fp8_load.c b/c/tests/test_fp8_load.c new file mode 100644 index 00000000..b57d9d1a --- /dev/null +++ b/c/tests/test_fp8_load.c @@ -0,0 +1,618 @@ +/* fmt=7 (native FP8-e4m3 passthrough) loader-seam tests. + * + * 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 and merged it into dev as a + * REAL fmt=6 (see quant.h's E8 constants and e8_ helper functions, and + * qt_resolve_fmt's ns==4-tag early check); 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. + * + * Part A: qt_resolve_fmt disambiguation suite -- THE DESIGN LANDMINE. fmt=7 + * weight bytes are byte-identical to fmt=1 (int8): both are O*I raw bytes. + * 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 -- 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. 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. 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 + * int8 control tensor of a DIFFERENT, non-colliding shape, loads both through + * qt_from_disk, and checks the byte-count/.qs-size inference picks fmt=7 vs + * 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, 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). 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 + * 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. */ +/* 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 +#include +#ifndef _WIN32 +#include +#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) + +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}; 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; + 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")); + + /* --- 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. 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_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_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. --- */ + 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;iO*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=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); + 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/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==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- + * 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(); + test_ue8m0_scale_refusal(); + test_loader_seam(); + check_fp8_bytes(2048,6144, "qt_bytes fmt=7 gate/up-shaped O=2048 I=6144 (spec example)"); + 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(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(); + if(fails){ printf("fp8 loader-seam tests: %d FAILED\n", fails); return 1; } + printf("fp8 loader-seam tests: ok\n"); + return 0; +} diff --git a/c/tests/test_fp8_passthrough.c b/c/tests/test_fp8_passthrough.c new file mode 100644 index 00000000..61fb3709 --- /dev/null +++ b/c/tests/test_fp8_passthrough.c @@ -0,0 +1,165 @@ +/* fmt=7 (native FP8-e4m3 passthrough) CPU kernel tests: LUT exactness (all 256 + * byte codes, incl. +-0, subnormals, NaN), matmul_fp8 vs a double-precision + * block-scale reference on block-edge (non-128-multiple O/I) shapes, a + * non-square-block-grid case with a distinct scale per block (catches a + * transposed/swapped block-index stride the same way test_backend_metal.mm's + * "stride-audit" case does for the GPU kernel), and the NaN-propagation policy + * (decode -> 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 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 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. 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, subprocess, 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) + + 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): + 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) + + +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/tests/test_qt_addrow.c b/c/tests/test_qt_addrow.c new file mode 100644 index 00000000..cc9f81c5 --- /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;i0) eoff+=(size_t)n; + close(pipefd[0]); + int status=0; waitpid(pid,&status,0); + if(WIFSIGNALED(status)){ + printf("FAIL %s: crashed with signal %d instead of refusing (exit(1)) -- this IS the " + "engine defect this test exists to catch\n", tag, WTERMSIG(status)); + return 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 + printf("skipped on Windows (no fork): %s\n", tag); + (void)fn; + return 1; +#endif +} + +static void test_refusals(void){ + CHECK(expect_refuse_call(call_addrow_fmt7, "qt_addrow refuses fmt=7 (was SIGSEGV)")); + CHECK(expect_refuse_call(call_addrow_fmt6, "qt_addrow refuses fmt=6")); + CHECK(expect_refuse_call(call_matvec_fmt7, "qt_matvec_rows refuses fmt=7")); + CHECK(expect_refuse_call(call_matvec_fmt6, "qt_matvec_rows refuses fmt=6")); +} + +int main(void){ + test_byte_identity_all_formats(); + test_refusals(); + if(fails){ printf("qt_addrow/qt_matvec_rows tests: %d FAILED\n", fails); return 1; } + printf("qt_addrow/qt_matvec_rows tests: ok\n"); + return 0; +} 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/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 new file mode 100644 index 00000000..d292270f --- /dev/null +++ b/c/tools/repack_fp8_passthrough.py @@ -0,0 +1,325 @@ +"""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 -- 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 +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 +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 +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. + + 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): + """--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) + # 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__": + main()