diff --git a/c/Makefile b/c/Makefile index 3ebf7d8..06fe176 100644 --- a/c/Makefile +++ b/c/Makefile @@ -171,7 +171,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(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_tok_o200k$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) tests/test_pipe_block$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(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_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_pipe_block$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -343,6 +343,12 @@ tests/test_sample_nan$(EXE): tests/test_sample_nan.c colibri.c st.h uring.h json tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c colibri.c st.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_int3$(EXE): tests/test_int3.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_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_logit_nan$(EXE): tests/test_logit_nan.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) diff --git a/c/colibri.c b/c/colibri.c index fe1b383..ef32b01 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -94,7 +94,13 @@ 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). q4 ospita sia int4 che int2 packed. */ +/* fmt: 0 F32, 1 INT8, 2 INT4 (2/byte), 3 INT2 (4/byte), 4 INT4-GROUPED, 5 INT3-G64. + * 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=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. */ 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 @@ -110,6 +116,10 @@ static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ if(t->fmt==4){ /* int4 grouped: packed nibbles + O*ceil(I/gs) scales */ int ng=(t->I+t->gs-1)/t->gs; return (int64_t)t->O*((t->I+1)/2) + (int64_t)t->O*ng*4; } + if(t->fmt==5){ /* int3-g64: 24B/group weights + one f32 scale per group (I3_* in quant.h, + * included below — keep the arithmetic literal here) */ + int64_t ng=((int64_t)t->I+63)/64; + return (int64_t)t->O*ng*24 + (int64_t)t->O*ng*4; } return (int64_t)t->O*((t->I+1)/2) + (int64_t)t->O*4; /* fmt=2 int4 per-row */ } @@ -267,6 +277,7 @@ static void qt_cuda_reset(QT *t){ t->cuda_failed=0; } static int qt_cuda_upload(QT *t){ + if(t->fmt==5) return 0; /* int3-g64: no CUDA kernel yet — tensor stays CPU-side */ 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 @@ -445,7 +456,7 @@ static void matmul_qt_ex(float *y, const float *x, QT *w, int S, int allow_idot) } #endif #ifdef COLI_CUDA - if(g_cuda_enabled && w->cuda_eligible && !w->cuda_failed && !omp_in_parallel()){ + if(g_cuda_enabled && w->cuda_eligible && !w->cuda_failed && w->fmt!=5 && !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)) return; @@ -467,6 +478,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 matmul_i4(y,x,w->q4,w->s,S,w->I,w->O); } @@ -659,18 +671,21 @@ static _Atomic int g_cur_moe_layer=-1; /* massimo layer moe in cui il MAIN e' static int g_pilot_inflight[256]; /* protected by g_pilot_mx; URING can load a layer concurrently */ static _Atomic long g_pilot_loads=0; /* load cross-layer VERI completati (banda spesa) */ static _Atomic long g_pilot_drops=0; /* predizioni scartate perche' il main possiede gia' il layer */ -/* sceglie il formato da `bits`: >=16 f32, 5..8 int8, <=4 int4-packed */ +/* format from `bits`: >=16 f32, 5..8 int8, 4 int4-packed, 3 int3-g64 (group scales), <=2 int2 */ static void qt_alloc(QT *t, int O, int I, int bits){ t->O=O; t->I=I; t->qf=NULL; t->q8=NULL; t->q4=NULL; t->s=NULL; if(bits>=16){ t->fmt=0; t->qf=falloc((int64_t)O*I); } else if(bits>=5 || g_nopack){ t->fmt=1; t->q8=qalloc((int64_t)O*I); t->s=qsalloc(O); } - else if(bits>=3){ t->fmt=2; t->q4=qalloc((int64_t)O*((I+1)/2)); t->s=qsalloc(O); } + else if(bits>=4){ t->fmt=2; t->q4=qalloc((int64_t)O*((I+1)/2)); t->s=qsalloc(O); } + else if(bits==3){ t->fmt=5; t->q4=qalloc((int64_t)O*i3_rowbytes(I)); + t->s=(float*)qalloc((size_t)O*i3_groups(I)*sizeof(float)); } else { t->fmt=3; t->q4=qalloc((int64_t)O*((I+3)/4)); t->s=qsalloc(O); } } static void qt_fill(QT *t, const float *w, int bits){ if(t->fmt==0) memcpy(t->qf, w, (int64_t)t->O*t->I*sizeof(float)); else if(t->fmt==1) quantize_rows(w, t->q8, t->s, t->O, t->I, bits); else if(t->fmt==3) pack_int2(w, t->q4, t->s, t->O, t->I, bits); + else if(t->fmt==5) pack_int3_g64(w, t->q4, t->s, t->O, t->I); else pack_int4(w, t->q4, t->s, t->O, t->I, bits); } @@ -840,16 +855,22 @@ 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) e scrive *gs. */ + * si termina invece di sforare. Ritorna fmt (1/2/3/4/5) e scrive *gs. */ 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); - int fmt = (nb==exp_i8)?1 : (nb==exp_i4)?2 : (nb==exp_i2)?3 : 0; + int64_t exp_i3=(int64_t)O*i3_rowbytes(I); /* int3-g64 (fmt=5): 24B per 64-input group */ + /* 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 layout for [%d,%d], refusing (untrusted container)\n", + fprintf(stderr,"%s: quantized weight is %lld bytes — no int8/int4/int2/int3-g64 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; } } - int64_t exp_scale = (fmt==4)? (int64_t)O*((I+*gs-1)/(*gs)) : (int64_t)O; /* in FLOAT */ + 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 */ 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); } @@ -873,6 +894,9 @@ static void qt_from_disk(Model *m, const char *name, int O, int I, int bits, int else if(fmt==4){ int ng=(I+gs-1)/gs; if(t->fmt!=4||!t->q4){ t->fmt=4; t->O=O; t->I=I; t->gs=gs; t->q4=qalloc(nb); t->s=falloc((int64_t)O*ng); } st_read_raw(&m->S,name,t->q4,drop); } + else if(fmt==5){ int64_t ng=i3_groups(I); /* int3-g64: 24B/group weights + O*ng group scales */ + if(t->fmt!=5||!t->q4){ t->fmt=5; t->O=O; t->I=I; t->gs=0; t->q4=qalloc(nb); t->s=falloc((int64_t)O*ng); } + st_read_raw(&m->S,name,t->q4,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); } st_read_f32(&m->S,sn,t->s,drop); } else { @@ -1104,6 +1128,13 @@ static void embed_row(Model *m, int tok, float *x){ for(int i=0;i>1]; x[i]=(float)((int)(byte&0xF)-8)*s; if(i+1>4)-8)*s; } return; } + if(e->fmt==5){ const uint8_t *q=e->q4+(int64_t)tok*i3_rowbytes(D); /* int3-g64 */ + const float *sr=e->s+(int64_t)tok*i3_groups(D); int64_t ng=i3_groups(D); + for(int64_t g=0; g>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); + x[base+k]=(float)((int)u-4)*sr[g]; } } + return; } const uint8_t *q=e->q4+(int64_t)tok*((D+3)/4); float s=e->s[tok]; /* int2 */ for(int i=0;i>2]; int sh=(i&3)*2; x[i]=(float)((int)((byte>>sh)&3)-2)*s; } } @@ -1598,8 +1629,11 @@ static int uring_finalize_load(UringBatch *b,int li,int publish_eid){ for(int k=0;k<3;k++){ fp[k]=s->fslab+fo; fo+=l->tq[k]->nbytes/4; int64_t nb=l->tw[k]->nbytes; - int fmt=(nb==(int64_t)OO[k]*II[k])?1:(nb==(int64_t)OO[k]*((II[k]+1)/2))?2:3; - qt[k]->fmt=fmt; qt[k]->O=OO[k]; qt[k]->I=II[k]; qt[k]->qf=NULL; + /* qt_resolve_fmt like the other two expert paths: the raw ?1:?2:3 inference here + * missed grouped int4 (fmt=4, gs never set) and would mis-tag int3-g64 as int2. */ + int gs=0; + int fmt=qt_resolve_fmt(l->tw[k]->name,OO[k],II[k],nb,l->tq[k]->nbytes,&gs); + qt[k]->fmt=fmt; qt[k]->O=OO[k]; qt[k]->I=II[k]; qt[k]->gs=gs; qt[k]->qf=NULL; qt[k]->q8=(int8_t*)(s->slab+l->pos[k]); qt[k]->q4=s->slab+l->pos[k]; qt[k]->s=fp[k]; } if(publish_eid) s->eid=l->eid; @@ -1830,6 +1864,14 @@ static void qt_addrow(const QT *t, int row, float coef, float *acc){ acc[i] +=coef*scl[i/gs] *((int)(b&0xF)-8); acc[i+1]+=coef*scl[(i+1)/gs]*((int)(b>>4)-8); } if(I&1){ uint8_t b=w[I>>1]; acc[I-1]+=coef*scl[(I-1)/gs]*((int)(b&0xF)-8); } return; } + /* fmt=5 likewise before c: int3-g64 scales are per-GROUP [O,ng], not s[row] */ + 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); int64_t ng=i3_groups(I); + for(int64_t g=0; g>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); + acc[base+k]+=cg*(float)((int)u-4); } } + return; } 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); @@ -1855,6 +1897,13 @@ static void qt_matvec_rows(const QT *t, int r0, int n, const float *x, float *y) for(int i=base;i>1]; acc+=(float)((i&1)?((int)(b>>4)-8):((int)(b&0xF)-8))*x[i]; } a+=(double)acc*scl[g]; } } + else 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); int64_t ng=i3_groups(I); + for(int64_t g=0; g>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; for(int i=0;i>2]; acc+=((int)((b>>((i&3)*2))&3)-2)*x[i]; } a=acc*s; } y[j]=(float)a; diff --git a/c/quant.h b/c/quant.h index eb7dba2..b375538 100644 --- a/c/quant.h +++ b/c/quant.h @@ -268,6 +268,68 @@ static void matmul_i2(float *y, const float *x, const uint8_t *q2, const float * y[(int64_t)s*O+o]=a*sc; } } } +/* ---- int3-g64 (fmt=5): 3-bit weights with ONE f32 scale per 64-input group - + * Per group: 16B low plane (2 bits/val, int2 layout) + 8B high plane (1 bit/val), + * values in [-4,3] stored v+4. 3.5 bits/weight effective — the quality/size point + * the #132 OLMoE ablation measured BEATING per-row int4. */ +#define I3_GROUP 64 +#define I3_GBYTES 24 /* 16B low plane + 8B high plane per group */ +static inline int64_t i3_groups(int I){ return ((int64_t)I + I3_GROUP - 1) / I3_GROUP; } +static inline int64_t i3_rowbytes(int I){ return i3_groups(I) * I3_GBYTES; } + +/* Dequant-on-use with PER-GROUP scale. Exact f32 path only (no IDOT in v1: int8 + * activations don't compose with per-group accumulation without a kernel + * restructure — follow-up). NEON: low plane = matmul_i2's unpack, high plane + * expanded via vtst on bit masks; x86 stays scalar for now (follow-up). */ +static void matmul_i3(float *y, const float *x, const uint8_t *q3, const float *scale, int S, int I, int O){ + int64_t ng=i3_groups(I), rb=i3_rowbytes(I); + #pragma omp parallel for schedule(static) + for(int o=0;o>2), 4); /* 4 bytes = 16 low-plane values */ + uint8x8_t by=vreinterpret_u8_u32(vdup_n_u32(wd)); + uint8x8x2_t z01=vzip_u8(vand_u8(by,m2v), vand_u8(vshr_n_u8(by,2),m2v)); + uint8x8x2_t z23=vzip_u8(vand_u8(vshr_n_u8(by,4),m2v), vshr_n_u8(by,6)); + uint16x4x2_t zz=vzip_u16(vreinterpret_u16_u8(z01.val[0]), vreinterpret_u16_u8(z23.val[0])); + uint8x16_t lov=vcombine_u8(vreinterpret_u8_u16(zz.val[0]), vreinterpret_u8_u16(zz.val[1])); + uint8x16_t hv=vcombine_u8(vdup_n_u8(hi[k>>3]), vdup_n_u8(hi[(k>>3)+1])); + uint8x16_t hb=vandq_u8(vtstq_u8(hv,bitm), fourq); /* 4 where high bit set */ + int8x16_t wq=vsubq_s8(vreinterpretq_s8_u8(vaddq_u8(lov,hb)), b4q); /* [-4,3] in order */ + int16x8_t w0=vmovl_s8(vget_low_s8(wq)), w1=vmovl_s8(vget_high_s8(wq)); + ac0=vfmaq_f32(ac0, vld1q_f32(xs+base+k), vcvtq_f32_s32(vmovl_s16(vget_low_s16(w0)))); + ac1=vfmaq_f32(ac1, vld1q_f32(xs+base+k+4), vcvtq_f32_s32(vmovl_s16(vget_high_s16(w0)))); + ac0=vfmaq_f32(ac0, vld1q_f32(xs+base+k+8), vcvtq_f32_s32(vmovl_s16(vget_low_s16(w1)))); + ac1=vfmaq_f32(ac1, vld1q_f32(xs+base+k+12), vcvtq_f32_s32(vmovl_s16(vget_high_s16(w1)))); + } + a=vaddvq_f32(vaddq_f32(ac0,ac1)); + } +#endif + for(;k>2]>>((k&3)*2))&3) | (((hi[k>>3]>>(k&7))&1)<<2); + a += xs[base+k]*(float)((int)u-4); + } + acc += a*srow[g]; + } + y[(int64_t)s*O+o]=acc; + } + } +} + /* ---- IDOT: integer dot kernels (int8-quantized activations) --------------- */ #if defined(__AVX512VNNI__) && defined(__AVX512BW__) #define IDOT_KERNEL "avx512-vnni" @@ -689,6 +751,34 @@ static void pack_int4(const float *w, uint8_t *q4, float *scale, int O, int I, i } } } +/* quantize w[O,I] f32 -> int3-g64 (fmt=5): per 64-input group, symmetric absmax + * (qmax=3, clamp [-4,3], stored v+4), 16B low plane + 8B high plane, ONE f32 scale + * per group. Same math as tools/quant_ablation.py `_quant_last_dim(bits=3, group=64)` + * (#132), here with real bit packing. */ +static void pack_int3_g64(const float *w, uint8_t *q3, float *scale, int O, int I){ + int64_t ng=i3_groups(I), rb=i3_rowbytes(I); + #pragma omp parallel for schedule(static) + for(int o=0;oamax)amax=a; } + float s=amax/3.f; if(s<1e-8f)s=1e-8f; sr[g]=s; + uint8_t *lo=qr+g*I3_GBYTES, *hi=lo+16; + memset(lo,0,I3_GBYTES); + for(int k=0;k3)v=3; if(v<-4)v=-4; + unsigned u=(unsigned)(v+4); /* 0..7 */ + lo[k>>2] |= (uint8_t)((u&3)<<((k&3)*2)); + hi[k>>3] |= (uint8_t)(((u>>2)&1)<<(k&7)); + } + } + } +} + static void pack_int2(const float *w, uint8_t *q2, float *scale, int O, int I, int bits){ int qmax=(1<<(bits-1))-1, rb=(I+3)/4; #pragma omp parallel for schedule(static) diff --git a/c/tests/test_int3.c b/c/tests/test_int3.c new file mode 100644 index 0000000..9fa7c73 --- /dev/null +++ b/c/tests/test_int3.c @@ -0,0 +1,147 @@ +/* int3-g64 (fmt=5) tests: pack layout, dequant round-trip vs plain-C reference, + * matmul_i3 (NEON + scalar tail) vs reference dequant-matmul, per-row helpers, + * the .qs-size format tag, and the quality claim in miniature (per-group int3 + * beats per-row int4 on rows with outliers — the #132 result this format ships). */ +#define main coli_glm_main_unused +#include "../colibri.c" +#undef main + +#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 = 0x9E3779B97F4A7C15ull; +static float rndf(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; + return ((int64_t)(rng & 0xFFFFF) - 0x80000) / (float)0x80000; } + +/* reference: quantize like pack_int3_g64 but keep dequantized f32 (mirrors + * quant_ablation._quant_last_dim(bits=3, group=64)) */ +static void ref_i3_dequant(const float *w, float *dq, int O, int I){ + int64_t ng=i3_groups(I); + for(int o=0;oamax)amax=a; } + float s=amax/3.f; if(s<1e-8f)s=1e-8f; + for(int k=0;k3)v=3; if(v<-4)v=-4; + dq[(int64_t)o*I+base+k]=(float)v*s; + } + } +} +static void unpack_i3(const uint8_t *q3, const float *s, float *dq, int O, int I){ + int64_t ng=i3_groups(I), rb=i3_rowbytes(I); + for(int o=0;o>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); + dq[(int64_t)o*I+base+k]=(float)((int)u-4)*s[(int64_t)o*ng+g]; + } + } +} + +int main(void){ + const int Is[]={64,128,192,100,65,7168}; /* incl. short tail groups and one real GLM dim */ + enum { O=7, MAXI=7168 }; + static float w[(int64_t)O*MAXI], dq_ref[(int64_t)O*MAXI], dq_pk[(int64_t)O*MAXI]; + static float x[4*MAXI], y_ref[4*O], y_ker[4*O]; + static uint8_t q3[(int64_t)O*(MAXI/64+1)*24]; + static float sc[(int64_t)O*(MAXI/64+1)]; + + for(unsigned c=0;c unpack == reference quantize-dequantize, bit for bit */ + pack_int3_g64(w, q3, sc, O, I); + ref_i3_dequant(w, dq_ref, O, I); + unpack_i3(q3, sc, dq_pk, O, I); + int bad=0; + for(int64_t i=0;i<(int64_t)O*I;i++) if(dq_pk[i]!=dq_ref[i]) bad++; + CHECK(bad==0); + + /* 2. matmul_i3 == matmul over the dequantized reference (fp tolerance: + * NEON fma order differs from the scalar reference loop) */ + for(int S=1;S<=4;S+=3){ + for(int64_t i=0;i<(int64_t)S*I;i++) x[i]=rndf(); + matmul_i3(y_ker, x, q3, sc, S, I, O); + for(int s=0;s1?fabsf(y_ref[i]):1; + if(d/m>2e-4f){ CHECK(!"matmul_i3 mismatch"); break; } + } + } + + /* 3. QT plumbing: qt_alloc(bits=3) -> qt_fill -> matmul_qt & qt_bytes & helpers */ + QT t; qt_alloc(&t, O, I, 3); + CHECK(t.fmt==5); + qt_fill(&t, w, 3); + CHECK(qt_bytes(&t)==(int64_t)O*i3_rowbytes(I)+(int64_t)O*i3_groups(I)*4); + matmul_qt(y_ker, x, &t, 1); + for(int o=0;o1?fabsf((float)a):1; + CHECK(d/m<=2e-4f); + } + float acc[MAXI]; memset(acc,0,I*sizeof(float)); + qt_addrow(&t, 2, 0.5f, acc); + for(int i=0;i1?fabsf((float)a):1; + CHECK(d/m<=2e-4f); + } + + /* 4. format resolution through the #413 gate: fmt=5 is tagged by its distinct + * WEIGHT byte count. int3-g64 and grouped-int4-at-gs=64 carry the SAME scale + * cardinality O*ceil(I/64), so the pair (weight bytes, scale bytes) must + * disambiguate: same scales, int4 weights -> fmt=4/gs=64; int3 weights -> fmt=5. + * Only well-posed for I > 256: below that, O row scales legitimately match a + * 1-group grouped layout too (detect_group_size probes gs up to 256), so + * per-row vs grouped is not distinguishable from byte counts alone. */ + if(I>256){ int gs=-1; + int64_t ns_g64=(int64_t)O*i3_groups(I)*4, ns_row=(int64_t)O*4; + CHECK(qt_resolve_fmt("t.i3", O, I, (int64_t)O*i3_rowbytes(I), ns_g64, &gs)==5); + CHECK(gs==0); + CHECK(qt_resolve_fmt("t.i8", O, I, (int64_t)O*I, ns_row, &gs)==1); + CHECK(qt_resolve_fmt("t.i4", O, I, (int64_t)O*((I+1)/2), ns_row, &gs)==2); + CHECK(qt_resolve_fmt("t.i4g", O, I, (int64_t)O*((I+1)/2), ns_g64, &gs)==4); + CHECK(gs==64); + CHECK(qt_resolve_fmt("t.i2", O, I, (int64_t)O*((I+3)/4), ns_row, &gs)==3); } + free(t.q4); free(t.s); + } + + /* 5. quality in miniature: on rows with outliers, per-group int3 must beat + * per-row int4 on reconstruction RMS (the #132 finding this format ships). */ + { + int I=1024; + for(int64_t i=0;i<(int64_t)O*I;i++) w[i]=rndf()*0.02f; + for(int o=0;o>1]; + int v=(i&1)?((int)(b>>4)-8):((int)(b&0xF)-8); w4=(float)v*t4.s[o]; } + double d3=w[(int64_t)o*I+i]-dq_ref[(int64_t)o*I+i], d4=w[(int64_t)o*I+i]-w4; + e3+=d3*d3; e4+=d4*d4; + } + CHECK(e3 < e4); + printf(" outlier-rows RMS: int3-g64 %.3e < int4-row %.3e (ratio %.2f)\n", + sqrt(e3/((double)O*I)), sqrt(e4/((double)O*I)), sqrt(e4/e3)); + free(t4.q4); free(t4.s); + } + + if(fails){ printf("int3-g64 tests: %d FAILED\n", fails); return 1; } + printf("int3-g64 tests: ok\n"); + return 0; +} diff --git a/c/tests/test_int3_convert.py b/c/tests/test_int3_convert.py new file mode 100644 index 0000000..2eaa7f7 --- /dev/null +++ b/c/tests/test_int3_convert.py @@ -0,0 +1,70 @@ +"""quant_int3_g64 (tools/convert_fp8_to_int4.py): pack layout + round-trip. + +Decodes the packed bytes with an independent NumPy decoder implementing the +fmt=5 spec (16B low plane / 8B high plane per 64-group, v+4, per-group f32 +scale) and checks the dequantized result equals the reference +quantize-dequantize (same math as quant_ablation._quant_last_dim(3, 64)). +The C side of the same layout is covered by tests/test_int3.c. +""" +import os, sys, unittest +try: + import numpy as np +except ImportError: + raise unittest.SkipTest("numpy not installed") + +sys.path.insert(0, os.path.join(os.path.dirname(__file__), "..", "tools")) +from convert_fp8_to_int4 import quant_int3_g64 + + +def decode(packed, scales, O, I, group=64): + ng = (I + group - 1) // group + b = packed.reshape(O, ng, 24) + lo, hi = b[:, :, :16], b[:, :, 16:] + k = np.arange(group) + lov = (lo[:, :, k >> 2] >> ((k & 3) * 2)[None, None, :]) & 3 + hiv = (hi[:, :, k >> 3] >> (k & 7)[None, None, :]) & 1 + v = (lov | (hiv << 2)).astype(np.int64) - 4 + dq = v.astype(np.float64) * scales.reshape(O, ng, 1).astype(np.float64) + return dq.reshape(O, ng * group)[:, :I] + + +def reference(w, group=64): + """same math as quant_int3_g64 (which works in f32), replayed exactly, then + dequantized in f64 so it matches decode() bit for bit""" + O, I = w.shape + ng = (I + group - 1) // group + pad = ng * group - I + wp = np.pad(w, ((0, 0), (0, pad))) if pad else w + g = wp.reshape(O, ng, group) + s = np.maximum(np.abs(g).max(axis=2, keepdims=True) / 3.0, 1e-8).astype(np.float32) + q = np.clip(np.rint(g / s), -4, 3).astype(np.int64) + return (q.astype(np.float64) * s.astype(np.float64)).reshape(O, ng * group)[:, :I] + + +class Int3ConvertTest(unittest.TestCase): + def test_round_trip(self): + rng = np.random.default_rng(7) + for I in (64, 128, 100, 65, 7168): + w = (rng.standard_normal((5, I)) * 0.05).astype(np.float32) + w[0, 3] = 1.7; w[2, min(5, I - 1)] = -2.2 + packed, scales = quant_int3_g64(w) + ng = (I + 63) // 64 + self.assertEqual(packed.size, 5 * ng * 24) + self.assertEqual(scales.size, 5 * ng) + np.testing.assert_allclose(decode(packed, scales, 5, I), + reference(w), rtol=0, atol=0) + + def test_outliers_beat_row_int4(self): + rng = np.random.default_rng(11) + w = (rng.standard_normal((8, 1024)) * 0.02).astype(np.float32) + for o in range(8): w[o, (o * 37) % 1024] = 1.5 + packed, scales = quant_int3_g64(w) + e3 = float(((decode(packed, scales, 8, 1024) - w) ** 2).mean()) + s4 = np.maximum(np.abs(w).max(axis=1, keepdims=True) / 7.0, 1e-8) + w4 = np.clip(np.rint(w / s4), -8, 7) * s4 + e4 = float(((w4 - w) ** 2).mean()) + self.assertLess(e3, e4) + + +if __name__ == "__main__": + unittest.main() diff --git a/c/tests/test_int3_load.c b/c/tests/test_int3_load.c new file mode 100644 index 0000000..57634c0 --- /dev/null +++ b/c/tests/test_int3_load.c @@ -0,0 +1,103 @@ +/* Loader-seam test for fmt=5: writes a real .safetensors file containing an + * int3-g64 tensor (U8 payload + per-GROUP .qs) next to an int4 control tensor + * (per-row .qs), indexes it with st_init, loads both through qt_from_disk, and + * checks the byte-count/.qs-size format inference picks fmt=5 vs fmt=2 correctly + * and the loaded weights dequantize identically to pack_int3_g64's output. */ +#define main coli_glm_main_unused +#include "../colibri.c" +#undef main + +#include +#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 = 0xA5A5A5A55A5A5A5Aull; +static float rndf(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; + return ((int64_t)(rng & 0xFFFFF) - 0x80000) / (float)0x80000; } + +static void deq4(const QT *t, float *dq){ + for(int o=0;oO;o++) for(int i=0;iI;i++){ + if(t->fmt==5){ + int64_t g=i/I3_GROUP; const uint8_t *lo=t->q4+(int64_t)o*i3_rowbytes(t->I)+g*I3_GBYTES, *hi=lo+16; + int k=i%I3_GROUP; + unsigned u=((lo[k>>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); + dq[(int64_t)o*t->I+i]=(float)((int)u-4)*t->s[(int64_t)o*i3_groups(t->I)+g]; + } else { /* fmt2 */ + uint8_t b=t->q4[(int64_t)o*((t->I+1)/2)+(i>>1)]; + int v=(i&1)?((int)(b>>4)-8):((int)(b&0xF)-8); + dq[(int64_t)o*t->I+i]=(float)v*t->s[o]; + } + } +} + +int main(void){ + enum { O=5, I=320 }; /* 5 groups per row; I > 256 so the per-row int4 + * control stays fmt=2 (detect_group_size probes + * gs up to 256: any smaller I would make O row + * scales match a legitimate 1-group layout) */ + int64_t ng=i3_groups(I), rb=i3_rowbytes(I); + static float w[O*I]; + for(int i=0;i (qbytes U8 [O*ceil(I/64)*24], scales f32 [O*ceil(I/64)]) + """int3 with PER-GROUP scales (fmt=5 in colibri.c): per 64-input group, symmetric absmax + (qmax=3, clamp [-4,3], stored v+4), packed as 16B low plane (2 bits/val, int2 layout) + + 8B high plane (1 bit/val). Same math as quant_ablation._quant_last_dim(bits=3, + group=64) (#132), here with real packing. 3.5 bits/weight effective.""" + O, I = w.shape + ng = (I + group - 1) // group + pad = ng * group - I + wp = np.pad(w, ((0, 0), (0, pad))) if pad else w + g = wp.reshape(O, ng, group) + amax = np.abs(g).max(axis=2, keepdims=True) + s = np.maximum(amax / 3.0, 1e-8) + q = (np.clip(np.rint(g / s), -4, 3).astype(np.int32) + 4).astype(np.uint8) # 0..7 + if pad: q[:, -1, group - pad:] = 4 # pad packs as 0 after -4 + lo = np.zeros((O, ng, 16), np.uint8) + for k in range(4): + lo |= ((q[:, :, k::4] & 3) << (k * 2)).astype(np.uint8) + hi = np.zeros((O, ng, 8), np.uint8) + for b in range(8): + hi |= (((q[:, :, b::8] >> 2) & 1) << b).astype(np.uint8) + out = np.concatenate([lo, hi], axis=2) # [O, ng, 24] + return out.reshape(-1), s[:, :, 0].astype(np.float32).reshape(-1) + def quant_int2(w, bits): # -> (qbytes U8 [O*ceil(I/4)], scale f32 [O]); 4/byte O, I = w.shape qmax = (1 << (bits - 1)) - 1 # bits=2 -> qmax=1, valori [-2,1] @@ -214,6 +237,14 @@ def dequant(f, name, keys): return (w * sc).numpy() return f.get_tensor(name).to(torch.float32).numpy() +# Per-projection bit overrides for ROUTED experts (gate_proj/up_proj/down_proj), set from +# --up-bits/--gate-bits/--down-bits in main(). Empty = uniform xbits. Motivated by the +# measured result that up_proj tolerates int3-g64 at ~zero quality cost while int2 craters +# (OLMoE ablation, PR #168 comment): up-only int3 drops ~8% of expert bytes for free. +# NB: the resume manifests (check_or_record_params and the --indir progress file) already +# record dict(PROJ_BITS) — this global is the definition those sites depend on. +PROJ_BITS = {} + def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits, keep_mtp=False, keep_idx=False, group_size=0, bits_map=None): from safetensors import safe_open @@ -235,9 +266,16 @@ def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits, # Any unknown kind that fell through classify as "q" if bits_map and kind not in bits_map and kind not in ("io", "x", "sh", "o", "kvb", "attn", "dmlp"): bits = ebits + # Per-projection override for routed experts, applied on top of the type-level bits. + if kind == "x" and PROJ_BITS: # e.g. up_proj -> 3 (int3-g64) while gate/down stay 4 + for proj, pb in PROJ_BITS.items(): + if f".{proj}.weight" in name: bits = pb; break if w.ndim != 2: # es. bias 1D non previsto come 'q' -> tienilo f32 out_dict[name] = w.astype(np.float32); continue - if group_size > 0 and bits <= 4: + if bits == 3: + # int3-g64 (fmt=5): inherently group-64, distinct from grouped-int4. + q, s = quant_int3_g64(w) + elif group_size > 0 and bits <= 4: q, s = quant_int4_grouped(w, bits, group_size) else: q, s = (quant_int2(w, bits) if bits <= 2 else @@ -295,6 +333,13 @@ def main(): help="bits for dense MLP (first 3 layers). Default=ebits") ap.add_argument("--group-size", type=int, default=0, # 0 = per-row (backward compat); 128 = group-scaled help="group size for int4 scales: 0=per-row (default), 128=one scale per 128 elements (much better quality)") + # Per-projection bit overrides for routed experts (orthogonal to the type-level flags above). + ap.add_argument("--up-bits", type=int, default=None, + help="bits for up_proj in routed experts (e.g. 3 = int3-g64). Default=xbits") + ap.add_argument("--gate-bits", type=int, default=None, + help="bits for gate_proj in routed experts. Default=xbits") + ap.add_argument("--down-bits", type=int, default=None, + help="bits for down_proj in routed experts. Default=xbits") ap.add_argument("--n-layers", type=int, default=78) ap.add_argument("--min-free-gb", type=float, default=20.0) ap.add_argument("--selftest", action="store_true") @@ -324,6 +369,10 @@ def main(): "embedding half -> MTP acceptance ~0% (issue #8). Use the default --ebits 8, " "or add --group-size 128 for group-scaled int4.") if a.xbits is None: a.xbits = a.ebits + for proj, val in (("gate_proj", a.gate_bits), ("up_proj", a.up_bits), ("down_proj", a.down_bits)): + if val is not None: PROJ_BITS[proj] = val + if PROJ_BITS: + print(f"[per-projection expert bits] {PROJ_BITS} (others -> xbits={a.xbits})") # Build per-type bits map. If a type-specific arg is set, use it; otherwise the # converter falls back to ebits for that type.