int3-g64 (fmt=5): per-group-scale 3-bit weight format — engine, converter, tests
New weight format: int3 with ONE f32 scale per 64-input group (3.5 bits/weight effective). Per group: 16B low plane (2 bits/val, int2 layout) + 8B high plane (1 bit/val), values [-4,3] stored v+4. Same quantization math as tools/quant_ablation.py _quant_last_dim(bits=3, group=64) from #132, whose OLMoE ablation measured int3-g64 BEATING the shipped per-row int4 on quality (-7.5 vs -9.3pp) at ~14% fewer bits. Engine (placed per the #391 split): matmul_i3 + pack_int3_g64 + I3_* layout helpers in quant.h next to their kernel family; fmt=5 branches in colibri.c's qt_bytes/qt_alloc/qt_fill/matmul_qt/embed_row/qt_addrow/qt_matvec_rows. Format detection now goes through #413's qt_resolve_fmt: fmt=5 registers its distinct weight-byte layout O*ceil(I/64)*24 and its scale cardinality O*ceil(I/64) there, validated against [O,I] like every other format. int3-g64 and grouped-int4-at-gs=64 carry the SAME scale count, so the weight bytes are the int3 tag; row formats keep precedence for the small-I shapes where byte counts coincide. The io_uring expert path still used the raw ?1:?2:3 byte inference (it missed fmt=4 grouping entirely and never set gs) — converted to qt_resolve_fmt like the other two expert paths. Backends: qt_cuda_upload returns 0 for fmt=5 (tensor stays CPU-side, the documented fallback), the dense CUDA matmul gate excludes fmt=5, and Metal's existing fmt gates (gemm fmt<=3, moe fmt 1/2) already reject it. Converter: quant_int3_g64 in convert_fp8_to_int4.py; --ebits 3/--xbits 3 now emit it (previously bits=3 silently produced int4). Tests: tests/test_int3.c (bit-exact pack/unpack vs reference, matmul_i3 vs dequant-matmul incl. short tail groups and the real GLM I=7168, QT plumbing, qt_resolve_fmt disambiguation incl. the same-scale-count fmt=4/fmt=5 pair, outlier-rows RMS: int3-g64 3.3x lower error than per-row int4), tests/test_int3_load.c (hand-rolled .safetensors fixture through st_init + qt_from_disk: fmt=5 detected and loaded next to an int4 control tensor), tests/test_int3_convert.py (NumPy pack round-trip vs independent decoder). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This commit is contained in:
FABIOTESS
+7
-1
@@ -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)
|
||||
|
||||
|
||||
+59
-10
@@ -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<D;i+=2){ uint8_t byte=q[i>>1]; x[i]=(float)((int)(byte&0xF)-8)*s;
|
||||
if(i+1<D) x[i+1]=(float)((int)(byte>>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<ng; g++){ const uint8_t *lo=q+g*I3_GBYTES, *hi=lo+16;
|
||||
int base=(int)(g*I3_GROUP), n=D-base<I3_GROUP?D-base:I3_GROUP;
|
||||
for(int k=0;k<n;k++){ unsigned u=((lo[k>>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<D;i++){ uint8_t byte=q[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<ng; g++){ const uint8_t *lo=w+g*I3_GBYTES, *hi=lo+16;
|
||||
float cg=coef*sr[g]; int base=(int)(g*I3_GROUP), n=I-base<I3_GROUP?I-base:I3_GROUP;
|
||||
for(int k=0;k<n;k++){ unsigned u=((lo[k>>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;i<I;i++) acc[i]+=c*(float)w[i]; return; }
|
||||
if(t->fmt==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<end;i++){ uint8_t b=w[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<ng; g++){ const uint8_t *lo=w+g*I3_GBYTES, *hi=lo+16;
|
||||
int base=(int)(g*I3_GROUP), n=I-base<I3_GROUP?I-base:I3_GROUP; float acc=0;
|
||||
for(int k=0;k<n;k++){ unsigned u=((lo[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;
|
||||
for(int i=0;i<I;i++){ uint8_t b=w[i>>2]; acc+=((int)((b>>((i&3)*2))&3)-2)*x[i]; } a=acc*s; }
|
||||
y[j]=(float)a;
|
||||
|
||||
@@ -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<O;o++){
|
||||
const uint8_t *wrow=q3+(int64_t)o*rb;
|
||||
const float *srow=scale+(int64_t)o*ng;
|
||||
for(int s=0;s<S;s++){
|
||||
const float *xs=x+(int64_t)s*I;
|
||||
float acc=0;
|
||||
for(int64_t g=0; g<ng; g++){
|
||||
const uint8_t *lo=wrow+g*I3_GBYTES, *hi=lo+16;
|
||||
int base=(int)(g*I3_GROUP), n = I-base < I3_GROUP ? I-base : I3_GROUP;
|
||||
float a=0; int k=0;
|
||||
#if defined(__ARM_NEON)
|
||||
if(n==I3_GROUP){
|
||||
const uint8x8_t m2v=vdup_n_u8(3); const int8x16_t b4q=vdupq_n_s8(4);
|
||||
const uint8x16_t bitm={1,2,4,8,16,32,64,128,1,2,4,8,16,32,64,128};
|
||||
const uint8x16_t fourq=vdupq_n_u8(4);
|
||||
float32x4_t ac0=vdupq_n_f32(0), ac1=vdupq_n_f32(0);
|
||||
for(;k+16<=I3_GROUP;k+=16){
|
||||
uint32_t wd; memcpy(&wd, lo+(k>>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<n;k++){
|
||||
unsigned u=((lo[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;o<O;o++){
|
||||
const float *wr=w+(int64_t)o*I;
|
||||
uint8_t *qr=q3+(int64_t)o*rb;
|
||||
float *sr=scale+(int64_t)o*ng;
|
||||
for(int64_t g=0; g<ng; g++){
|
||||
int base=(int)(g*I3_GROUP), n = I-base < I3_GROUP ? I-base : I3_GROUP;
|
||||
float amax=0;
|
||||
for(int k=0;k<n;k++){ float a=fabsf(wr[base+k]); if(a>amax)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;k<n;k++){
|
||||
int v=(int)lrintf(wr[base+k]/s); if(v>3)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)
|
||||
|
||||
@@ -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 <stdio.h>
|
||||
#include <string.h>
|
||||
#include <math.h>
|
||||
|
||||
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;o<O;o++) for(int64_t g=0;g<ng;g++){
|
||||
int base=(int)(g*I3_GROUP), n=I-base<I3_GROUP?I-base:I3_GROUP;
|
||||
float amax=0; for(int k=0;k<n;k++){ float a=fabsf(w[(int64_t)o*I+base+k]); if(a>amax)amax=a; }
|
||||
float s=amax/3.f; if(s<1e-8f)s=1e-8f;
|
||||
for(int k=0;k<n;k++){
|
||||
int v=(int)lrintf(w[(int64_t)o*I+base+k]/s); if(v>3)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<O;o++) for(int64_t g=0;g<ng;g++){
|
||||
const uint8_t *lo=q3+(int64_t)o*rb+g*I3_GBYTES, *hi=lo+16;
|
||||
int base=(int)(g*I3_GROUP), n=I-base<I3_GROUP?I-base:I3_GROUP;
|
||||
for(int k=0;k<n;k++){
|
||||
unsigned u=((lo[k>>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<sizeof Is/sizeof *Is;c++){
|
||||
int I=Is[c];
|
||||
for(int64_t i=0;i<(int64_t)O*I;i++) w[i]=rndf()*0.05f;
|
||||
w[3]=1.7f; w[(int64_t)2*I+5]=-2.2f; /* outliers */
|
||||
|
||||
/* 1. pack -> 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;s<S;s++) for(int o=0;o<O;o++){
|
||||
double a=0; for(int i=0;i<I;i++) a+=(double)dq_ref[(int64_t)o*I+i]*x[(int64_t)s*I+i];
|
||||
y_ref[s*O+o]=(float)a;
|
||||
}
|
||||
for(int i=0;i<S*O;i++){
|
||||
float d=fabsf(y_ker[i]-y_ref[i]), m=fabsf(y_ref[i])>1?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;o<O;o++){
|
||||
double a=0; for(int i=0;i<I;i++) a+=(double)dq_ref[(int64_t)o*I+i]*x[i];
|
||||
float d=fabsf(y_ker[o]-(float)a), m=fabsf((float)a)>1?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;i<I;i++) CHECK(fabsf(acc[i]-0.5f*dq_ref[(int64_t)2*I+i])<=1e-6f);
|
||||
float yr[3];
|
||||
qt_matvec_rows(&t, 1, 3, x, yr);
|
||||
for(int j=0;j<3;j++){
|
||||
double a=0; for(int i=0;i<I;i++) a+=(double)dq_ref[(int64_t)(1+j)*I+i]*x[i];
|
||||
float d=fabsf(yr[j]-(float)a), m=fabsf((float)a)>1?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<O;o++) w[(int64_t)o*I+(o*37)%I]=1.5f; /* one outlier per row */
|
||||
ref_i3_dequant(w, dq_ref, O, I);
|
||||
QT t4; qt_alloc(&t4, O, I, 4); qt_fill(&t4, w, 4);
|
||||
double e3=0, e4=0;
|
||||
for(int o=0;o<O;o++) for(int i=0;i<I;i++){
|
||||
float w4; { const uint8_t *q=t4.q4+(int64_t)o*((I+1)/2); uint8_t b=q[i>>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;
|
||||
}
|
||||
@@ -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()
|
||||
@@ -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 <stdio.h>
|
||||
#include <string.h>
|
||||
#include <sys/stat.h>
|
||||
#include <unistd.h>
|
||||
|
||||
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;o<t->O;o++) for(int i=0;i<t->I;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<O*I;i++) w[i]=rndf()*0.05f;
|
||||
w[7]=1.9f;
|
||||
|
||||
static uint8_t q3[O*(I/64)*24]; static float s3[O*(I/64)];
|
||||
pack_int3_g64(w, q3, s3, O, I);
|
||||
static uint8_t q4b[O*((I+1)/2)]; static float s4[O];
|
||||
pack_int4(w, q4b, s4, O, I, 4);
|
||||
|
||||
/* write a minimal single-shard safetensors file */
|
||||
const char *dir="tests/tmp_int3_snap";
|
||||
#ifdef _WIN32
|
||||
mkdir(dir);
|
||||
#else
|
||||
mkdir(dir, 0755);
|
||||
#endif
|
||||
char path[256]; snprintf(path,sizeof path,"%s/model.safetensors",dir);
|
||||
int64_t nb3=(int64_t)O*rb, ns3=(int64_t)O*ng*4, nb4=(int64_t)O*((I+1)/2), ns4=(int64_t)O*4;
|
||||
char hdr[1024];
|
||||
int hl=snprintf(hdr,sizeof hdr,
|
||||
"{\"w3\":{\"dtype\":\"U8\",\"shape\":[%lld],\"data_offsets\":[0,%lld]},"
|
||||
"\"w3.qs\":{\"dtype\":\"F32\",\"shape\":[%lld],\"data_offsets\":[%lld,%lld]},"
|
||||
"\"w4\":{\"dtype\":\"U8\",\"shape\":[%lld],\"data_offsets\":[%lld,%lld]},"
|
||||
"\"w4.qs\":{\"dtype\":\"F32\",\"shape\":[%lld],\"data_offsets\":[%lld,%lld]}}",
|
||||
(long long)nb3,(long long)nb3,
|
||||
(long long)(O*ng),(long long)nb3,(long long)(nb3+ns3),
|
||||
(long long)nb4,(long long)(nb3+ns3),(long long)(nb3+ns3+nb4),
|
||||
(long long)O,(long long)(nb3+ns3+nb4),(long long)(nb3+ns3+nb4+ns4));
|
||||
FILE *f=fopen(path,"wb");
|
||||
if(!f){ printf("FAIL: cannot create %s (run from c/, like tools/run_tests.py does)\n", path); return 1; }
|
||||
uint64_t hlen=(uint64_t)hl;
|
||||
fwrite(&hlen,8,1,f); fwrite(hdr,1,hl,f);
|
||||
fwrite(q3,1,(size_t)nb3,f); fwrite(s3,1,(size_t)ns3,f);
|
||||
fwrite(q4b,1,(size_t)nb4,f); fwrite(s4,1,(size_t)ns4,f);
|
||||
fclose(f);
|
||||
|
||||
static Model gm; /* only gm.S is used by qt_from_disk */
|
||||
st_init(&gm.S, dir);
|
||||
|
||||
QT t3; memset(&t3,0,sizeof t3);
|
||||
qt_from_disk(&gm,"w3",O,I,8,0,&t3);
|
||||
CHECK(t3.fmt==5);
|
||||
static float dq_load[O*I], dq_ref[O*I];
|
||||
deq4(&t3,dq_load);
|
||||
QT tr={.fmt=5,.q4=q3,.s=s3,.O=O,.I=I};
|
||||
deq4(&tr,dq_ref);
|
||||
CHECK(memcmp(dq_load,dq_ref,sizeof dq_ref)==0);
|
||||
|
||||
QT t4; memset(&t4,0,sizeof t4);
|
||||
qt_from_disk(&gm,"w4",O,I,8,0,&t4);
|
||||
CHECK(t4.fmt==2); /* control: row-scale int4 still detected */
|
||||
deq4(&t4,dq_load);
|
||||
QT tr4={.fmt=2,.q4=q4b,.s=s4,.O=O,.I=I};
|
||||
deq4(&tr4,dq_ref);
|
||||
CHECK(memcmp(dq_load,dq_ref,sizeof dq_ref)==0);
|
||||
|
||||
unlink(path); rmdir(dir);
|
||||
if(fails){ printf("int3 loader tests: %d FAILED\n", fails); return 1; }
|
||||
printf("int3 loader tests: ok\n");
|
||||
return 0;
|
||||
}
|
||||
@@ -83,6 +83,29 @@ def quant_int4_grouped(w, bits, gs=128):
|
||||
s_flat = s[:, :, 0].astype(np.float32).reshape(-1)
|
||||
return out.reshape(-1), s_flat
|
||||
|
||||
def quant_int3_g64(w, bits=3, group=64): # -> (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]
|
||||
@@ -237,7 +260,10 @@ def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits,
|
||||
bits = ebits
|
||||
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
|
||||
|
||||
Reference in New Issue
Block a user