Merge dev into dual-ssd-mirror: resolve #298/#192-era conflicts
Combined resolution: mir_pread/st_prefetch_rep (this PR) now carry dev's DISK-CLASS accounting unwind (dc_wall_exit) and O_DIRECT prefetch skip (g_direct); direct-path keeps dc_direct=1 plus the mirror read counters. Verified: clean build, token-exact tiny models unchanged, test_st_mirror and test_st_pread pass. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
This commit is contained in:
JustVugg
@@ -0,0 +1,98 @@
|
||||
# Efficiency suite — regression tests + optimization dossier
|
||||
|
||||
Two layers:
|
||||
|
||||
1. **`test_inefficiency.py`** — tiny-model *asserted* regression tests. Fast
|
||||
(~0.15s/run), gate CI, catch breakage. Run as part of `make test`.
|
||||
2. **`test_efficiency_report.py`** — an *opt-in optimization dossier* for a real
|
||||
model. Runs every instrumentation flag, prints a 9-section report answering
|
||||
*what is doing what, when, with what, is it inefficient, how to improve*.
|
||||
Never fails CI (it's a report, not a gate).
|
||||
|
||||
## The dossier (what you run when optimizing)
|
||||
|
||||
```bash
|
||||
# CPU-only (safe, fast to validate):
|
||||
COLI_EFFICIENCY_MODEL=../glm52_i4_g64 make efficiency-report
|
||||
|
||||
# CUDA (dense + expert tiers — needs a CUDA build, see below):
|
||||
COLI_EFFICIENCY_MODEL=../glm52_i4_g64 COLI_EFFICIENCY_CUDA=1 make efficiency-report
|
||||
```
|
||||
|
||||
It turns ON every observability flag the engine supports — `PROF=1`,
|
||||
`COLI_CUDA_PROFILE=1`, `CACHE_ROUTE=1` (auto-unlocks `route_agree`/`route_kl`),
|
||||
`DISK_SPLIT=1`, `LOOKA=1` — so nothing the engine can tell you is left dark.
|
||||
None of these change the computed output; they only add telemetry.
|
||||
|
||||
The 9 sections, and the question each answers:
|
||||
|
||||
| § | section | answers |
|
||||
|---|---|---|
|
||||
| 1 | PROVENANCE | what is running, on what CPU/backend, with what effective config |
|
||||
| 2 | THROUGHPUT | tok/s + forward-latency p50/p90/p99/max (is the tail healthy?) |
|
||||
| 3 | WHERE TIME GOES | the 5 PROFILE phases as % of decode + absolute seconds + verdict |
|
||||
| 3a | ATTENTION BREAKDOWN | attention split into projection/RoPE, score-softmax-value, output |
|
||||
| 4 | EXPERT CACHE | hit %, experts-loaded/token vs baseline topk |
|
||||
| 5 | DISK I/O | GB fetched, MB/token, GB/s, read-service vs felt-wait, phase split |
|
||||
| 5a | DISK-LOAD SPLIT | loads by decode phase (draft/absorb/verify) + MTP-vs-main bytes |
|
||||
| 6 | ROUTING QUALITY | route_agree %, route_kl, cache swaps |
|
||||
| 6a | ROUTING PREDICTABILITY | LOOKAHEAD recall per predictor (which prefetch wins) |
|
||||
| 7 | SPECULATION | tokens/forward, MTP acceptance % |
|
||||
| 8 | GPU TIERS | resident tensors, expert tier (count/GB/calls), H2D/kernel/D2H ms |
|
||||
|
||||
Every line that crosses an advisory threshold is marked `[FLAG]` with the
|
||||
concrete lever to pull (raise RAM_GB, add PIN_GB, try DIRECT=1, lower CTX, …),
|
||||
and all flags repeat in a summary at the end.
|
||||
|
||||
## Tunable thresholds
|
||||
|
||||
The `IS IT INEFFICIENT?` lines are advisory constants at the top of
|
||||
`test_efficiency_report.py`:
|
||||
|
||||
| constant | default | meaning |
|
||||
|---|---|---|
|
||||
| `DISK_WAIT_DOMINANT` | 0.40 | >40% decode waiting on expert reads → I/O-bound |
|
||||
| `LOW_HIT_RATE` | 0.30 | <30% cache hit → thrashing |
|
||||
| `LOW_ROUTE_AGREE` | 0.80 | <80% routing overlap → prefetch guessing wrong |
|
||||
| `HIGH_TAIL_RATIO` | 3.0 | p99 > 3× p50 → decode stalls |
|
||||
| `LOW_MTP_ACCEPT` | 0.20 | <20% MTP acceptance → draft decoder is dead weight |
|
||||
|
||||
The tiny-model asserted floors live in `tools/efficiency.py` (`TINY_TOK_S_FLOOR`,
|
||||
`MAX_DISK_WAIT_SHARE`, `MIN_CPU_CUDA_AGREEMENT`).
|
||||
|
||||
## The regression tests (what gates CI)
|
||||
|
||||
`test_inefficiency.py` runs on the bundled `glm_tiny` model and asserts:
|
||||
|
||||
- telemetry parses (no format drift)
|
||||
- tiny tok/s ≥ floor (throughput regression)
|
||||
- PROFILE phases present and non-negative (accounting sanity)
|
||||
- disk-wait not dominant on a resident model (I/O-path regression)
|
||||
- CPU determinism (two greedy runs agree)
|
||||
- **CUDA** (skip unless CUDA built): init path, dense uploads VRAM, CPU-vs-CUDA
|
||||
argmax agreement ≥ 70% (kernel-correctness guard)
|
||||
|
||||
```bash
|
||||
make efficiency # tiny CPU tests
|
||||
make efficiency-cuda # tiny CUDA tests (needs CUDA build)
|
||||
```
|
||||
|
||||
## CUDA build prerequisite
|
||||
|
||||
The default `make glm.exe` builds **without** CUDA. The CUDA tests and the CUDA
|
||||
dossier need a host built with `-DCOLI_CUDA` plus the runtime DLL:
|
||||
|
||||
```bash
|
||||
make clean && make glm.exe CUDA_DLL=1 && make cuda-dll
|
||||
```
|
||||
|
||||
`make efficiency-cuda` auto-skips with a clear message if the host is CPU-only
|
||||
(it scans the binary for the "CPU-only" marker the engine embeds).
|
||||
|
||||
## Files
|
||||
|
||||
- `tools/efficiency.py` — shared harness: `parse_run()` (captures every
|
||||
telemetry signal), `run_engine()`, thresholds. Reuses `PROFILE_RE`/`SPEED_RE`
|
||||
from `tools/benchmark_cuda_fixture.py`.
|
||||
- `tests/test_inefficiency.py` — tiny-model asserted tests (CPU + CUDA).
|
||||
- `tests/test_efficiency_report.py` — the opt-in optimization dossier.
|
||||
@@ -0,0 +1,92 @@
|
||||
/* Microbenchmark: old (single-accumulator) vs new (independent-accumulator) AVX-VNNI
|
||||
* int8/int4 dot kernels (quant.h). NOT a unit test -- test_idot.c proves correctness.
|
||||
* This measures the headline claim: breaking the serial vpdpbusd->acc chain lifts
|
||||
* per-core kernel throughput, the same win the NEON path already took ("26->63 GB/s, 2.4x").
|
||||
*
|
||||
* It re-implements the OLD single-acc AVX-VNNI kernels inline and calls the REAL (new)
|
||||
* ones via the include-colibri.c pattern -- one process, identical frozen inputs, warm
|
||||
* caches. Reports median ns/call + GB/s-of-weights + new/old ratio. Measures the kernel's
|
||||
* compute ceiling (warm caches), NOT end-to-end tok/s.
|
||||
*
|
||||
* Run: make tests/bench_idot ARCH=native && ./tests/bench_idot (not in TEST_BINS -- not a gate)
|
||||
*/
|
||||
#define main coli_glm_main_unused
|
||||
#include "../colibri.c"
|
||||
#undef main
|
||||
#include <stdint.h>
|
||||
#include <string.h>
|
||||
|
||||
static uint32_t rs=0x2545F491u;
|
||||
static uint32_t xr(void){ rs^=rs<<13; rs^=rs>>17; rs^=rs<<5; return rs; }
|
||||
|
||||
/* ---- OLD kernels: verbatim copies of the pre-change __AVXVNNI__ branches (single acc) ---- */
|
||||
#if defined(__AVXVNNI__) && defined(__AVX2__)
|
||||
static int32_t dot_i8i8_old(const int8_t *w, const int8_t *x, int I){
|
||||
int32_t sum=0; int i=0;
|
||||
__m128i acc=_mm_setzero_si128();
|
||||
for(;i+16<=I;i+=16){
|
||||
__m128i wv=_mm_loadu_si128((const __m128i*)(w+i));
|
||||
__m128i xv=_mm_loadu_si128((const __m128i*)(x+i));
|
||||
__m128i xs=_mm_sign_epi8(xv,wv);
|
||||
acc=_mm_dpbusd_epi32(acc,_mm_abs_epi8(wv),xs);
|
||||
}
|
||||
sum=hsum128_i32(acc);
|
||||
for(;i<I;i++) sum+=(int32_t)w[i]*x[i];
|
||||
return sum;
|
||||
}
|
||||
static int32_t dot_i4i8_old(const uint8_t *w4, const int8_t *x, int I){
|
||||
int32_t sum=0; int i=0;
|
||||
const __m128i m4=_mm_set1_epi8(0x0F); const __m128i b8=_mm_set1_epi8(8);
|
||||
__m128i acc=_mm_setzero_si128();
|
||||
for(;i+32<=I;i+=32){
|
||||
__m128i by=_mm_loadu_si128((const __m128i*)(w4+(i>>1)));
|
||||
__m128i lo=_mm_and_si128(by,m4), hi=_mm_and_si128(_mm_srli_epi16(by,4),m4);
|
||||
__m128i n0=_mm_unpacklo_epi8(lo,hi), n1=_mm_unpackhi_epi8(lo,hi);
|
||||
__m128i w0=_mm_sub_epi8(n0,b8), w1=_mm_sub_epi8(n1,b8);
|
||||
__m128i x0=_mm_loadu_si128((const __m128i*)(x+i));
|
||||
__m128i x1=_mm_loadu_si128((const __m128i*)(x+i+16));
|
||||
acc=_mm_dpbusd_epi32(acc,_mm_abs_epi8(w0),_mm_sign_epi8(x0,w0));
|
||||
acc=_mm_dpbusd_epi32(acc,_mm_abs_epi8(w1),_mm_sign_epi8(x1,w1));
|
||||
}
|
||||
sum=hsum128_i32(acc);
|
||||
for(;i<I;i++){ uint8_t b=w4[i>>1]; int v=(i&1)?((int)(b>>4)-8):((int)(b&0xF)-8); sum+=v*x[i]; }
|
||||
return sum;
|
||||
}
|
||||
#else
|
||||
#error "bench_idot requires an AVX-VNNI build: make tests/bench_idot ARCH=native on an AVX-VNNI CPU"
|
||||
#endif
|
||||
|
||||
#define I_DIM 6144
|
||||
#define N_REPEAT 20000
|
||||
static int cmp_d(const void*a,const void*b){ double x=*(const double*)a,y=*(const double*)b; return x<y?-1:x>y?1:0; }
|
||||
|
||||
int main(void){
|
||||
static int8_t w8[I_DIM], x8[I_DIM]; static uint8_t w4[I_DIM/2];
|
||||
for(int i=0;i<I_DIM;i++){ w8[i]=(int8_t)(xr()&0xFF); x8[i]=(int8_t)((int)(xr()%255)-127); }
|
||||
for(int i=0;i<I_DIM/2;i++) w4[i]=(uint8_t)(xr()&0xFF);
|
||||
|
||||
/* correctness sanity: new must equal old (bit-exact) */
|
||||
if(dot_i8i8(w8,x8,I_DIM)!=dot_i8i8_old(w8,x8,I_DIM)){ fprintf(stderr,"MISMATCH i8i8\n"); return 1; }
|
||||
if(dot_i4i8(w4,x8,I_DIM)!=dot_i4i8_old(w4,x8,I_DIM)){ fprintf(stderr,"MISMATCH i4i8\n"); return 1; }
|
||||
|
||||
static double t[N_REPEAT]; volatile int32_t sink=0;
|
||||
const char *names[4]={"i8i8 old","i8i8 new","i4i8 old","i4i8 new"};
|
||||
double gbs[4];
|
||||
for(int k=0;k<4;k++){
|
||||
for(int wi=0;wi<200;wi++) sink+= (k==0)?dot_i8i8_old(w8,x8,I_DIM):(k==1)?dot_i8i8(w8,x8,I_DIM)
|
||||
:(k==2)?dot_i4i8_old(w4,x8,I_DIM):dot_i4i8(w4,x8,I_DIM); /* warmup */
|
||||
for(int r=0;r<N_REPEAT;r++){
|
||||
double t0=now_s();
|
||||
sink+= (k==0)?dot_i8i8_old(w8,x8,I_DIM):(k==1)?dot_i8i8(w8,x8,I_DIM)
|
||||
:(k==2)?dot_i4i8_old(w4,x8,I_DIM):dot_i4i8(w4,x8,I_DIM);
|
||||
t[r]=(now_s()-t0)*1e9;
|
||||
}
|
||||
qsort(t,N_REPEAT,sizeof(double),cmp_d);
|
||||
double med=t[N_REPEAT/2];
|
||||
double bytes=(k<2)?I_DIM:(double)I_DIM/2; /* weight bytes touched */
|
||||
gbs[k]=bytes/med;
|
||||
printf("%-9s %8.1f ns/call %6.2f GB/s\n", names[k], med, gbs[k]);
|
||||
}
|
||||
printf("ratio i8i8 new/old: %.2fx | ratio i4i8 new/old: %.2fx\n", gbs[1]/gbs[0], gbs[3]/gbs[2]);
|
||||
(void)sink; return 0;
|
||||
}
|
||||
@@ -30,9 +30,9 @@ int main(){
|
||||
for(int i=0;i<D;i++)ds[i]=0.006f+(i%7)*0.0002f;
|
||||
for(size_t i=0;i<x.size();i++)x[i]=std::sin((float)(i+1)*0.013f)*2.f;
|
||||
ColiCudaTensor *g=nullptr,*u=nullptr,*d=nullptr;
|
||||
if(!coli_cuda_tensor_upload(&g,hidden.data(),hs.data(),2,D,I,device)||
|
||||
!coli_cuda_tensor_upload(&u,hidden.data(),hs.data(),2,D,I,device)||
|
||||
!coli_cuda_tensor_upload(&d,down.data(),ds.data(),2,I,D,device))return 2;
|
||||
if(!coli_cuda_tensor_upload(&g,hidden.data(),hs.data(),2,D,I,device,0)||
|
||||
!coli_cuda_tensor_upload(&u,hidden.data(),hs.data(),2,D,I,device,0)||
|
||||
!coli_cuda_tensor_upload(&d,down.data(),ds.data(),2,I,D,device,0))return 2;
|
||||
for(int rows: {1,2,4,8}){
|
||||
double scalar=run(g,u,d,x.data(),a.data(),rows,3,0);
|
||||
double packed=run(g,u,d,x.data(),b.data(),rows,3,1);
|
||||
|
||||
@@ -50,10 +50,56 @@ int main(int argc, char **argv) {
|
||||
if (coli_cuda_tensor_upload(&t8, q8, s8, 1, 5, 2, d0)) return 1;
|
||||
if (ndev > 1 && coli_cuda_tensor_upload(&t8, q8, s8, 1, 4, 2, d1)) return 1;
|
||||
if (!coli_cuda_matmul(&t8, got, x, q8, s8, 1, 2, 4, 2, d0) || !close_enough(got, want8, 4)) return 1;
|
||||
/* Cached tensor must stay callable without live host pointers
|
||||
* (CUDA_RELEASE_HOST slots null theirs after upload) — including
|
||||
* SUSTAINED reuse, not just the first call. */
|
||||
for (int rep = 0; rep < 64; rep++)
|
||||
if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
|
||||
!close_enough(got, want8, 4)) return 1;
|
||||
/* A tensor uploaded from a TEMPORARY host buffer must survive the buffer
|
||||
* being scribbled and freed (the release-host lifecycle). */
|
||||
{
|
||||
int8_t *tmpw = static_cast<int8_t *>(std::malloc(8));
|
||||
float *tmps = static_cast<float *>(std::malloc(2 * sizeof(float)));
|
||||
if (!tmpw || !tmps) return 2;
|
||||
for (int i = 0; i < 8; i++) tmpw[i] = q8[i];
|
||||
tmps[0] = s8[0]; tmps[1] = s8[1];
|
||||
ColiCudaTensor *tt = nullptr;
|
||||
if (!coli_cuda_tensor_upload(&tt, tmpw, tmps, 1, 4, 2, d0)) return 1;
|
||||
for (int i = 0; i < 8; i++) tmpw[i] = 99;
|
||||
std::free(tmpw); std::free(tmps);
|
||||
if (!coli_cuda_matmul(&tt, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
|
||||
!close_enough(got, want8, 4)) return 1;
|
||||
coli_cuda_tensor_free(tt);
|
||||
}
|
||||
/* Upload failures must be graceful and must not corrupt accounting —
|
||||
* and must not poison LATER healthy launches (sticky-error regression). */
|
||||
{
|
||||
size_t c0 = 0, b0 = 0, c1 = 0, b1 = 0;
|
||||
coli_cuda_stats(-1, &c0, &b0);
|
||||
ColiCudaTensor *bad = nullptr;
|
||||
if (coli_cuda_tensor_upload(&bad, q8, s8, 1, 4, 2, 9999)) return 1;
|
||||
if (coli_cuda_tensor_upload(&bad, q8, s8, 7, 4, 2, d0)) return 1;
|
||||
if (coli_cuda_tensor_upload(&bad, q8, nullptr, 1, 4, 2, d0)) return 1;
|
||||
if (coli_cuda_tensor_upload(&bad, nullptr, s8, 1, 4, 2, d0)) return 1;
|
||||
if (coli_cuda_tensor_upload(&bad, q8, s8, 1, 1 << 20, 1 << 24, d0)) return 1; /* ~16 TB */
|
||||
if (bad) return 1;
|
||||
coli_cuda_stats(-1, &c1, &b1);
|
||||
if (c0 != c1 || b0 != b1) return 1;
|
||||
/* healthy launch immediately after the failed allocation */
|
||||
if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
|
||||
!close_enough(got, want8, 4)) return 1;
|
||||
}
|
||||
/* Fault injection hook: on/off, restores cleanly. */
|
||||
if (setenv("COLI_GPU_FAIL_AFTER", "0", 1)) return 2;
|
||||
if (coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0)) return 1;
|
||||
if (unsetenv("COLI_GPU_FAIL_AFTER")) return 2;
|
||||
if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) ||
|
||||
!close_enough(got, want8, 4)) return 1;
|
||||
const int8_t q8b[8]={-1,-2,-3,-4, 1,-2,3,-4};
|
||||
const float s8b[2]={1.f,.5f},want8b[4]={10.f,15.f,-3.f,-2.5f};
|
||||
if(!coli_cuda_tensor_update(t8,q8b,s8b)||
|
||||
!coli_cuda_matmul(&t8,got,x,q8b,s8b,1,2,4,2,d0)||
|
||||
!coli_cuda_matmul(&t8,got,x,q8b,s8b,1,2,4,2,d0,0)||
|
||||
!close_enough(got,want8b,4))return 1;
|
||||
|
||||
/* Rows [-8,-1,0,7] and [1,2,3,4], packed low nibble first. */
|
||||
@@ -61,26 +107,26 @@ int main(int argc, char **argv) {
|
||||
const float s4[2] = {1.0f, 0.25f};
|
||||
const float want4[2] = {-34.0f, -2.5f};
|
||||
ColiCudaTensor *t4 = nullptr;
|
||||
if (!coli_cuda_matmul(&t4, got, x, q4, s4, 2, 1, 4, 2, d1) || !close_enough(got, want4, 2)) return 1;
|
||||
if (!coli_cuda_matmul(&t4, got, x, q4, s4, 2, 1, 4, 2, d1, 0) || !close_enough(got, want4, 2)) return 1;
|
||||
|
||||
const uint8_t q2[2] = {0xe4, 0x1b};
|
||||
const float s2[2] = {0.5f, 2.0f};
|
||||
const float want2[2] = {-2.0f, 12.0f};
|
||||
ColiCudaTensor *t2 = nullptr;
|
||||
if (!coli_cuda_matmul(&t2, got, x, q2, s2, 3, 1, 4, 2, d1) || !close_enough(got, want2, 2)) return 1;
|
||||
if (!coli_cuda_matmul(&t2, got, x, q2, s2, 3, 1, 4, 2, d1, 0) || !close_enough(got, want2, 2)) return 1;
|
||||
|
||||
const float wf[8] = {1, 0, -1, 2, 0.5f, 0.5f, 0.5f, 0.5f};
|
||||
const float wantf[2] = {-10.0f, -1.0f};
|
||||
ColiCudaTensor *tf = nullptr;
|
||||
if (!coli_cuda_matmul(&tf, got, x, wf, nullptr, 0, 1, 4, 2, d0) || !close_enough(got, wantf, 2)) return 1;
|
||||
if (!coli_cuda_matmul(&tf, got, x, wf, nullptr, 0, 1, 4, 2, d0, 0) || !close_enough(got, wantf, 2)) return 1;
|
||||
|
||||
const float eg[8] = {1,0,0,0, 0,1,0,0};
|
||||
const float eu[8] = {1,0,0,0, 0,1,0,0};
|
||||
const float ed[8] = {1,0, 0,1, 1,1, 1,-1};
|
||||
ColiCudaTensor *tg=nullptr,*tu=nullptr,*td=nullptr;
|
||||
if (!coli_cuda_tensor_upload(&tg,eg,nullptr,0,4,2,d0) ||
|
||||
!coli_cuda_tensor_upload(&tu,eu,nullptr,0,4,2,d0) ||
|
||||
!coli_cuda_tensor_upload(&td,ed,nullptr,0,2,4,d0)) return 1;
|
||||
if (!coli_cuda_tensor_upload(&tg,eg,nullptr,0,4,2,d0,0) ||
|
||||
!coli_cuda_tensor_upload(&tu,eu,nullptr,0,4,2,d0,0) ||
|
||||
!coli_cuda_tensor_upload(&td,ed,nullptr,0,2,4,d0,0)) return 1;
|
||||
float expert[8], want_expert[8];
|
||||
for(int s=0;s<2;s++){
|
||||
float a=x[s*4], b=x[s*4+1];
|
||||
@@ -98,7 +144,7 @@ int main(int argc, char **argv) {
|
||||
const float aw[16]={1,0,0,0, 0,1,0,0, 0,0,1,0, 0,0,0,1};
|
||||
const float aq[4]={1,2,.5f,-.5f},al[12]={1,0,0,0, 0,1,0,0, 0,0,1,0};
|
||||
const float ar[6]={1,0, 0,1, 1,1};float actx[2],aref[2];
|
||||
ColiCudaTensor *at=nullptr;if(!coli_cuda_tensor_upload(&at,aw,nullptr,0,4,4,d0))return 1;
|
||||
ColiCudaTensor *at=nullptr;if(!coli_cuda_tensor_upload(&at,aw,nullptr,0,4,4,d0,0))return 1;
|
||||
float score[3];for(int t=0;t<3;t++)score[t]=aq[0]*al[t*4]+aq[1]*al[t*4+1]+aq[2]*ar[t*2]+aq[3]*ar[t*2+1];
|
||||
float mx=score[0],z=0;for(int t=1;t<3;t++)mx=score[t]>mx?score[t]:mx;
|
||||
for(int t=0;t<3;t++){score[t]=std::exp(score[t]-mx);z+=score[t];}for(int t=0;t<3;t++)score[t]/=z;
|
||||
@@ -117,9 +163,9 @@ int main(int argc, char **argv) {
|
||||
for(int i=0;i<32;i++)ws4[i]=0.01f+(i%5)*0.002f;
|
||||
for(int i=0;i<64;i++)gx4[i]=std::sin((float)(i+1)*0.17f)*2.f;
|
||||
ColiCudaTensor *g4=nullptr,*u4=nullptr,*d4=nullptr;
|
||||
if(!coli_cuda_tensor_upload(&g4,w4,ws4,2,32,32,d0)||
|
||||
!coli_cuda_tensor_upload(&u4,w4,ws4,2,32,32,d0)||
|
||||
!coli_cuda_tensor_upload(&d4,w4,ws4,2,32,32,d0))return 1;
|
||||
if(!coli_cuda_tensor_upload(&g4,w4,ws4,2,32,32,d0,0)||
|
||||
!coli_cuda_tensor_upload(&u4,w4,ws4,2,32,32,d0,0)||
|
||||
!coli_cuda_tensor_upload(&d4,w4,ws4,2,32,32,d0,0))return 1;
|
||||
ColiCudaTensor *gg4[2]={g4,g4},*ug4[2]={u4,u4},*dg4[2]={d4,d4};
|
||||
if(!coli_cuda_expert_group(gg4,ug4,dg4,group_rows,2,scalar4,gx4))return 1;
|
||||
setenv("COLI_CUDA_TC_INT4","1",1);
|
||||
|
||||
@@ -177,6 +177,68 @@ static int run_attn(int S, int pos_base, const char* name){
|
||||
return pass?0:1;
|
||||
}
|
||||
|
||||
// serial r_top8 vs parallel r_top8_par on the ENGINE build's own compiled shaders — the
|
||||
// exact-match contract (same indices, same order, same weights bitwise, same keff)
|
||||
// enforced with memcmp, per adversarial input family. `mode` selects the input
|
||||
// construction; see the inventory at the call sites in main(). E is a parameter (not
|
||||
// hardcoded 256) so the same helper drives both the original E=256 fuzz and the
|
||||
// expert-count-generality cases (E=24 <32-lane-width, E=168 REAP-pruned, E=200
|
||||
// lane-straddling boundary, E=257 out-of-contract auto-serial-fallback proof).
|
||||
static int run_rtop8(int mode, int S, int E, float topp, int normk, float rscale, const char *name) {
|
||||
const int K=8, Ksel=8;
|
||||
std::vector<float> sig((size_t)S*E), bias(E);
|
||||
srand(4242+mode*17+S+E);
|
||||
for (int e=0;e<E;e++) bias[e]=((rand()%2001)-1000)/1000.f;
|
||||
for (int s=0;s<S;s++) for (int e=0;e<E;e++) {
|
||||
float *v=&sig[(size_t)s*E+e];
|
||||
switch (mode) {
|
||||
case 0: *v=(float)(rand()%10000)/10000.f; break; // generic sigmoid-like
|
||||
case 1: *v=0.5f; break; // ALL EQUAL: pure tie-break test
|
||||
case 2: *v=(float)((e/2)%8)/8.f; break; // massed duplicates (paired+cyclic ties)
|
||||
case 3: *v=(e%2)?1e-40f:2e-40f; break; // denormal logits (flush behavior must match)
|
||||
case 4: *v=(float)(rand()%3)/2.f; break; // 3-level ties across the whole row
|
||||
// boundary-forcing: elevate the LAST 4 valid experts (E-4..E-1) to near-max choice
|
||||
// so they are guaranteed in the top-8. For an E whose per-lane block size doesn't
|
||||
// divide E evenly, E-1's lane straddles the E boundary (real indices below E,
|
||||
// sentinel -1e30f at/above E in the SAME ch[] block) -- e.g. E=200: per=ceil(200/
|
||||
// 32)=7, lane 28 owns indices 196..202, of which 196-199 are real and 200-202 are
|
||||
// sentinel. Forcing selection onto 196-199 exercises exactly that lane's per-index
|
||||
// e<E boundary check, rather than hoping random data happens to land there.
|
||||
case 5: *v=(e>=E-4)?1.0f:(float)(rand()%10000)/10000.f; break;
|
||||
default: *v=(float)(rand()%10000)/10000.f; break;
|
||||
}
|
||||
}
|
||||
if (mode==1) for (int e=0;e<E;e++) bias[e]=0.25f; // choice fully tied too
|
||||
if (mode==3) for (int e=0;e<E;e++) bias[e]=(e%3)?3e-40f:-3e-40f; // denormal bias as well
|
||||
if (mode==5) for (int e=E-4;e<E;e++) bias[e]=1.0f; // combined choice = 2.0, max possible
|
||||
std::vector<int> is((size_t)S*K), ip((size_t)S*K); std::vector<float> ws((size_t)S*K), wp((size_t)S*K);
|
||||
std::vector<int> ks(S), kp(S);
|
||||
if (!coli_metal_rtop8(0,sig.data(),bias.data(),S,E,K,Ksel,topp,normk,rscale,is.data(),ws.data(),ks.data()) ||
|
||||
!coli_metal_rtop8(1,sig.data(),bias.data(),S,E,K,Ksel,topp,normk,rscale,ip.data(),wp.data(),kp.data())) {
|
||||
printf(" %-34s FAIL (rtop8 runner returned 0)\n", name); return 1; }
|
||||
int ok = memcmp(is.data(),ip.data(),(size_t)S*K*4)==0 &&
|
||||
memcmp(ws.data(),wp.data(),(size_t)S*K*4)==0 && // bitwise: same ops, same order
|
||||
memcmp(ks.data(),kp.data(),(size_t)S*4)==0;
|
||||
if (mode==5 && ok) {
|
||||
// Don't just trust the input design -- confirm the straddling lane's valid segment
|
||||
// (E-4..E-1) was actually selected, in EVERY row, so this case can't silently
|
||||
// degrade into an unrelated pass if the input construction above ever changes.
|
||||
for (int s=0;s<S;s++) { int seen=0;
|
||||
for (int k=0;k<K;k++) if (ip[(size_t)s*K+k]>=E-4 && ip[(size_t)s*K+k]<E) seen++;
|
||||
if (seen<4) { printf(" %-34s *** boundary segment not exercised (row %d saw %d/4) -- test setup bug\n", name, s, seen); return 1; }
|
||||
}
|
||||
}
|
||||
if (!ok) {
|
||||
printf(" %-34s *** MISMATCH\n", name);
|
||||
for (int s=0;s<S;s++){ printf(" row %d keff %d/%d:",s,ks[s],kp[s]);
|
||||
for(int k=0;k<K;k++) printf(" [%d]%d/%d %.6g/%.6g",k,is[s*K+k],ip[s*K+k],ws[s*K+k],wp[s*K+k]);
|
||||
printf("\n"); }
|
||||
return 1;
|
||||
}
|
||||
printf(" %-34s ok (serial==parallel bitwise, S=%d E=%d)\n", name, S, E);
|
||||
return 0;
|
||||
}
|
||||
|
||||
int main(void) {
|
||||
if (!coli_metal_init()) { printf("Metal unavailable (skipping)\n"); return 0; }
|
||||
printf("Metal backend kernel tests:\n");
|
||||
@@ -215,6 +277,32 @@ int main(void) {
|
||||
fail |= run_attn(1, 37, "attn S=1 pos=37");
|
||||
fail |= run_attn(4, 12, "attn S=4 pos=12 (MTP)");
|
||||
fail |= run_attn(3, 0, "attn S=3 pos=0");
|
||||
printf("Metal top-8 select serial-vs-parallel tests (exact-match contract, E=256):\n");
|
||||
fail |= run_rtop8(0, 1, 256, 0.0f, 1, 1.0f, "top8 generic S=1");
|
||||
fail |= run_rtop8(0, 4, 256, 0.0f, 1, 1.0f, "top8 generic S=4");
|
||||
fail |= run_rtop8(1, 1, 256, 0.0f, 1, 1.0f, "top8 ALL-EQUAL ties");
|
||||
fail |= run_rtop8(2, 4, 256, 0.0f, 1, 1.0f, "top8 massed dup ties S=4");
|
||||
fail |= run_rtop8(4, 2, 256, 0.0f, 0, 2.5f, "top8 3-level ties rscale");
|
||||
fail |= run_rtop8(3, 1, 256, 0.0f, 1, 1.0f, "top8 denormal logits");
|
||||
fail |= run_rtop8(0, 1, 256, 0.01f, 1, 1.0f, "top8 topp=0.01 (Ke=1 edge)");
|
||||
fail |= run_rtop8(2, 1, 256, 0.6f, 1, 1.0f, "top8 topp=0.6 tied weights");
|
||||
fail |= run_rtop8(0, 4, 256, 0.999f,1, 1.75f, "top8 topp=0.999 S=4");
|
||||
fail |= run_rtop8(1, 2, 256, 0.5f, 0, 1.0f, "top8 topp on ALL-EQUAL");
|
||||
printf("Metal top-8 select expert-count-generality tests (E!=256, REAP/#428 motivated):\n");
|
||||
fail |= run_rtop8(0, 1, 168, 0.0f, 1, 1.0f, "top8 E=168 (REAP) generic S=1");
|
||||
fail |= run_rtop8(2, 4, 168, 0.0f, 1, 1.0f, "top8 E=168 (REAP) massed dup ties S=4");
|
||||
fail |= run_rtop8(0, 1, 24, 0.0f, 1, 1.0f, "top8 E=24 (<32 lane width) generic");
|
||||
fail |= run_rtop8(1, 1, 24, 0.0f, 1, 1.0f, "top8 E=24 (<32 lane width) ALL-EQUAL ties");
|
||||
// E=200: per-lane block size ceil(200/32)=7, and 200 is NOT a multiple of 7, so lane 28
|
||||
// (indices 196..202) straddles the boundary -- 196-199 real, 200-202 sentinel -1e30f in
|
||||
// the SAME ch[] block. E=24 and E=168 above both happen to divide evenly by their own
|
||||
// per (24/1, 168/6), so no case before this one exercised a lane whose ch[] mixes real
|
||||
// and sentinel indices. mode 5 deterministically forces indices 196-199 into the top-8
|
||||
// (see run_rtop8) and asserts they were actually selected, rather than hoping random
|
||||
// data lands there -- proving by TEST what the per-index `e<E` check was proven by
|
||||
// reading (both kernels agree bitwise on a selection that requires that check to fire).
|
||||
fail |= run_rtop8(5, 4, 200, 0.0f, 1, 1.0f, "top8 E=200 (lane straddles E boundary)");
|
||||
fail |= run_rtop8(0, 1, 257, 0.0f, 1, 1.0f, "top8 E=257 (>256, auto-serial-fallback)");
|
||||
printf(fail? "metal backend tests: FAILED\n" : "metal backend tests: ok\n");
|
||||
coli_metal_shutdown();
|
||||
return fail;
|
||||
|
||||
@@ -44,6 +44,12 @@ static void test_submit_header(void)
|
||||
assert(coli_submit_parse("SUBMIT 1 0 16777216 3 1 1", &sub));
|
||||
assert(!coli_submit_parse("SUBMIT 1 0 16777217 3 1 1", &sub));
|
||||
assert(!coli_submit_parse("SUBMIT 1 0 2 3 1 1 trailing", &sub));
|
||||
/* optional 7th field: per-request grammar length (0 when absent) */
|
||||
assert(coli_submit_parse("SUBMIT 42 3 17 64 0.7 0.95", &sub) && sub.gbytes == 0);
|
||||
assert(coli_submit_parse("SUBMIT 42 3 17 64 0.7 0.95 512", &sub) && sub.gbytes == 512);
|
||||
assert(coli_submit_parse("SUBMIT 42 3 17 64 0.7 0.95 1048576", &sub));
|
||||
assert(!coli_submit_parse("SUBMIT 42 3 17 64 0.7 0.95 1048577", &sub));
|
||||
assert(!coli_submit_parse("SUBMIT 42 3 17 64 0.7 0.95 512 extra", &sub));
|
||||
}
|
||||
|
||||
int main(void)
|
||||
|
||||
@@ -0,0 +1,315 @@
|
||||
#!/usr/bin/env python3
|
||||
"""Exhaustive optimization dossier for a colibri engine run.
|
||||
|
||||
This is NOT a pass/fail test. It runs the engine with every instrumentation flag
|
||||
on (PROF, COLI_CUDA_PROFILE, CACHE_ROUTE, DISK_SPLIT, LOOKA) and prints a section-
|
||||
by-section report answering, for each subsystem:
|
||||
|
||||
WHAT is doing it — which phase/kernel/tier
|
||||
WHEN it is doing it — how much of decode wall-time it owns
|
||||
WITH WHAT — the config/weights/tier it used
|
||||
IS IT INEFFICIENT? — a verdict, with the threshold
|
||||
HOW TO IMPROVE — the concrete knob, named
|
||||
|
||||
Activation (opt-in only — NOT in `make test`):
|
||||
COLI_EFFICIENCY_MODEL=<model_dir> python tests/test_efficiency_report.py
|
||||
|
||||
Optional env:
|
||||
COLI_EFFICIENCY_CUDA=1 also exercise the CUDA dense/expert tiers
|
||||
COLI_EFFICIENCY_NGEN=N decode tokens (default 24)
|
||||
COLI_EFFICIENCY_RAM_GB=N RAM budget (default 28)
|
||||
COLI_EFFICIENCY_VRAM_GB=N CUDA expert-tier budget GB (default 4)
|
||||
COLI_EFFICIENCY_PROMPT=... prompt (default: a code-gen prompt)
|
||||
|
||||
Exit code is always 0 (it's a dossier, not a gate). Lines marked FLAG point at
|
||||
the most likely lever to move tok/s for the observed bottleneck.
|
||||
"""
|
||||
import os
|
||||
import sys
|
||||
import time
|
||||
from pathlib import Path
|
||||
|
||||
sys.path.insert(0, str(Path(__file__).resolve().parent.parent))
|
||||
from tools.efficiency import run_engine, disk_wait_share # noqa: E402
|
||||
|
||||
|
||||
# --- advisory thresholds (the "IS IT INEFFICIENT?" lines) ---
|
||||
DISK_WAIT_DOMINANT = 0.40 # >40% decode waiting on expert reads -> I/O-bound
|
||||
LOW_HIT_RATE = 0.30 # <30% cache hit -> thrashing (cap too small)
|
||||
LOW_ROUTE_AGREE = 0.80 # <80% routing overlap -> prefetch is guessing wrong
|
||||
HIGH_TAIL_RATIO = 3.0 # p99 > 3x p50 -> decode stalls (I/O hiccups / KV grow)
|
||||
LOW_MTP_ACCEPT = 0.20 # <20% MTP acceptance -> draft decoder is dead weight
|
||||
VRAM_WASTE_CALLS = 0 # experts pinned in VRAM but 0 calls served
|
||||
|
||||
|
||||
def _flag(ok): return "OK " if ok else "FLAG"
|
||||
|
||||
|
||||
def _bar(frac, width=24):
|
||||
"""A simple ASCII bar for share visualization."""
|
||||
n = max(0, min(width, round(frac * width)))
|
||||
return "#" * n + "." * (width - n)
|
||||
|
||||
|
||||
def _line(label, value, flag=None, note=""):
|
||||
tag = f" [{flag}]" if flag else ""
|
||||
print(f" {label:<22} {value}{tag} {note}" if note else f" {label:<22} {value}{tag}")
|
||||
|
||||
|
||||
def main() -> int:
|
||||
model = os.environ.get("COLI_EFFICIENCY_MODEL")
|
||||
if not model:
|
||||
print(__doc__)
|
||||
print("\nNot activated: set COLI_EFFICIENCY_MODEL=<model_dir> to run.")
|
||||
return 0
|
||||
model = str(Path(model).resolve())
|
||||
if not Path(model).is_dir():
|
||||
print(f"ERROR: {model} is not a directory", file=sys.stderr)
|
||||
return 0
|
||||
|
||||
ngen = int(os.environ.get("COLI_EFFICIENCY_NGEN", "24"))
|
||||
ram_gb = os.environ.get("COLI_EFFICIENCY_RAM_GB", "28")
|
||||
vram_gb = os.environ.get("COLI_EFFICIENCY_VRAM_GB", "4")
|
||||
prompt = os.environ.get(
|
||||
"COLI_EFFICIENCY_PROMPT",
|
||||
"Write a Python function that computes the factorial of a number. "
|
||||
"Include error handling and a docstring.")
|
||||
use_cuda = os.environ.get("COLI_EFFICIENCY_CUDA") == "1"
|
||||
|
||||
# Turn ON every instrumentation flag so the dossier has maximum detail.
|
||||
# These are all observability toggles (PROF/COLI_CUDA_PROFILE/CACHE_ROUTE/
|
||||
# DISK_SPLIT/LOOKA); none change the computed output.
|
||||
overlay = dict(
|
||||
NGEN=str(ngen), TEMP="0", RAM_GB=ram_gb, PROMPT=prompt,
|
||||
PROF="1", CACHE_ROUTE="1", DISK_SPLIT="1", LOOKA="1", ROUTE_AGREE="1",
|
||||
)
|
||||
if use_cuda:
|
||||
overlay.update(COLI_CUDA="1", COLI_GPU="0", CUDA_DENSE="1",
|
||||
COLI_CUDA_PROFILE="1", CUDA_EXPERT_GB=vram_gb)
|
||||
|
||||
print("=" * 78)
|
||||
print(f"OPTIMIZATION DOSSIER — {Path(model).name}")
|
||||
print(f" mode : {'CUDA (dense+expert tiers)' if use_cuda else 'CPU-only'} "
|
||||
f"ngen : {ngen} ram : {ram_gb} GB" +
|
||||
(f" vram : {vram_gb} GB" if use_cuda else ""))
|
||||
print("=" * 78)
|
||||
|
||||
t0 = time.time()
|
||||
t, proc = run_engine(overlay, snap=model, timeout=3600.0)
|
||||
wall = time.time() - t0
|
||||
flags = [] # collected FLAG lines for the summary
|
||||
|
||||
print(f"\n[0] RUN")
|
||||
_line("wall clock", f"{wall:.0f}s")
|
||||
_line("exit code", proc.returncode,
|
||||
None if proc.returncode == 0 else "FLAG",
|
||||
"" if proc.returncode == 0 else "non-zero exit")
|
||||
if proc.returncode != 0:
|
||||
print(" stderr tail:")
|
||||
for ln in proc.stderr.strip().splitlines()[-8:]:
|
||||
print(f" {ln}")
|
||||
return 0
|
||||
|
||||
# ---------------------------------------------------------------- [1] WHO ----
|
||||
print(f"\n[1] PROVENANCE — what is running, on what, with what config")
|
||||
if t.get("machine"):
|
||||
m = t["machine"]
|
||||
_line("CPU", m["cpu"])
|
||||
_line("cores / omp", f"{m['cores']} cores")
|
||||
_line("backend", m["backend"])
|
||||
if t.get("load"):
|
||||
ld = t["load"]
|
||||
_line("model load time", f"{ld['load_s']:.2f}s")
|
||||
_line("resident dense", f"{ld['resident_dense_mb']:.1f} MB")
|
||||
_line("layers / experts", f"{ld['layers']} layers, {ld['experts']} experts")
|
||||
_line("MTP", f"{ld['mtp_status']} (draft={ld['draft']})")
|
||||
if t.get("config_str"):
|
||||
_line("resolved config", t["config_str"])
|
||||
print(" (this is the EFFECTIVE config after auto-budgeting — not your env verbatim)")
|
||||
|
||||
# ---------------------------------------------------------------- [2] SPEED --
|
||||
print(f"\n[2] THROUGHPUT — is it fast, is the tail healthy")
|
||||
if t.get("tok_s") is not None:
|
||||
_line("tok/s", f"{t['tok_s']:.3f}")
|
||||
else:
|
||||
flags.append("throughput line missing — engine output format may have changed")
|
||||
if t.get("latency"):
|
||||
la = t["latency"]
|
||||
_line("decode forwards", f"{int(la['forwards'])}")
|
||||
_line("p50 / p90", f"{la['p50_ms']:.1f} / {la['p90_ms']:.1f} ms")
|
||||
_line("p99 / max", f"{la['p99_ms']:.1f} / {la['max_ms']:.1f} ms")
|
||||
tail_ok = la["p99_ms"] <= HIGH_TAIL_RATIO * la["p50_ms"]
|
||||
_line("tail ratio (p99/p50)", f"{la['p99_ms']/max(la['p50_ms'],1e-9):.2f}x",
|
||||
_flag(tail_ok),
|
||||
"high tail = decode stalls (I/O hiccups, KV growth, re-pin)")
|
||||
if not tail_ok:
|
||||
flags.append(f"tail latency p99={la['p99_ms']:.1f}ms >> p50={la['p50_ms']:.1f}ms "
|
||||
"(look for REPIN swaps or disk stalls)")
|
||||
|
||||
# ---------------------------------------------------------------- [3] TIME ---
|
||||
print(f"\n[3] WHERE TIME GOES — what is doing it, when (share of decode)")
|
||||
ts = t.get("time_shares")
|
||||
prof = t.get("profile")
|
||||
if ts:
|
||||
order = [("io", "expert-disk I/O", DISK_WAIT_DOMINANT, "the cache is too small / disk is slow"),
|
||||
("matmul", "expert matmul", 0.40, "compute-bound; more cores or a GPU expert tier"),
|
||||
("attention", "attention", 0.35, "context length is the cost; lower CTX"),
|
||||
("head", "lm_head", 0.10, "vocab projection; unusual to dominate"),
|
||||
("other", "other", 0.30, "scheduling / KV bookkeeping overhead")]
|
||||
for key, name, thresh, lever in order:
|
||||
f = ts[key]
|
||||
ok = f < thresh
|
||||
_line(name, f"{f:5.0%} {_bar(f)}", _flag(ok),
|
||||
"" if ok else f"->{lever}")
|
||||
if not ok:
|
||||
flags.append(f"{name} dominates ({f:.0%}) -> {lever}")
|
||||
if t.get("verdict"):
|
||||
print(f" engine verdict : {t['verdict']}")
|
||||
elif prof:
|
||||
print(" (no [PROF] time shares — set PROF=1 for phase percentages)")
|
||||
if prof:
|
||||
print(" absolute seconds :")
|
||||
for k in ("disk", "expert_matmul", "attention", "lm_head", "other"):
|
||||
_line(k, f"{prof[k]:.3f}s")
|
||||
|
||||
# attention sub-breakdown: how is attention being read
|
||||
ab = t.get("attn_breakdown")
|
||||
if ab:
|
||||
print(f"\n[3a] ATTENTION BREAKDOWN — how the attention phase is spent")
|
||||
atot = sum(ab.values()) or 1.0
|
||||
for k, label in (("proj_rope", "projection + RoPE"),
|
||||
("score_sm_value", "score-softmax-value"),
|
||||
("out_proj", "output projection")):
|
||||
_line(label, f"{ab[k]:.3f}s ({ab[k]/atot:.0%} of attn)")
|
||||
|
||||
# ---------------------------------------------------------------- [4] CACHE --
|
||||
print(f"\n[4] EXPERT CACHE — is the cache efficient")
|
||||
hit = t.get("hit_pct")
|
||||
if hit is not None:
|
||||
ok = hit >= LOW_HIT_RATE * 100
|
||||
_line("hit rate", f"{hit:.1f}%", _flag(ok),
|
||||
"" if ok else "<30% = thrashing; raise RAM_GB or cap")
|
||||
if not ok:
|
||||
flags.append(f"cache hit {hit:.1f}% is low -> raise RAM_GB (or cap), add PIN_GB")
|
||||
el = t.get("experts_loaded")
|
||||
if el:
|
||||
per_tok = el["per_tok"]
|
||||
_line("experts loaded/token", f"{per_tok:.1f}")
|
||||
_line(" per-layer", f"{el['per_layer']:.2f} across {el['n_sparse_layers']} sparse layers")
|
||||
_line(" baseline", f"topk={el['baseline_topk']} active experts/token")
|
||||
base_topk = el["baseline_topk"]
|
||||
if base_topk > 0 and per_tok > 2 * base_topk:
|
||||
flags.append(f"loading {per_tok:.0f} experts/token vs topk={base_topk} "
|
||||
"-> redundant I/O; cache is re-fetching evicted experts")
|
||||
|
||||
# ---------------------------------------------------------------- [5] DISK ---
|
||||
print(f"\n[5] DISK I/O — is I/O the bottleneck, and where")
|
||||
eio = t.get("expert_io")
|
||||
if eio:
|
||||
_line("total fetched", f"{eio['gb_fetched']:.3f} GB")
|
||||
_line("per token", f"{eio['mb_per_tok']:.1f} MB/token")
|
||||
_line("disk throughput", f"{eio['gb_per_s']:.2f} GB/s over the run")
|
||||
_line("read service", f"{eio['read_service_s']:.2f}s (on I/O threads)")
|
||||
_line("felt wait", f"{eio['felt_wait_s']:.2f}s (stall compute felt)")
|
||||
if eio["felt_wait_s"] > eio["read_service_s"] * 0.5 and eio["read_service_s"] > 0:
|
||||
flags.append("felt wait is a large fraction of read service -> PIPE=1 may not be "
|
||||
"overlapping fully, or DIRECT=1 on NVMe")
|
||||
ds = t.get("disk_split")
|
||||
if ds:
|
||||
print(f"\n[5a] DISK-LOAD SPLIT — which decode phase reads the bytes")
|
||||
_line("draft phase", f"{ds['draft']} loads")
|
||||
_line("absorb phase", f"{ds['absorb']} loads")
|
||||
_line("verify/main", f"{ds['verify_main']} loads")
|
||||
_line("MTP-layer bytes", f"{ds['mtp_loads']} loads, {ds['mtp_gb']:.2f} GB")
|
||||
_line("main-layer bytes", f"{ds['main_loads']} loads, {ds['main_gb']:.2f} GB")
|
||||
if ds.get("mtp_bytes_pct") is not None:
|
||||
_line("MTP share of bytes", f"{ds['mtp_bytes_pct']:.1f}%")
|
||||
share = disk_wait_share(t)
|
||||
if share is not None:
|
||||
ok = share < DISK_WAIT_DOMINANT
|
||||
_line("disk-wait share", f"{share:.0%}", _flag(ok),
|
||||
"" if ok else "I/O-bound (see levers in [3])")
|
||||
|
||||
# ---------------------------------------------------------------- [6] ROUTE --
|
||||
print(f"\n[6] ROUTING QUALITY — is the router / prefetch accurate")
|
||||
ra = t.get("route_agree")
|
||||
if ra:
|
||||
ok = ra["agree_pct"] >= LOW_ROUTE_AGREE * 100
|
||||
_line("route_agree", f"{ra['agree_pct']:.1f}% overlap with true top-K",
|
||||
_flag(ok),
|
||||
"" if ok else "prefetch is guessing wrong; CACHE_ROUTE params may need tuning")
|
||||
_line("route_kl", f"{ra['kl']:.4f} mean KL (lower = closer to true routing)")
|
||||
if not ok:
|
||||
flags.append(f"route_agree {ra['agree_pct']:.1f}% low -> tune ROUTE_J/M/P, "
|
||||
"or prefetch is hurting more than helping")
|
||||
sw = t.get("swap")
|
||||
if sw:
|
||||
_line("cache swaps", f"{sw['swaps']}/{sw['slots']} ({sw['pct']:.1f}%)",
|
||||
None, "high swap = churn between turns")
|
||||
la = t.get("lookahead")
|
||||
if la:
|
||||
print(f"\n[6a] ROUTING PREDICTABILITY — recall of true experts in predicted top-8")
|
||||
print(" (which predictor should drive prefetch? highest recall wins)")
|
||||
for row in la:
|
||||
_line(row["predictor"][:34], f"{row['pct']:5.1f}% ({row['hit']}/{row['tot']})")
|
||||
|
||||
# ---------------------------------------------------------------- [7] SPEC ---
|
||||
print(f"\n[7] SPECULATION — is the draft decoder pulling weight")
|
||||
sp = t.get("speculation")
|
||||
if sp:
|
||||
_line("tokens/forward", f"{sp['tok_per_fw']:.2f} (>1.0 means speculation helps)")
|
||||
_line("forwards/tokens", f"{sp['forwards']} forwards for {sp['tokens']} tokens")
|
||||
# Speculation helps only if acceptance is high enough that tok/forward > 1.
|
||||
# tok_per_fw already == 1.0 when nothing verifies, so judge by acceptance.
|
||||
ok = sp["mtp_accept_pct"] >= LOW_MTP_ACCEPT * 100
|
||||
_line("MTP acceptance", f"{sp['mtp_accept_pct']:.0f}%", _flag(ok),
|
||||
"" if ok else "<20% -> drafts rarely verify; DRAFT=0 may be faster")
|
||||
if not ok:
|
||||
flags.append(f"MTP acceptance {sp['mtp_accept_pct']:.0f}% low -> "
|
||||
"drafts cost more I/O than they save; try DRAFT=0")
|
||||
|
||||
# ---------------------------------------------------------------- [8] GPU ----
|
||||
print(f"\n[8] GPU TIERS — is the GPU actually used")
|
||||
cuda = t.get("cuda") or {}
|
||||
if not cuda.get("enabled"):
|
||||
print(" (CUDA not enabled — CPU-only run)")
|
||||
else:
|
||||
if cuda.get("resident_tensors") is not None:
|
||||
_line("resident dense tensors", f"{cuda['resident_tensors']} tensors, "
|
||||
f"{cuda['resident_gb']:.2f} GB")
|
||||
if cuda.get("expert_count") is not None:
|
||||
waste = cuda["calls_served"] <= VRAM_WASTE_CALLS
|
||||
_line("expert tier", f"{cuda['expert_count']} experts pinned "
|
||||
f"({cuda['expert_gb']:.2f} GB)", _flag(not waste))
|
||||
_line(" calls served", f"{cuda['calls_served']} from VRAM",
|
||||
_flag(not waste),
|
||||
"" if not waste else "WASTE: pinned but never routed -> lower CUDA_EXPERT_GB")
|
||||
if waste:
|
||||
flags.append("VRAM expert tier has 0 calls served -> experts pinned but unused; "
|
||||
"PIN stats may not match this workload")
|
||||
if cuda.get("groups"):
|
||||
g = cuda["groups"]
|
||||
_line("expert groups", f"{g['calls']} calls, {g['experts']} experts, "
|
||||
f"{g['rows']} rows ({g['experts_per_call']:.1f} experts/call)")
|
||||
if cuda.get("groups_timing"):
|
||||
gt = cuda["groups_timing"]
|
||||
_line("GPU timing", f"H2D {gt['h2d_ms']:.1f} ms | kernel {gt['kernel_ms']:.1f} ms | "
|
||||
f"D2H {gt['d2h_ms']:.1f} ms")
|
||||
if gt["h2d_ms"] + gt["d2h_ms"] > gt["kernel_ms"]:
|
||||
flags.append("CUDA H2D+D2H > kernel time -> transfer-bound; "
|
||||
"consider larger expert tier to keep weights resident")
|
||||
|
||||
# ---------------------------------------------------------------- summary ----
|
||||
print("\n" + "=" * 78)
|
||||
if flags:
|
||||
print(f" {len(flags)} FLAG(s) — the most likely levers to move tok/s:")
|
||||
for f in flags:
|
||||
print(f" - {f}")
|
||||
else:
|
||||
print(" no flags — every measured subsystem is within advisory thresholds.")
|
||||
print("=" * 78)
|
||||
return 0
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
raise SystemExit(main())
|
||||
@@ -32,9 +32,16 @@ def args(**over):
|
||||
|
||||
|
||||
class EnvDefaultsTest(unittest.TestCase):
|
||||
def env_for_with(self, environ, platform):
|
||||
def env_for_with(self, environ, platform, cuda=False):
|
||||
"""Run env_for on a bare-chat args() under a faked env + platform.
|
||||
|
||||
cuda=False by default so the existing default-I/O tests stay
|
||||
deterministic: the Windows auto-enable branch calls cuda_binary() and
|
||||
(if True) discover_gpus(), both of which reach the real machine — faking
|
||||
False keeps these tests independent of the host's GPU."""
|
||||
with mock.patch.dict(os.environ, environ, clear=True), \
|
||||
mock.patch.object(sys, "platform", platform):
|
||||
mock.patch.object(sys, "platform", platform), \
|
||||
mock.patch.object(coli, "cuda_binary", return_value=cuda):
|
||||
return coli.env_for(args())
|
||||
|
||||
def test_win32_sets_measured_defaults(self):
|
||||
@@ -64,5 +71,80 @@ class EnvDefaultsTest(unittest.TestCase):
|
||||
self.assertNotIn(k, e)
|
||||
|
||||
|
||||
class CudaAutoEnableTest(unittest.TestCase):
|
||||
"""Windows bare `coli chat` (no --gpu/--vram/--auto-tier) used to ALWAYS run
|
||||
CPU-only even on a CUDA build with a GPU present. env_for now auto-enables
|
||||
CUDA on win32 when cuda_binary() is True and a GPU is discoverable; falls
|
||||
back to CPU with a warning if nvidia-smi (discover_gpus) is missing; stays
|
||||
silent on a CPU build; and never touches the Linux path."""
|
||||
|
||||
def _env_for(self, platform, cuda, gpus, plan=None):
|
||||
# Patch discover_gpus / build_plan / environment_for_plan at the
|
||||
# resource_plan module (env_for imports them lazily on each call, so the
|
||||
# patches are live when those imports run). Stubbing the planner keeps
|
||||
# the test independent of a real model dir (args().model == "X").
|
||||
import resource_plan
|
||||
a = args()
|
||||
GPB = 1024 ** 3
|
||||
if plan is None:
|
||||
plan = {"tiers": {"ram": {"budget_bytes": 16 * GPB, "cache_slots_per_layer": 4},
|
||||
"vram": {"budget_bytes": int(8.0 * GPB), "devices": gpus}}}
|
||||
|
||||
def fake_environment_for_plan(p, env, cuda_enabled=True):
|
||||
# Mirror the real contract: size CUDA_EXPERT_GB from the plan's VRAM
|
||||
# budget (this is the value env_for propagates into the engine env).
|
||||
r = dict(env)
|
||||
if cuda_enabled and p["tiers"]["vram"]["devices"] and p["tiers"]["vram"]["budget_bytes"] > 0:
|
||||
r["CUDA_EXPERT_GB"] = f"{p['tiers']['vram']['budget_bytes'] / GPB:.3f}"
|
||||
return r
|
||||
|
||||
with mock.patch.dict(os.environ, {}, clear=True), \
|
||||
mock.patch.object(sys, "platform", platform), \
|
||||
mock.patch.object(coli, "cuda_binary", return_value=cuda), \
|
||||
mock.patch.object(resource_plan, "discover_gpus", return_value=gpus), \
|
||||
mock.patch.object(resource_plan, "build_plan", return_value=plan), \
|
||||
mock.patch.object(resource_plan, "environment_for_plan",
|
||||
side_effect=fake_environment_for_plan):
|
||||
return coli.env_for(a)
|
||||
|
||||
def _fake_gpu(self, index=0, name="NVIDIA GeForce RTX 5070 Ti",
|
||||
total_mib=16384, free_mib=15000):
|
||||
return {"index": index, "name": name,
|
||||
"total_bytes": total_mib * 1024 * 1024,
|
||||
"free_bytes": free_mib * 1024 * 1024}
|
||||
|
||||
def test_win32_auto_enables_cuda_when_gpu_present(self):
|
||||
e = self._env_for("win32", cuda=True, gpus=[self._fake_gpu()])
|
||||
self.assertEqual(e["COLI_CUDA"], "1")
|
||||
self.assertEqual(e["COLI_GPUS"], "0")
|
||||
# VRAM budget is sized from free VRAM by build_plan (real minus reserve),
|
||||
# so it must be present and positive — never a guess or zero.
|
||||
self.assertIn("CUDA_EXPERT_GB", e)
|
||||
self.assertGreater(float(e["CUDA_EXPERT_GB"]), 0.0)
|
||||
# Dense offload is an explicit opt-in (matches --auto-tier): not set here.
|
||||
self.assertNotIn("CUDA_DENSE", e)
|
||||
|
||||
def test_win32_falls_back_to_cpu_when_nvidia_smi_missing(self):
|
||||
# coli_cuda.dll present (cuda=True) but nvidia-smi absent (no GPUs found)
|
||||
# -> warn + CPU-only, never crash, never set COLI_CUDA.
|
||||
e = self._env_for("win32", cuda=True, gpus=[])
|
||||
self.assertNotIn("COLI_CUDA", e)
|
||||
self.assertNotIn("COLI_GPUS", e)
|
||||
self.assertNotIn("CUDA_EXPERT_GB", e)
|
||||
|
||||
def test_win32_cpu_build_stays_silent(self):
|
||||
# No coli_cuda.dll (cuda=False) -> CPU build, nothing GPU-related emitted.
|
||||
e = self._env_for("win32", cuda=False, gpus=[self._fake_gpu()])
|
||||
self.assertNotIn("COLI_CUDA", e)
|
||||
self.assertNotIn("COLI_GPUS", e)
|
||||
|
||||
def test_linux_bare_chat_not_auto_enabled(self):
|
||||
# The auto-enable is scoped to win32: a Linux bare chat with a GPU
|
||||
# present must NOT turn CUDA on (Linux keeps the explicit-flag UX).
|
||||
e = self._env_for("linux", cuda=True, gpus=[self._fake_gpu()])
|
||||
self.assertNotIn("COLI_CUDA", e)
|
||||
self.assertNotIn("CUDA_EXPERT_GB", e)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
unittest.main()
|
||||
|
||||
@@ -0,0 +1,108 @@
|
||||
/* Grouped-int4 (fmt=4) CUDA kernel oracle (#334).
|
||||
*
|
||||
* Feeds random offset-binary nibble weights + [O, ng] group scales through
|
||||
* grouped_hidden_g4_dual / grouped_down_g4 and checks against a CPU reference
|
||||
* that replicates matmul_i4_grouped's semantics (value = nibble - 8, per-group
|
||||
* partial dot x scale). Covers gs=64, a non-divisible tail group, and a
|
||||
* per-row (gs=0) member riding in the same launch — the fmt=2-compat case.
|
||||
*
|
||||
* The device buffers get the same XOR 0x88 offset->signed conversion the
|
||||
* upload path applies, so the kernels are exercised exactly as deployed.
|
||||
*
|
||||
* Build: nvcc -O2 -std=c++17 -arch=native tests/test_grouped_g4_cuda.cu -o tests/test_grouped_g4
|
||||
*/
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <cmath>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include "../backend_cuda.cu"
|
||||
|
||||
static void cpu_gemv_g4(const uint8_t *q,const float *sc,int K,int O,int gs,
|
||||
const float *x,float *y){
|
||||
int rb=(K+1)/2, ng=gs>0?(K+gs-1)/gs:1, egs=gs>0?gs:K;
|
||||
for(int o=0;o<O;o++){
|
||||
const uint8_t *row=q+(size_t)o*rb; const float *scl=sc+(size_t)o*ng;
|
||||
double a=0;
|
||||
for(int g=0; g*egs<K; g++){
|
||||
int base=g*egs, glen=egs; if(base+glen>K) glen=K-base;
|
||||
double p=0;
|
||||
for(int i=base;i<base+glen;i++){
|
||||
uint8_t v=row[i>>1]; int n=(i&1)?(v>>4):(v&15);
|
||||
p+=(double)x[i]*(n-8);
|
||||
}
|
||||
a+=p*scl[g];
|
||||
}
|
||||
y[o]=(float)a;
|
||||
}
|
||||
}
|
||||
|
||||
int main(void){
|
||||
srand(7);
|
||||
const int D=200, I=96, gs=64; /* tail group: 200 % 64 = 8 */
|
||||
const int COUNT=3; /* expert 0,1: fmt4 gs=64; expert 2: per-row (gs=0) */
|
||||
const int rbD=(D+1)/2, rbI=(I+1)/2;
|
||||
const int ngD=(D+gs-1)/gs, ngI=(I+gs-1)/gs;
|
||||
int trials=50, bad=0;
|
||||
for(int t=0;t<trials;t++){
|
||||
GroupDesc host[COUNT]; float *xs; cudaMallocManaged(&xs,(size_t)COUNT*D*4);
|
||||
float *gate,*up,*y; cudaMallocManaged(&gate,(size_t)COUNT*I*4);
|
||||
cudaMallocManaged(&up,(size_t)COUNT*I*4); cudaMallocManaged(&y,(size_t)COUNT*D*4);
|
||||
uint8_t *qg[COUNT],*qu[COUNT],*qd[COUNT]; float *sg[COUNT],*su[COUNT],*sd[COUNT];
|
||||
uint8_t *hg[COUNT],*hu[COUNT],*hd[COUNT]; float *hgs[COUNT],*hus[COUNT],*hds[COUNT];
|
||||
for(int c=0;c<COUNT;c++){
|
||||
int cgs = c==2 ? 0 : gs;
|
||||
int cngD = cgs? ngD:1, cngI = cgs? ngI:1;
|
||||
hg[c]=(uint8_t*)malloc((size_t)I*rbD); hu[c]=(uint8_t*)malloc((size_t)I*rbD);
|
||||
hd[c]=(uint8_t*)malloc((size_t)D*rbI);
|
||||
hgs[c]=(float*)malloc((size_t)I*cngD*4); hus[c]=(float*)malloc((size_t)I*cngD*4);
|
||||
hds[c]=(float*)malloc((size_t)D*cngI*4);
|
||||
for(size_t i=0;i<(size_t)I*rbD;i++){ hg[c][i]=rand()&255; hu[c][i]=rand()&255; }
|
||||
for(size_t i=0;i<(size_t)D*rbI;i++) hd[c][i]=rand()&255;
|
||||
for(size_t i=0;i<(size_t)I*cngD;i++){ hgs[c][i]=.01f+.05f*(rand()/(float)RAND_MAX);
|
||||
hus[c][i]=.01f+.05f*(rand()/(float)RAND_MAX); }
|
||||
for(size_t i=0;i<(size_t)D*cngI;i++) hds[c][i]=.01f+.05f*(rand()/(float)RAND_MAX);
|
||||
cudaMalloc(&qg[c],(size_t)I*rbD); cudaMalloc(&qu[c],(size_t)I*rbD); cudaMalloc(&qd[c],(size_t)D*rbI);
|
||||
cudaMalloc(&sg[c],(size_t)I*cngD*4); cudaMalloc(&su[c],(size_t)I*cngD*4); cudaMalloc(&sd[c],(size_t)D*cngI*4);
|
||||
cudaMemcpy(qg[c],hg[c],(size_t)I*rbD,cudaMemcpyHostToDevice);
|
||||
cudaMemcpy(qu[c],hu[c],(size_t)I*rbD,cudaMemcpyHostToDevice);
|
||||
cudaMemcpy(qd[c],hd[c],(size_t)D*rbI,cudaMemcpyHostToDevice);
|
||||
offset_to_signed_s4<<<64,256>>>(qg[c],(size_t)I*rbD);
|
||||
offset_to_signed_s4<<<64,256>>>(qu[c],(size_t)I*rbD);
|
||||
offset_to_signed_s4<<<64,256>>>(qd[c],(size_t)D*rbI);
|
||||
cudaMemcpy(sg[c],hgs[c],(size_t)I*cngD*4,cudaMemcpyHostToDevice);
|
||||
cudaMemcpy(su[c],hus[c],(size_t)I*cngD*4,cudaMemcpyHostToDevice);
|
||||
cudaMemcpy(sd[c],hds[c],(size_t)D*cngI*4,cudaMemcpyHostToDevice);
|
||||
host[c]={qg[c],qu[c],qd[c],sg[c],su[c],sd[c],4,4,4,1,c,cgs,cgs,cgs};
|
||||
}
|
||||
for(size_t i=0;i<(size_t)COUNT*D;i++) xs[i]=(rand()/(float)RAND_MAX-.5f)*2.f;
|
||||
GroupDesc *ddesc; cudaMalloc(&ddesc,sizeof(host));
|
||||
cudaMemcpy(ddesc,host,sizeof(host),cudaMemcpyHostToDevice);
|
||||
dim3 hgd((unsigned)I,1,(unsigned)COUNT),ogd((unsigned)D,1,(unsigned)COUNT);
|
||||
grouped_hidden_g4_dual<<<hgd,256>>>(gate,up,xs,ddesc,I,D);
|
||||
grouped_down_g4<<<ogd,256>>>(y,gate,ddesc,D,I);
|
||||
if(cudaDeviceSynchronize()!=cudaSuccess){ printf("FAIL cuda\n"); return 1; }
|
||||
for(int c=0;c<COUNT;c++){
|
||||
int cgs=c==2?0:gs;
|
||||
float rg[512],ru[512],ry[512];
|
||||
cpu_gemv_g4(hg[c],hgs[c],D,I,cgs,xs+(size_t)c*D,rg);
|
||||
cpu_gemv_g4(hu[c],hus[c],D,I,cgs,xs+(size_t)c*D,ru);
|
||||
for(int o=0;o<I;o++){
|
||||
if(fabsf(gate[(size_t)c*I+o]-rg[o])>1e-3f*(fabsf(rg[o])+1e-3f)||
|
||||
fabsf(up[(size_t)c*I+o]-ru[o])>1e-3f*(fabsf(ru[o])+1e-3f)) bad++;
|
||||
}
|
||||
cpu_gemv_g4(hd[c],hds[c],I,D,cgs,(float*)gate+(size_t)c*I,ry);
|
||||
for(int o=0;o<D;o++)
|
||||
if(fabsf(y[(size_t)c*D+o]-ry[o])>1e-3f*(fabsf(ry[o])+1e-3f)) bad++;
|
||||
}
|
||||
for(int c=0;c<COUNT;c++){ cudaFree(qg[c]);cudaFree(qu[c]);cudaFree(qd[c]);
|
||||
cudaFree(sg[c]);cudaFree(su[c]);cudaFree(sd[c]);
|
||||
free(hg[c]);free(hu[c]);free(hd[c]);free(hgs[c]);free(hus[c]);free(hds[c]); }
|
||||
cudaFree(ddesc);cudaFree(xs);cudaFree(gate);cudaFree(up);cudaFree(y);
|
||||
}
|
||||
printf("grouped-g4 oracle: %d trials x %d experts (gs=64 + tail + per-row member), %d mismatches\n",
|
||||
trials,COUNT,bad);
|
||||
if(bad){ printf("FAIL\n"); return 1; }
|
||||
printf("OK\n"); return 0;
|
||||
}
|
||||
@@ -0,0 +1,257 @@
|
||||
"""Inefficiency / regression tests for the colibri engine (tiny model, asserted).
|
||||
|
||||
These run against the bundled glm_tiny model (~0.6 MB resident, ~0.1s/run) and
|
||||
gate CI: a regression here means something broke. They run on a plain CPU-only
|
||||
`glm.exe` build; the CUDA_* tests auto-skip if the engine wasn't built with
|
||||
CUDA_DLL=1 (see tests/README_efficiency.md for the build command).
|
||||
|
||||
The signals under test, and the inefficiency each catches:
|
||||
- tok/s floor : a throughput regression (broken build / bad config)
|
||||
- profile phases sum : telemetry accounting bug (other balloons)
|
||||
- disk-wait not dominant: a tiny resident model should never be I/O-bound
|
||||
- CPU determinism : greedy decode is reproducible (no stray RNG/threading)
|
||||
- CUDA init path : COLI_CUDA=1 initializes and does not silently exit 2
|
||||
- CUDA dense uses VRAM : CUDA_DENSE=1 actually uploads tensors (no silent CPU fallback)
|
||||
- CPU vs CUDA TF-match : identical weights+inputs → identical argmax (kernel bug guard)
|
||||
"""
|
||||
import os
|
||||
import shutil
|
||||
import unittest
|
||||
from pathlib import Path
|
||||
|
||||
from tools.efficiency import (
|
||||
parse_run, run_engine, disk_wait_share, tf_agreement,
|
||||
TINY_TOK_S_FLOOR, MAX_DISK_WAIT_SHARE, MIN_CPU_CUDA_AGREEMENT,
|
||||
)
|
||||
|
||||
HERE = Path(__file__).resolve().parent
|
||||
C_DIR = HERE.parent
|
||||
ENGINE = C_DIR / "glm.exe"
|
||||
TINY = C_DIR / "glm_tiny"
|
||||
|
||||
|
||||
def _engine_present() -> bool:
|
||||
"""True iff BOTH the built engine AND the tiny fixture are available.
|
||||
|
||||
These tests need glm.exe (a build artifact) AND glm_tiny/ (a generated
|
||||
fixture, gitignored). CI runs `make check` = "dependency-free tests, no
|
||||
model downloads" (workflow .github/workflows/check.yml, by design #140), so
|
||||
neither is present there and these tests must SKIP rather than fail. They
|
||||
run locally after `make glm.exe` (glm_tiny ships alongside the source, or
|
||||
is regenerated by tools/make_glm_oracle.py).
|
||||
"""
|
||||
return ENGINE.exists() and (TINY / "config.json").exists()
|
||||
|
||||
|
||||
def _skip_reason() -> str:
|
||||
"""Name exactly which prerequisite is missing, so the skip is actionable."""
|
||||
if not ENGINE.exists():
|
||||
return "glm.exe not built (run: make glm.exe)"
|
||||
if not (TINY / "config.json").exists():
|
||||
return "glm_tiny fixture absent (gitignored; ship it locally or run tools/make_glm_oracle.py)"
|
||||
return ""
|
||||
|
||||
|
||||
def _cuda_available() -> bool:
|
||||
"""True iff the engine binary has the CUDA loader compiled in AND the DLL is present.
|
||||
|
||||
The host is built with -DCOLI_CUDA only when CUDA_DLL=1 (Makefile). A binary
|
||||
built without it embeds the string "this binary is CPU-only; rebuild" and
|
||||
exits 2 on any CUDA env var — so we detect the CPU-only build by scanning
|
||||
the binary for that marker (avoids a slow ldd/strings on every import; we
|
||||
only read enough to find it). On Windows the DLL is also required
|
||||
(backend_loader.c loads it at runtime); on Linux it's direct-linked.
|
||||
"""
|
||||
if not _engine_present():
|
||||
return False
|
||||
try:
|
||||
# Read once; the marker is near the read-only string table. 256 KB is
|
||||
# plenty for this string and avoids loading a 1 MB binary into memory.
|
||||
with open(ENGINE, "rb") as f:
|
||||
blob = f.read(2 * 1024 * 1024)
|
||||
if b"this binary is CPU-only" in blob:
|
||||
return False # CPU-only build: COLI_CUDA=1 would exit 2.
|
||||
except OSError:
|
||||
pass
|
||||
# Windows: DLL required at runtime. Linux: direct-linked (no DLL).
|
||||
if os.name == "nt":
|
||||
return (C_DIR / "coli_cuda.dll").exists()
|
||||
import subprocess
|
||||
try:
|
||||
out = subprocess.run(["ldd", str(ENGINE)], capture_output=True, text=True)
|
||||
return "libcudart" in out.stdout
|
||||
except (FileNotFoundError, OSError):
|
||||
return False
|
||||
|
||||
|
||||
@unittest.skipUnless(_engine_present(), _skip_reason() or "glm.exe + glm_tiny required")
|
||||
class TinyEfficiencyTest(unittest.TestCase):
|
||||
"""Asserted regression tests on the resident tiny model. Gates CI."""
|
||||
|
||||
def _run(self, **overlay):
|
||||
return run_engine(overlay, engine=str(ENGINE), snap=str(TINY))[0]
|
||||
|
||||
# -- telemetry contract ---------------------------------------------------
|
||||
|
||||
def test_telemetry_parses(self):
|
||||
"""A REPLAY run must emit the throughput + PROFILE lines the suite keys on.
|
||||
|
||||
If this fails, either the engine changed its output format (update the
|
||||
parsers in tools/efficiency.py) or the run crashed early."""
|
||||
t = self._run(REPLAY="1", TEMP="0", NGEN="4")
|
||||
self.assertEqual(t["returncode"], 0, f"engine exited non-zero:\n{t['stderr']}")
|
||||
self.assertIn("tok_s", t["parsed"], f"missing tok/s line:\n{t['stderr']}")
|
||||
self.assertIn("profile", t["parsed"], f"missing PROFILE line:\n{t['stderr']}")
|
||||
self.assertIn("hit_pct", t["parsed"], f"missing expert hit line:\n{t['stderr']}")
|
||||
self.assertIsNotNone(t["tok_s"])
|
||||
self.assertIsNotNone(t["profile"])
|
||||
|
||||
# -- throughput floor -----------------------------------------------------
|
||||
|
||||
def test_tiny_tok_s_floor(self):
|
||||
"""Tiny decode must beat TINY_TOK_S_FLOOR.
|
||||
|
||||
The tiny model is fully resident and runs ~200 tok/s; the 20 tok/s
|
||||
default floor is a 10x margin that catches broken builds or a pathological
|
||||
config cascade (the cap=1 trap from ISSUE_new_model_resource_regression.md)
|
||||
without flapping on machine noise."""
|
||||
t = self._run(REPLAY="1", TEMP="0", NGEN="8")
|
||||
self.assertEqual(t["returncode"], 0, f"engine exited non-zero:\n{t['stderr']}")
|
||||
self.assertGreaterEqual(
|
||||
t["tok_s"], TINY_TOK_S_FLOOR,
|
||||
f"tok/s {t['tok_s']:.1f} below floor {TINY_TOK_S_FLOOR} "
|
||||
f"(regression, or a config cascade starving the cache)",
|
||||
)
|
||||
|
||||
# -- accounting sanity ----------------------------------------------------
|
||||
|
||||
def test_profile_phases_present_and_nonneg(self):
|
||||
"""Every PROFILE phase must be present and non-negative.
|
||||
|
||||
`other` can go slightly negative from timer overhead (the engine allows
|
||||
it), but a large negative means the timers are double-counting."""
|
||||
t = self._run(REPLAY="1", TEMP="0", NGEN="4")
|
||||
p = t["profile"]
|
||||
for phase in ("disk", "expert_matmul", "attention", "lm_head"):
|
||||
self.assertGreaterEqual(p[phase], -0.01, f"{phase} went negative: {p}")
|
||||
# 'other' is a residual; allow a small negative from timer overlap.
|
||||
self.assertGreaterEqual(p["other"], -0.05, f"other too negative (double-count): {p}")
|
||||
|
||||
# -- disk-wait not dominant on a resident model ---------------------------
|
||||
|
||||
def test_disk_wait_not_dominant(self):
|
||||
"""A fully-resident tiny model must NOT be I/O-bound.
|
||||
|
||||
Everything fits in RAM; the expert-disk wait share should be ~0. If it
|
||||
exceeds MAX_DISK_WAIT_SHARE, the cache/I/O path regressed — on a real
|
||||
model this same regression would make decode I/O-bound (the exact
|
||||
failure mode the [PROF] verdict flags)."""
|
||||
t = self._run(REPLAY="1", TEMP="0", NGEN="8", PROF="1")
|
||||
self.assertIn("time_shares", t["parsed"], f"missing [PROF] time shares:\n{t['stderr']}")
|
||||
share = disk_wait_share(t)
|
||||
self.assertIsNotNone(share)
|
||||
self.assertLess(
|
||||
share, MAX_DISK_WAIT_SHARE,
|
||||
f"expert-I/O share {share:.0%} on a resident model — I/O path regressed",
|
||||
)
|
||||
|
||||
# -- determinism ----------------------------------------------------------
|
||||
|
||||
def test_cpu_vs_cpu_determinism(self):
|
||||
"""Two greedy REPLAY runs with the same seed produce identical telemetry.
|
||||
|
||||
TEMP=0 = greedy (no sampling), so tok/s and hit-rate must be reproducible.
|
||||
A drift here means non-determinism crept into the decode path (stray
|
||||
threading, uninitialized state) — which on a real model would make A/B
|
||||
comparisons meaningless."""
|
||||
a = self._run(REPLAY="1", TEMP="0", NGEN="8", SEED="1")
|
||||
b = self._run(REPLAY="1", TEMP="0", NGEN="8", SEED="1")
|
||||
self.assertEqual(a["returncode"], 0)
|
||||
self.assertEqual(b["returncode"], 0)
|
||||
self.assertEqual(a["hit_pct"], b["hit_pct"], "greedy hit-rate drifted between runs")
|
||||
# tok/s within 25% — exact equality is too strict across scheduler noise.
|
||||
self.assertLess(abs(a["tok_s"] - b["tok_s"]) / max(a["tok_s"], b["tok_s"]), 0.25)
|
||||
|
||||
|
||||
@unittest.skipUnless(_cuda_available(),
|
||||
_skip_reason() or "CUDA build not present (run: make clean && make glm.exe CUDA_DLL=1 && make cuda-dll)")
|
||||
class TinyCudaEfficiencyTest(unittest.TestCase):
|
||||
"""CUDA-path regression tests on the tiny model. Skip unless CUDA built.
|
||||
|
||||
glm_tiny is small and fully resident, so CUDA here is fast and exercises the
|
||||
real GPU code path (init, dense upload, kernel correctness) without the
|
||||
long load time or memory pressure of the full model. These guard the
|
||||
silent-failure modes that are otherwise invisible:
|
||||
- COLI_CUDA=1 silently falling back to CPU (loader/DLL missing)
|
||||
- CUDA_DENSE=1 uploading nothing
|
||||
- a CUDA kernel producing different argmax than CPU on identical inputs
|
||||
"""
|
||||
|
||||
def _run(self, **overlay):
|
||||
return run_engine(overlay, engine=str(ENGINE), snap=str(TINY))[0]
|
||||
|
||||
def test_cuda_init_path(self):
|
||||
"""COLI_CUDA=1 must initialize the device and NOT exit 2.
|
||||
|
||||
Exit 2 is the engine's "requested backend is unavailable" path
|
||||
(glm.c: g_cuda_enabled check). A clean init prints the [CUDA] device
|
||||
banner to stderr. If this fails, the DLL is broken or the loader can't
|
||||
resolve symbols (ABI drift between backend_cuda.h and the dll)."""
|
||||
t = self._run(COLI_CUDA="1", COLI_GPU="0", REPLAY="1", TEMP="0", NGEN="4")
|
||||
self.assertNotEqual(t["returncode"], 2,
|
||||
f"engine refused CUDA backend:\n{t['stderr']}")
|
||||
self.assertTrue(t["cuda"]["enabled"],
|
||||
f"no [CUDA] device banner on stderr:\n{t['stderr']}")
|
||||
|
||||
def test_cuda_dense_uses_vram(self):
|
||||
"""CUDA_DENSE=1 must actually upload dense tensors to VRAM.
|
||||
|
||||
The minimal GPU-exercising config (per backend_loader.c analysis):
|
||||
COLI_CUDA=1 + CUDA_DENSE=1. Without CUDA_DENSE the dense path stays on
|
||||
CPU and [CUDA] resident set reports 0 tensors — a silent no-op. This
|
||||
catches that regression: after the run, resident_tensors > 0."""
|
||||
t = self._run(COLI_CUDA="1", COLI_GPU="0", CUDA_DENSE="1",
|
||||
REPLAY="1", TEMP="0", NGEN="4")
|
||||
self.assertEqual(t["returncode"], 0, f"engine exited non-zero:\n{t['stderr']}")
|
||||
rt = t["cuda"]["resident_tensors"]
|
||||
self.assertIsNotNone(rt, f"no [CUDA] resident set line:\n{t['stderr']}")
|
||||
self.assertGreater(rt, 0,
|
||||
f"CUDA_DENSE=1 but {rt} tensors resident — silent CPU fallback")
|
||||
|
||||
def test_cpu_vs_cuda_tf_match(self):
|
||||
"""CPU and CUDA teacher-forcing must agree on most positions (DIRECTLY).
|
||||
|
||||
Both paths prefill the SAME oracle sequence on the SAME weights, so their
|
||||
argmaxes should match position-for-position — but not exactly: the two
|
||||
backends accumulate dot-products in different orders (x86 SIMD vs CUDA
|
||||
kernel), so a few near-tied logits flip. That divergence is expected
|
||||
numeric behavior, not a kernel bug. A *catastrophic* kernel regression
|
||||
(wrong GEMM, wrong scale, wrong fmt) would collapse agreement toward
|
||||
random (~1/vocab = ~4%); the MIN_CPU_CUDA_AGREEMENT floor (default 70%)
|
||||
catches that while tolerating harmless drift.
|
||||
|
||||
We compare CPU-vs-CUDA *directly* (not via the oracle match-counts),
|
||||
because both backends differ from the oracle at different positions and
|
||||
the summary line can't tell "CPU≠CUDA" from "CPU≠oracle"."""
|
||||
import json
|
||||
ref = json.loads((C_DIR / "ref_glm.json").read_text())
|
||||
oracle = ref["tf_pred"]
|
||||
|
||||
cpu = self._run(TF="1", TEMP="0")
|
||||
cuda = self._run(COLI_CUDA="1", COLI_GPU="0", CUDA_DENSE="1", TF="1", TEMP="0")
|
||||
self.assertEqual(cpu["returncode"], 0, f"CPU TF run failed:\n{cpu['stderr']}")
|
||||
self.assertEqual(cuda["returncode"], 0, f"CUDA TF run failed:\n{cuda['stderr']}")
|
||||
self.assertIn("tf_mismatches", cpu["parsed"], "CPU run missing per-position mismatches")
|
||||
self.assertIn("tf_mismatches", cuda["parsed"], "CUDA run missing per-position mismatches")
|
||||
|
||||
agree, diff = tf_agreement(cpu, cuda, oracle)
|
||||
self.assertGreaterEqual(
|
||||
agree, MIN_CPU_CUDA_AGREEMENT,
|
||||
f"CPU-vs-CUDA argmax agreement {agree:.0%} below floor "
|
||||
f"{MIN_CPU_CUDA_AGREEMENT:.0%} — CUDA kernel likely regressed. "
|
||||
f"Differing positions: {diff[:10]}{'...' if len(diff)>10 else ''}",
|
||||
)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
unittest.main()
|
||||
@@ -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;
|
||||
}
|
||||
@@ -0,0 +1,96 @@
|
||||
"""fmt=5 codec oracle (#452 ladder step 2).
|
||||
|
||||
Checks the properties the container, the converter and the decode kernels all
|
||||
depend on: exact byte budget, deterministic encode, decode agreeing with a
|
||||
straight-from-the-spec reader, sign parity closure, and reconstruction quality
|
||||
matching the ablation that chose this scheme (#453).
|
||||
"""
|
||||
import os
|
||||
import sys
|
||||
import unittest
|
||||
|
||||
# The runtime path is dependency-free by design and CI keeps it that way, so the
|
||||
# offline-tooling tests skip rather than fail where numpy is absent.
|
||||
np = None
|
||||
P = None
|
||||
try:
|
||||
import numpy as np
|
||||
|
||||
sys.path.insert(0, os.path.join(os.path.dirname(__file__), "..", "tools"))
|
||||
import iq3_pack as P
|
||||
except ImportError: # pragma: no cover - exercised only on dependency-free CI
|
||||
pass
|
||||
|
||||
|
||||
def ref_decode(packed, K):
|
||||
"""Independent reader written straight from the layout comment — deliberately
|
||||
naive and loop-based, so a shared bug in the vectorized path shows up."""
|
||||
g = P.grid()
|
||||
nsb = K // P.QK
|
||||
rows = packed.reshape(-1, nsb * P.BLOCK_BYTES)
|
||||
out = np.zeros((len(rows), K), dtype=np.float32)
|
||||
for r in range(len(rows)):
|
||||
for sb in range(nsb):
|
||||
base = sb * P.BLOCK_BYTES
|
||||
d = float(rows[r, base + 96:base + 98].copy().view(np.float16)[0])
|
||||
for ib in range(P.QK // P.SUB):
|
||||
word = int(np.ascontiguousarray(
|
||||
rows[r, base + 64 + ib * 4:base + 64 + ib * 4 + 4]).view(np.uint32)[0])
|
||||
db = d * (0.5 + ((word >> 28) & 0xF)) * 0.5
|
||||
for l in range(4):
|
||||
seven = (word >> (7 * l)) & 0x7F
|
||||
bits = [(seven >> j) & 1 for j in range(7)]
|
||||
bits.append(sum(bits) & 1) # odd parity closes the 8th
|
||||
idx = int(rows[r, base + ib * 8 + l * 2 + 0])
|
||||
idx2 = int(rows[r, base + ib * 8 + l * 2 + 1])
|
||||
mags = list(g[idx]) + list(g[idx2])
|
||||
for j in range(8):
|
||||
pos = sb * P.QK + ib * P.SUB + l * 8 + j
|
||||
out[r, pos] = mags[j] * db * (-1.0 if bits[j] else 1.0)
|
||||
return out
|
||||
|
||||
|
||||
@unittest.skipIf(P is None, "numpy not available (offline-tooling test)")
|
||||
class TestIq3Pack(unittest.TestCase):
|
||||
def setUp(self):
|
||||
np.random.seed(4242)
|
||||
self.x = (np.random.randn(6, 1024) * 0.05).astype(np.float32)
|
||||
|
||||
def test_byte_budget(self):
|
||||
self.assertEqual(P.BLOCK_BYTES, 98)
|
||||
self.assertAlmostEqual(P.bpw(), 3.0625, places=6)
|
||||
packed = P.encode(self.x)
|
||||
self.assertEqual(packed.shape, (6, 1024 // P.QK * 98))
|
||||
self.assertEqual(packed.dtype, np.uint8)
|
||||
|
||||
def test_encode_is_deterministic(self):
|
||||
self.assertTrue(np.array_equal(P.encode(self.x), P.encode(self.x)))
|
||||
|
||||
def test_decode_matches_spec_reader(self):
|
||||
packed = P.encode(self.x)
|
||||
fast = P.decode(packed, 1024)
|
||||
slow = ref_decode(packed, 1024)
|
||||
self.assertTrue(np.allclose(fast, slow, rtol=1e-6, atol=1e-8),
|
||||
f"max |Δ| = {np.abs(fast - slow).max()}")
|
||||
|
||||
def test_sign_parity_closes(self):
|
||||
"""Every 8-weight block must have an even number of negatives — that is
|
||||
what lets the 8th sign be derived instead of stored."""
|
||||
y = P.decode(P.encode(self.x), 1024)
|
||||
neg = (y < 0).reshape(-1, 8).sum(-1)
|
||||
self.assertTrue(np.all(neg % 2 == 0), "a block stored odd negatives")
|
||||
|
||||
def test_reconstruction_quality(self):
|
||||
y = P.decode(P.encode(self.x), 1024)
|
||||
rel = np.sqrt(((y - self.x) ** 2).mean()) / np.sqrt((self.x ** 2).mean())
|
||||
# the torch model that won the #453 A/B measures ~0.195 on this input class
|
||||
self.assertLess(rel, 0.25, f"rel-RMSE {rel:.4f} — worse than the chosen scheme")
|
||||
self.assertGreater(rel, 0.05, f"rel-RMSE {rel:.4f} — implausibly good, check the test")
|
||||
|
||||
def test_shape_guard(self):
|
||||
with self.assertRaises(ValueError):
|
||||
P.encode(np.zeros((2, 300), dtype=np.float32))
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
unittest.main()
|
||||
@@ -20,8 +20,8 @@ class FakeEngine:
|
||||
self.calls = []
|
||||
|
||||
def generate(self, prompt, maximum, temperature, top_p, on_text, cache_slot=0,
|
||||
cancelled=None):
|
||||
self.calls.append((prompt, maximum, temperature, top_p, cache_slot))
|
||||
cancelled=None, grammar=None):
|
||||
self.calls.append((prompt, maximum, temperature, top_p, cache_slot, grammar))
|
||||
on_text("Hé")
|
||||
on_text("llo")
|
||||
return {"prompt_tokens": 7, "completion_tokens": 2, "length_limited": False}
|
||||
@@ -34,7 +34,7 @@ class BlockingEngine(FakeEngine):
|
||||
self.release = threading.Event()
|
||||
|
||||
def generate(self, prompt, maximum, temperature, top_p, on_text, cache_slot=0,
|
||||
cancelled=None):
|
||||
cancelled=None, grammar=None):
|
||||
self.entered.set()
|
||||
self.release.wait(2)
|
||||
return super().generate(prompt, maximum, temperature, top_p, on_text, cache_slot,
|
||||
@@ -71,11 +71,11 @@ class TemplateTest(unittest.TestCase):
|
||||
|
||||
def test_validates_generation_limits(self):
|
||||
self.assertEqual(generation_options({"max_tokens": 4, "temperature": 0, "top_p": 1}, 8),
|
||||
(4, 0.0, 1.0))
|
||||
(4, 0.0, 1.0, None))
|
||||
# max_tokens above the server cap is clamped, not rejected (#260): OpenAI
|
||||
# clients default to large values; erroring breaks them.
|
||||
self.assertEqual(generation_options({"max_tokens": 9, "temperature": 0, "top_p": 1}, 8),
|
||||
(8, 0.0, 1.0))
|
||||
(8, 0.0, 1.0, None))
|
||||
# non-positive / non-int max_tokens is still a hard error
|
||||
with self.assertRaises(APIError):
|
||||
generation_options({"max_tokens": 0}, 8)
|
||||
@@ -84,7 +84,31 @@ class TemplateTest(unittest.TestCase):
|
||||
with self.assertRaises(APIError):
|
||||
generation_options({"top_p": math.inf}, 8)
|
||||
self.assertEqual(generation_options({"temperature": None, "top_p": None}, 8),
|
||||
(8, 0.7, 0.9))
|
||||
(8, 0.7, 0.9, None))
|
||||
# response_format -> grammar plumbing (draft source, never a constraint)
|
||||
opts = generation_options({"max_tokens": 4, "response_format": {"type": "json_object"}}, 8)
|
||||
self.assertIn("root ::=", opts[3])
|
||||
schema = {"type": "object", "properties": {"a": {"type": "string"}}, "required": ["a"]}
|
||||
opts = generation_options({"max_tokens": 4, "response_format":
|
||||
{"type": "json_schema", "json_schema": {"schema": schema}}}, 8)
|
||||
self.assertEqual(json.loads(opts[3]), schema)
|
||||
opts = generation_options({"max_tokens": 4, "response_format":
|
||||
{"type": "gbnf", "grammar": 'root ::= "x"'}}, 8)
|
||||
self.assertEqual(opts[3], 'root ::= "x"')
|
||||
with self.assertRaises(APIError):
|
||||
generation_options({"response_format": {"type": "yaml"}}, 8)
|
||||
with self.assertRaises(APIError):
|
||||
generation_options({"response_format": {"type": "json_schema", "json_schema": {}}}, 8)
|
||||
with self.assertRaises(APIError): # non-dict response_format
|
||||
generation_options({"response_format": "json"}, 8)
|
||||
with self.assertRaises(APIError): # empty gbnf
|
||||
generation_options({"response_format": {"type": "gbnf", "grammar": " "}}, 8)
|
||||
with self.assertRaises(APIError): # oversized grammar (> 1 MiB pre-check)
|
||||
generation_options({"response_format": {"type": "gbnf", "grammar": "x" * ((1 << 20) + 1)}}, 8)
|
||||
# malformed GBNF passes the gateway by design: the ENGINE fail-softs it
|
||||
# (draft source only — bad grammar costs the speedup, never the request)
|
||||
opts = generation_options({"response_format": {"type": "gbnf", "grammar": "not a grammar ::="}}, 8)
|
||||
self.assertEqual(opts[3], "not a grammar ::=")
|
||||
|
||||
|
||||
class ProtocolTest(unittest.TestCase):
|
||||
|
||||
@@ -0,0 +1,169 @@
|
||||
/* COLI_PIPE_BLOCK: the pipe pool's condvar waiter must be observably
|
||||
* equivalent to the sched_yield spin it replaces — same bytes land in the
|
||||
* same ws[] slots, and no interleaving loses a wakeup (the worker RELEASE-
|
||||
* stores ready[] BEFORE taking mx to broadcast; the waiter re-checks under
|
||||
* the lock, so a flag set between its fast-path check and the wait cannot
|
||||
* be missed). Both waiters are exercised against the same on-disk fixture,
|
||||
* alternating parked waits (wait issued before the load finishes) with
|
||||
* fast-path waits (load already done), across enough generations to cycle
|
||||
* the pool's gen-tagged cursor.
|
||||
*
|
||||
* Also pins the PIPE_WORKERS => PIPE implication table: fires ONLY when
|
||||
* PIPE is unset in the env AND the platform default left the pipe off AND
|
||||
* PIPE_WORKERS parses positive (PIPE_WORKERS=0/empty/negative must NOT
|
||||
* silently enable a clamped 1-worker pipe). */
|
||||
#include <errno.h>
|
||||
#include <fcntl.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
#include <unistd.h>
|
||||
#define main coli_glm_main_unused
|
||||
#include "../colibri.c"
|
||||
#undef main
|
||||
|
||||
static int fail(const char *s){ fprintf(stderr,"FAIL: %s\n",s); return 1; }
|
||||
|
||||
enum { NE=8, LAYER=1 }; /* experts 0..NE-1 on one MoE layer */
|
||||
/* per-expert file image: [gate 12][up 12][down 12][gate.qs 12][up.qs 12][down.qs 16] */
|
||||
enum { WB=12, QS_G=12, QS_U=12, QS_D=16, ESZ=3*WB+QS_G+QS_U+QS_D };
|
||||
|
||||
static unsigned char wbyte(int e,int j){ return (unsigned char)(e*31+j+1); }
|
||||
static float scale(int e,int i){ return (float)(e*8+i)+0.5f; }
|
||||
|
||||
#define TMPF "test_pipe_block.tmp"
|
||||
|
||||
static int write_fixture(void){
|
||||
FILE *w=fopen(TMPF,"wb"); if(!w) return fail("create temp");
|
||||
for(int e=0;e<NE;e++){
|
||||
unsigned char img[ESZ];
|
||||
for(int j=0;j<3*WB;j++) img[j]=wbyte(e,j);
|
||||
float sc[(QS_G+QS_U+QS_D)/4];
|
||||
for(int i=0;i<(int)(sizeof(sc)/sizeof(sc[0]));i++) sc[i]=scale(e,i);
|
||||
memcpy(img+3*WB,sc,sizeof(sc));
|
||||
if(fwrite(img,1,ESZ,w)!=ESZ){ fclose(w); return fail("expert fixture write"); }
|
||||
}
|
||||
fclose(w);
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int build_fixture(Model *m,int fd){
|
||||
m->c.hidden=4; m->c.moe_inter=3; m->ebits=8;
|
||||
m->S.n=NE*6; m->S.cap=NE*6; m->S.t=calloc(NE*6,sizeof(st_tensor));
|
||||
if(!m->S.t) return fail("tensor metadata allocation");
|
||||
const char *proj[3]={"gate_proj","up_proj","down_proj"};
|
||||
int sbytes[3]={QS_G,QS_U,QS_D};
|
||||
for(int e=0;e<NE;e++){
|
||||
int64_t wo=(int64_t)e*ESZ, so=wo+3*WB;
|
||||
for(int k=0;k<3;k++){
|
||||
char name[300];
|
||||
snprintf(name,sizeof(name),"model.layers.%d.mlp.experts.%d.%s.weight",LAYER,e,proj[k]);
|
||||
m->S.t[e*6+k]=(st_tensor){strdup(name),fd,wo,WB,3,WB}; wo+=WB;
|
||||
size_t n=strlen(name); memcpy(name+n,".qs",4);
|
||||
m->S.t[e*6+3+k]=(st_tensor){strdup(name),fd,so,sbytes[k],2,sbytes[k]/4}; so+=sbytes[k];
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int check_slot(ESlot *s,int e){
|
||||
if(s->eid!=e || s->g.fmt!=1 || s->u.fmt!=1 || s->d.fmt!=1){
|
||||
fprintf(stderr," slot: eid=%d (want %d) fmt g/u/d=%d/%d/%d (want 1/1/1)\n",
|
||||
s->eid,e,s->g.fmt,s->u.fmt,s->d.fmt);
|
||||
return 1;
|
||||
}
|
||||
const unsigned char *g=(const unsigned char*)s->g.q8,
|
||||
*u=(const unsigned char*)s->u.q8,
|
||||
*d=(const unsigned char*)s->d.q8; /* q8 is int8_t; compare raw bytes */
|
||||
for(int j=0;j<WB;j++)
|
||||
if(g[j]!=wbyte(e,j) || u[j]!=wbyte(e,WB+j) || d[j]!=wbyte(e,2*WB+j)){
|
||||
fprintf(stderr," slot e=%d weight byte %d: g=%d/%d u=%d/%d d=%d/%d (got/want)\n",e,j,
|
||||
g[j],wbyte(e,j),u[j],wbyte(e,WB+j),d[j],wbyte(e,2*WB+j));
|
||||
return 1;
|
||||
}
|
||||
for(int i=0;i<3;i++)
|
||||
if(s->g.s[i]!=scale(e,i) || s->u.s[i]!=scale(e,3+i)){
|
||||
fprintf(stderr," slot e=%d scale %d: g=%g/%g u=%g/%g (got/want)\n",e,i,
|
||||
(double)s->g.s[i],(double)scale(e,i),(double)s->u.s[i],(double)scale(e,3+i));
|
||||
return 1;
|
||||
}
|
||||
for(int i=0;i<4;i++)
|
||||
if(s->d.s[i]!=scale(e,6+i)){
|
||||
fprintf(stderr," slot e=%d scale %d: d=%g/%g (got/want)\n",e,i,
|
||||
(double)s->d.s[i],(double)scale(e,6+i));
|
||||
return 1;
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int run_generations(Model *m,int block,int gens){
|
||||
g_pipe_block=block;
|
||||
for(int gen=0;gen<gens;gen++){
|
||||
int eids[NE];
|
||||
for(int q=0;q<NE;q++) eids[q]=(gen*3+q)%NE; /* deterministic shuffle across gens */
|
||||
pipe_dispatch(m,LAYER,eids,NE);
|
||||
if(gen%4==0) usleep(300); /* let loads finish → fast-path wait */
|
||||
for(int i=0;i<NE;i++){
|
||||
/* odd gens wait on the LAST-dispatched slot first: with jobs this
|
||||
* small, in-order waits mostly find ready already set — reverse
|
||||
* order is what actually parks the waiter on the condvar. */
|
||||
int q=(gen&1)?NE-1-i:i;
|
||||
pipe_wait(q);
|
||||
if(!atomic_load_explicit(&g_pp.ready[q],memory_order_acquire))
|
||||
return fail(block?"blocking wait returned before ready":"spin wait returned before ready");
|
||||
if(check_slot(&m->ws[q],eids[q])) return fail(block?"slot contents (block)":"slot contents (spin)");
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int test_implication_table(void){
|
||||
struct { const char *pipe_env,*pw_env; int pipe_now,want; } T[]={
|
||||
{NULL,"4",0,1}, /* pool sized, pipe off, PIPE unset → imply */
|
||||
{NULL,"16",0,1},
|
||||
{NULL,"0",0,0}, /* PIPE_WORKERS=0 must NOT enable a clamped pipe */
|
||||
{NULL,"",0,0},
|
||||
{NULL,"-3",0,0},
|
||||
{"0","4",0,0}, /* explicit PIPE=0 always wins */
|
||||
{"1","4",1,0}, /* explicit PIPE=1: nothing to imply */
|
||||
{NULL,"4",1,0}, /* platform default already ON (win32) */
|
||||
{NULL,NULL,0,0},
|
||||
};
|
||||
for(size_t i=0;i<sizeof(T)/sizeof(T[0]);i++)
|
||||
if(pipe_workers_imply_pipe(T[i].pipe_env,T[i].pw_env,T[i].pipe_now)!=T[i].want){
|
||||
fprintf(stderr,"FAIL: implication row %zu (PIPE=%s PIPE_WORKERS=%s pipe_now=%d)\n",
|
||||
i,T[i].pipe_env?T[i].pipe_env:"<unset>",T[i].pw_env?T[i].pw_env:"<unset>",T[i].pipe_now);
|
||||
return 1;
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
int main(void){
|
||||
if(test_implication_table()) return 1;
|
||||
|
||||
/* Relative to the CWD, like test_compat_direct's TMPF — NOT "/tmp/...":
|
||||
* the windows job builds native .exe files and "/tmp" is not a Windows
|
||||
* path. fwrite then reopen read-only: Windows compat has pread, not pwrite. */
|
||||
if(write_fixture()) return 1;
|
||||
int fd=open(TMPF,COMPAT_O_RDONLY);
|
||||
if(fd<0) return fail("open temp");
|
||||
|
||||
static Model m; /* zeroed: buffered pread path, no mmap/cuda */
|
||||
if(build_fixture(&m,fd)){ close(fd); remove(TMPF); return 1; }
|
||||
|
||||
g_pipe=1; g_pipe_nw=4;
|
||||
pipe_init(&m);
|
||||
|
||||
/* spin waiter first (control), then the condvar waiter under the same
|
||||
* dispatch pattern; 200 generations each cycles the gen-tagged cursor
|
||||
* and alternates parked/fast-path waits. */
|
||||
if(run_generations(&m,0,200) || run_generations(&m,1,200)){ close(fd); remove(TMPF); return 1; }
|
||||
|
||||
for(int q=0;q<NE;q++){ compat_aligned_free(m.ws[q].slab); free(m.ws[q].fslab); }
|
||||
for(int i=0;i<m.S.n;i++) free(m.S.t[i].name);
|
||||
free(m.S.t);
|
||||
close(fd);
|
||||
remove(TMPF);
|
||||
puts("test_pipe_block: ok");
|
||||
return 0;
|
||||
}
|
||||
@@ -119,7 +119,7 @@ int main(void){
|
||||
for(int i=0;i<O;i++) sc[i]=0.01f+0.001f*(i%7);
|
||||
for(size_t i=0;i<(size_t)S*K;i++) x[i]=rndf();
|
||||
ColiCudaTensor *t=NULL;
|
||||
ok&=coli_cuda_matmul(&t,ref,x,w4,sc,2,S,K,O,0); /* host path = reference */
|
||||
ok&=coli_cuda_matmul(&t,ref,x,w4,sc,2,S,K,O,0,0); /* host path = reference */
|
||||
float *xd=(float*)coli_cuda_pipe_alloc(0,(size_t)S*K*4);
|
||||
float *yd=(float*)coli_cuda_pipe_alloc(0,(size_t)S*O*4);
|
||||
ok&=coli_cuda_pipe_upload(0,xd,x,(size_t)S*K*4);
|
||||
|
||||
@@ -5,6 +5,7 @@ import sys
|
||||
import tempfile
|
||||
import unittest
|
||||
from pathlib import Path
|
||||
from unittest import mock
|
||||
|
||||
from resource_plan import (
|
||||
GB,
|
||||
@@ -14,6 +15,7 @@ from resource_plan import (
|
||||
environment_for_plan,
|
||||
format_plan,
|
||||
memory_available,
|
||||
physical_cpu_count,
|
||||
)
|
||||
|
||||
|
||||
@@ -87,6 +89,79 @@ class ResourcePlanTest(unittest.TestCase):
|
||||
self.assertIn("clamped", plan["warnings"][0])
|
||||
self.assertIn("0:test-gpu", format_plan(plan))
|
||||
|
||||
def test_auto_tier_thread_count_uses_physical_cores(self):
|
||||
# End-to-end for #325: build_plan + environment_for_plan must export the
|
||||
# physical (not logical SMT) core count as OMP_NUM_THREADS. The original
|
||||
# suite passed physical_cpus=24 explicitly, so it never exercised the
|
||||
# real physical_cpu_count() probe whose single-core failure pinned decode.
|
||||
def lscpu(stdout):
|
||||
return subprocess.CompletedProcess(args=[], returncode=0,
|
||||
stdout=stdout, stderr="")
|
||||
# 1 socket, 12 cores, 2 SMT siblings -> 24 threads, 12 physical cores.
|
||||
|
||||
# The parser must return 12 physical cores under BOTH lscpu layouts:
|
||||
# - 2-col: `lscpu -p=core,socket` emits exactly [core,socket] (this is
|
||||
# what the probe actually requests; the previous fields[1]/[2]
|
||||
# indexing skipped every line here and fell through to the
|
||||
# logical count -> the regression JustVugg caught).
|
||||
# - 3-col: bare `lscpu -p` prepends a CPU column -> [cpu,core,socket].
|
||||
# Taking the last two fields is correct in both cases.
|
||||
layouts = {
|
||||
"2-col (-p=core,socket)": (
|
||||
"# core,socket\n" +
|
||||
"\n".join(f"{core},0" for core in range(12) for _ in range(2))),
|
||||
"3-col (bare -p, CPU prefix)": (
|
||||
"# CPU,Core,Socket\n" +
|
||||
"\n".join(f"{cpu},{core},0" for core in range(12) for cpu in range(2))),
|
||||
}
|
||||
for label, blob in layouts.items():
|
||||
with mock.patch("resource_plan.subprocess.run",
|
||||
return_value=lscpu(blob)), \
|
||||
mock.patch.object(sys, "platform", "linux"):
|
||||
plan = build_plan(self.model, available_memory=16 * GB,
|
||||
available_disk=1, gpus=[])
|
||||
env = environment_for_plan(plan)
|
||||
self.assertEqual(plan["cpu"]["physical_cores"], 12, label)
|
||||
self.assertEqual(env["OMP_NUM_THREADS"], "12", label)
|
||||
|
||||
def test_plan_does_not_set_omp_affinity_vars(self):
|
||||
# The real #325 regression: --auto-tier set OMP_PROC_BIND=spread +
|
||||
# OMP_PLACES=cores, which ran before the engine's overwrite=0 setenv and
|
||||
# so won, collapsing the OpenMP team to one CPU on the reporter's 64-core
|
||||
# Linux box even though OMP_NUM_THREADS was correct. The plan must leave
|
||||
# affinity to the engine's own hot-thread tuning (which prefers 'close').
|
||||
plan = build_plan(self.model, available_memory=16 * GB,
|
||||
available_disk=1, gpus=[], physical_cpus=64)
|
||||
env = environment_for_plan(plan)
|
||||
self.assertEqual(env["OMP_NUM_THREADS"], "64")
|
||||
self.assertNotIn("OMP_PROC_BIND", env)
|
||||
self.assertNotIn("OMP_PLACES", env)
|
||||
|
||||
def test_plan_conserves_budget_and_experts_above_256gb(self):
|
||||
# Regression for #325's reporter: a 512 GB machine loading the whole
|
||||
# model into RAM. Verify the budget math stays exact at large RAM sizes
|
||||
# (no integer truncation, no over-allocation, no experts lost between
|
||||
# tiers). Checked at 256/512/800 GB to bracket the reporter's box.
|
||||
for ram_gb in (256, 512, 800):
|
||||
plan = build_plan(self.model, ram_gb=ram_gb, available_disk=1,
|
||||
gpus=[], physical_cpus=64)
|
||||
ram = plan["tiers"]["ram"]
|
||||
# RAM budget never over-allocated: dense + runtime + cache <= budget.
|
||||
allocated = (ram["dense_bytes"] + ram["runtime_bytes"]
|
||||
+ ram["expert_cache_bytes"])
|
||||
self.assertLessEqual(allocated, ram["budget_bytes"],
|
||||
f"over-allocated RAM at {ram_gb} GB")
|
||||
# Every expert byte is accounted for exactly once across the tiers.
|
||||
tiers = plan["tiers"]
|
||||
tiered = (tiers["vram"]["hot_expert_bytes"]
|
||||
+ ram["warm_expert_bytes"]
|
||||
+ tiers["disk"]["cold_expert_bytes"])
|
||||
self.assertEqual(tiered, plan["model"]["expert_bytes"],
|
||||
f"expert bytes lost/duplicated at {ram_gb} GB")
|
||||
# A positive RAM budget yields a non-negative cache and a sensible cap.
|
||||
self.assertGreaterEqual(ram["expert_cache_bytes"], 0)
|
||||
self.assertGreaterEqual(ram["cache_slots_per_layer"], 0)
|
||||
|
||||
def test_filters_requested_devices(self):
|
||||
gpus = [{"index": 0, "name": "a", "total_bytes": 8 * GB, "free_bytes": 8 * GB}]
|
||||
plan = build_plan(self.model, available_memory=16 * GB, available_disk=1,
|
||||
@@ -117,13 +192,13 @@ class ResourcePlanTest(unittest.TestCase):
|
||||
self.assertEqual(env["COLI_CUDA"], "1")
|
||||
self.assertEqual(env["COLI_GPUS"], "1")
|
||||
self.assertEqual(env["OMP_NUM_THREADS"], str(plan["cpu"]["physical_cores"]))
|
||||
if sys.platform == "win32":
|
||||
# MinGW libgomp: niente affinity su Windows, le chiavi non vanno emesse
|
||||
self.assertNotIn("OMP_PROC_BIND", env)
|
||||
self.assertNotIn("OMP_PLACES", env)
|
||||
else:
|
||||
self.assertEqual(env["OMP_PROC_BIND"], "spread")
|
||||
self.assertEqual(env["OMP_PLACES"], "cores")
|
||||
# The plan must NOT set OMP_PROC_BIND / OMP_PLACES on any platform:
|
||||
# the engine's own hot-thread tuning owns affinity (it prefers
|
||||
# OMP_PROC_BIND=close for the back-to-back per-expert matmuls). Setting
|
||||
# spread + cores here ran before the engine's overwrite=0 setenv and so
|
||||
# won, collapsing the team to one CPU on some libgomp topologies (#325).
|
||||
self.assertNotIn("OMP_PROC_BIND", env)
|
||||
self.assertNotIn("OMP_PLACES", env)
|
||||
self.assertEqual(env["PIN_GB"], env["CUDA_EXPERT_GB"])
|
||||
|
||||
explicit_threads = environment_for_plan(plan, {"OMP_NUM_THREADS": "7",
|
||||
@@ -140,6 +215,81 @@ class ResourcePlanTest(unittest.TestCase):
|
||||
plan = build_plan(self.model, available_memory=16 * GB, available_disk=1,
|
||||
gpus=[], physical_cpus=8, cpu_sockets=1)
|
||||
self.assertNotIn("COLI_NUMA", environment_for_plan(plan))
|
||||
|
||||
def test_auto_tune_mtp_off_when_compute_bound(self):
|
||||
# Tiny model with 64 GB RAM and no GPU: all experts fit in RAM with no
|
||||
# warm tier, so the plan classifies as compute-bound.
|
||||
plan = build_plan(self.model, ram_gb=64, available_memory=64 * GB,
|
||||
available_disk=100 * GB, gpus=[], physical_cpus=24,
|
||||
cpu_sockets=2)
|
||||
# With such a small model fully in RAM and no GPU, bottleneck is compute
|
||||
self.assertEqual(plan["bottleneck_class"], "compute")
|
||||
self.assertIn("DRAFT", plan["tune"])
|
||||
self.assertEqual(plan["tune"]["DRAFT"]["value"], "0")
|
||||
env = environment_for_plan(plan)
|
||||
self.assertEqual(env["DRAFT"], "0")
|
||||
explicit = environment_for_plan(plan, {"DRAFT": "3"})
|
||||
self.assertEqual(explicit["DRAFT"], "3")
|
||||
|
||||
def test_auto_tune_mtp_off_when_disk_low_hit(self):
|
||||
# Use a model large enough that 8 GB RAM can't hold all experts.
|
||||
big = tempfile.TemporaryDirectory()
|
||||
bigmodel = Path(big.name)
|
||||
(bigmodel / "config.json").write_text(json.dumps({
|
||||
"num_hidden_layers": 2, "n_routed_experts": 4,
|
||||
"kv_lora_rank": 4, "qk_rope_head_dim": 2,
|
||||
"qk_nope_head_dim": 3, "v_head_dim": 5, "num_attention_heads": 2,
|
||||
}))
|
||||
expert_size = 3 * GB # each expert 3 GB → 12 GB total, won't fit in 8 GB budget
|
||||
write_shard(bigmodel / "out-00000.safetensors", [
|
||||
("model.embed_tokens.weight", 100),
|
||||
("model.layers.0.self_attn.q_a_proj.weight", 200),
|
||||
])
|
||||
for i in range(4):
|
||||
write_shard(bigmodel / f"out-{i+1:05d}.safetensors", [
|
||||
(f"model.layers.1.mlp.experts.{i}.gate_proj.weight", expert_size),
|
||||
])
|
||||
plan = build_plan(bigmodel, ram_gb=0, available_memory=4 * GB,
|
||||
available_disk=100 * GB, gpus=[], physical_cpus=8,
|
||||
cpu_sockets=1)
|
||||
big.cleanup()
|
||||
self.assertEqual(plan["bottleneck_class"], "disk")
|
||||
self.assertLess(plan["projected_hit_rate"], 0.90)
|
||||
self.assertEqual(plan["tune"]["DRAFT"]["value"], "0")
|
||||
|
||||
def test_auto_tune_pipe_multi_gpu(self):
|
||||
gpus = [
|
||||
{"index": 0, "name": "a", "total_bytes": 32 * GB, "free_bytes": 30 * GB},
|
||||
{"index": 1, "name": "b", "total_bytes": 32 * GB, "free_bytes": 30 * GB},
|
||||
]
|
||||
plan = build_plan(self.model, ram_gb=16, available_memory=32 * GB,
|
||||
available_disk=1, gpus=gpus, cpu_sockets=2)
|
||||
self.assertEqual(plan["tune"]["COLI_CUDA_PIPE"]["value"], "2")
|
||||
env = environment_for_plan(plan)
|
||||
self.assertEqual(env["COLI_CUDA_PIPE"], "2")
|
||||
|
||||
def test_auto_tune_pipe_single_gpu(self):
|
||||
gpus = [{"index": 0, "name": "a", "total_bytes": 12 * GB, "free_bytes": 10 * GB}]
|
||||
plan = build_plan(self.model, ram_gb=16, available_memory=32 * GB,
|
||||
available_disk=1, gpus=gpus, cpu_sockets=1)
|
||||
self.assertEqual(plan["tune"]["COLI_CUDA_PIPE"]["value"], "1")
|
||||
|
||||
def test_auto_tune_numa_hint_for_cpu_only(self):
|
||||
plan = build_plan(self.model, ram_gb=64, available_memory=64 * GB,
|
||||
available_disk=1, gpus=[], physical_cpus=64, cpu_sockets=2)
|
||||
self.assertIn("_numa_hint", plan["tune"])
|
||||
self.assertIn("numactl", plan["tune"]["_numa_hint"])
|
||||
self.assertIn("auto-tune", format_plan(plan))
|
||||
|
||||
def test_format_plan_shows_tune_and_hit_rate(self):
|
||||
plan = build_plan(self.model, ram_gb=64, available_memory=64 * GB,
|
||||
available_disk=100 * GB, gpus=[], physical_cpus=24,
|
||||
cpu_sockets=1)
|
||||
text = format_plan(plan)
|
||||
self.assertIn("hit", text)
|
||||
self.assertIn("auto-tune", text)
|
||||
self.assertIn("DRAFT", text)
|
||||
|
||||
def test_cpu_binary_does_not_apply_gpu_tier(self):
|
||||
plan = build_plan(self.model, available_memory=16 * GB, available_disk=1,
|
||||
gpus=[{"index": 0, "name": "a", "total_bytes": 8 * GB,
|
||||
@@ -178,5 +328,73 @@ class ResourcePlanTest(unittest.TestCase):
|
||||
self.assertIn("expected_bottleneck", plan)
|
||||
|
||||
|
||||
class PhysicalCpuCountTest(unittest.TestCase):
|
||||
"""Regression for #325: --auto-tier pinned decode to one core because
|
||||
physical_cpu_count() silently returned 1.
|
||||
|
||||
Two root causes this locks down:
|
||||
1. lscpu -p prepends a CPU column, so `-p=core,socket` emits
|
||||
CPU,Core,Socket; counting rows counted logical SMT siblings.
|
||||
2. any probe failure fell through to ``os.cpu_count() or 1`` and the
|
||||
``or 1`` could pin a constrained/cgroup'd box to a single core.
|
||||
"""
|
||||
|
||||
def _lscpu(self, stdout):
|
||||
return subprocess.CompletedProcess(args=[], returncode=0,
|
||||
stdout=stdout, stderr="")
|
||||
|
||||
def _lscpu_topology(self, sockets, cores_per_socket, threads_per_core):
|
||||
# Real lscpu shape: socket-local core IDs repeat across sockets; the
|
||||
# CPU column (always prepended) is a unique logical-CPU index.
|
||||
rows, cpu = [], 0
|
||||
for sock in range(sockets):
|
||||
for core in range(cores_per_socket):
|
||||
for _ in range(threads_per_core):
|
||||
rows.append(f"{cpu},{core},{sock}")
|
||||
cpu += 1
|
||||
return "# CPU,Core,Socket\n" + "\n".join(rows)
|
||||
|
||||
def test_counts_physical_cores_not_smt_threads(self):
|
||||
blob = self._lscpu_topology(sockets=2, cores_per_socket=16, threads_per_core=2)
|
||||
with mock.patch("resource_plan.subprocess.run", return_value=self._lscpu(blob)), \
|
||||
mock.patch.object(sys, "platform", "linux"):
|
||||
self.assertEqual(physical_cpu_count(), 32)
|
||||
|
||||
def test_single_socket_no_smt(self):
|
||||
blob = self._lscpu_topology(sockets=1, cores_per_socket=8, threads_per_core=1)
|
||||
with mock.patch("resource_plan.subprocess.run", return_value=self._lscpu(blob)), \
|
||||
mock.patch.object(sys, "platform", "linux"):
|
||||
self.assertEqual(physical_cpu_count(), 8)
|
||||
|
||||
def test_skips_offline_core_socket_fields(self):
|
||||
# VMs / large NUMA boxes emit "-" for offline core or socket IDs; that
|
||||
# used to raise ValueError, discard the whole parse, and fall through
|
||||
# to the single-core fallback.
|
||||
blob = "# CPU,Core,Socket\n0,0,0\n1,-,0\n2,1,0\n3,1,0\n"
|
||||
with mock.patch("resource_plan.subprocess.run", return_value=self._lscpu(blob)), \
|
||||
mock.patch.object(sys, "platform", "linux"):
|
||||
self.assertEqual(physical_cpu_count(), 2)
|
||||
|
||||
def test_lscpu_missing_falls_back_to_logical_not_silent_one(self):
|
||||
# The bug: lscpu absent -> os.cpu_count() or 1. On a constrained box
|
||||
# os.cpu_count() can be 1. We still must never silently pick 1 without
|
||||
# a warning, and when logical cores exist they must be used.
|
||||
import os
|
||||
with mock.patch("resource_plan.subprocess.run", side_effect=FileNotFoundError), \
|
||||
mock.patch.object(sys, "platform", "linux"), \
|
||||
mock.patch("resource_plan.os.cpu_count", return_value=16), \
|
||||
mock.patch("sys.stderr"):
|
||||
self.assertEqual(physical_cpu_count(), 16)
|
||||
|
||||
def test_zero_logical_cores_warns_and_returns_one(self):
|
||||
# The genuine degenerate case: no probe works and os.cpu_count() is
|
||||
# None/1. Must return 1 (engine needs a positive team size) but warn.
|
||||
with mock.patch("resource_plan.subprocess.run", side_effect=FileNotFoundError), \
|
||||
mock.patch.object(sys, "platform", "linux"), \
|
||||
mock.patch("resource_plan.os.cpu_count", return_value=None), \
|
||||
mock.patch("sys.stderr"):
|
||||
self.assertEqual(physical_cpu_count(), 1)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
unittest.main()
|
||||
|
||||
@@ -0,0 +1,79 @@
|
||||
/* Device-router kernel oracle (#431 PR-A).
|
||||
*
|
||||
* Feeds random activations/router weights through pipe_router_logits +
|
||||
* pipe_router_select and checks against a CPU reference that replicates
|
||||
* moe()'s plain routing path verbatim (sigmoid -> bias-augmented top-K by
|
||||
* `choice`, weights from raw `logit`, route-level TOPP truncation, norm_topk,
|
||||
* routed_scale). The dot/expf rounding may differ from libm at ~1e-6 rel, so
|
||||
* a handful of near-tie index flips across trials is tolerated; the weight
|
||||
* math itself must agree to 1e-4 rel on matching selections.
|
||||
*
|
||||
* Build: nvcc -O2 -std=c++17 -arch=native tests/test_router_cuda.cu -o tests/test_router_cuda
|
||||
*/
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <cmath>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
/* pull in the kernel definitions (same idiom as the CPU tests' #include "../colibri.c") */
|
||||
#include "../backend_cuda.cu"
|
||||
|
||||
static void cpu_ref(const float *x,const float *W,const float *bias,int D,int E,
|
||||
int Ksel,float topp,int norm_topk,float rscale,
|
||||
int *idx,float *w,int *keff){
|
||||
float *logit=(float*)malloc(E*sizeof(float)),*choice=(float*)malloc(E*sizeof(float));
|
||||
for(int e=0;e<E;e++){ double a=0; const float *r=W+(size_t)e*D;
|
||||
for(int i=0;i<D;i++) a+=(double)x[i]*r[i];
|
||||
float lg=1.f/(1.f+expf(-(float)a)); logit[e]=lg; choice[e]=lg+bias[e]; }
|
||||
for(int kk=0;kk<Ksel;kk++){ int best=-1; float bv=-1e30f;
|
||||
for(int e=0;e<E;e++){ int tk=0; for(int j=0;j<kk;j++) if(idx[j]==e){tk=1;break;}
|
||||
if(!tk && choice[e]>bv){bv=choice[e];best=e;} }
|
||||
idx[kk]=best; w[kk]=logit[best]; }
|
||||
int Ke=Ksel;
|
||||
if(topp>0.f && topp<1.f){
|
||||
for(int a=1;a<Ksel;a++){ int ii=idx[a]; float ww=w[a]; int b=a-1;
|
||||
while(b>=0 && w[b]<ww){ w[b+1]=w[b]; idx[b+1]=idx[b]; b--; } w[b+1]=ww; idx[b+1]=ii; }
|
||||
float tot=1e-20f; for(int kk=0;kk<Ksel;kk++) tot+=w[kk];
|
||||
float cum=0; for(int kk=0;kk<Ksel;kk++){ cum+=w[kk]; if(cum>=topp*tot){ Ke=kk+1; break; } } }
|
||||
if(norm_topk){ float sm=0; for(int kk=0;kk<Ke;kk++) sm+=w[kk]; sm+=1e-20f;
|
||||
for(int kk=0;kk<Ke;kk++) w[kk]/=sm; }
|
||||
for(int kk=0;kk<Ke;kk++) w[kk]*=rscale;
|
||||
*keff=Ke; free(logit); free(choice);
|
||||
}
|
||||
|
||||
int main(void){
|
||||
const int D=6144,E=256,K=8,TRIALS=200;
|
||||
srand(42);
|
||||
float *x,*W,*b; cudaMallocManaged(&x,D*4); cudaMallocManaged(&W,(size_t)E*D*4);
|
||||
cudaMallocManaged(&b,E*4);
|
||||
float *lg,*ch; char *out;
|
||||
cudaMalloc(&lg,E*4); cudaMalloc(&ch,E*4); cudaMalloc(&out,K*8+4);
|
||||
int flips=0, bad=0;
|
||||
for(int t=0;t<TRIALS;t++){
|
||||
float topp = (t%3==1)?0.7f:0.f;
|
||||
int nt = (t%2);
|
||||
float rs = 1.0f+(t%5)*0.25f;
|
||||
for(int i=0;i<D;i++) x[i]=(rand()/(float)RAND_MAX-.5f)*2.f;
|
||||
for(size_t i=0;i<(size_t)E*D;i++) W[i]=(rand()/(float)RAND_MAX-.5f)*.06f;
|
||||
for(int e=0;e<E;e++) b[e]=(rand()/(float)RAND_MAX-.5f)*.02f;
|
||||
pipe_router_logits<<<E,128>>>(x,W,b,D,lg,ch);
|
||||
pipe_router_select<<<1,1>>>(lg,ch,E,K,topp,nt,rs,out);
|
||||
char pack[K*8+4];
|
||||
if(cudaMemcpy(pack,out,sizeof(pack),cudaMemcpyDeviceToHost)!=cudaSuccess){
|
||||
printf("FAIL cuda\n"); return 1; }
|
||||
int gidx[K],gkeff; float gw[K];
|
||||
memcpy(gidx,pack,K*4); memcpy(gw,pack+K*4,K*4); memcpy(&gkeff,pack+K*8,4);
|
||||
int ridx[K],rkeff; float rw[K];
|
||||
cpu_ref(x,W,b,D,E,K,topp,nt,rs,ridx,rw,&rkeff);
|
||||
int mism=0; for(int k2=0;k2<K;k2++) if(gidx[k2]!=ridx[k2]) mism++;
|
||||
if(mism||gkeff!=rkeff){ flips++; continue; } /* near-tie flip: counted, tolerated */
|
||||
for(int k2=0;k2<gkeff;k2++){
|
||||
float ref=rw[k2], d=fabsf(gw[k2]-ref);
|
||||
if(d>1e-4f*(fabsf(ref)+1e-6f)+1e-6f){ bad++; break; }
|
||||
}
|
||||
}
|
||||
printf("router oracle: %d trials, %d near-tie flips, %d weight mismatches\n",TRIALS,flips,bad);
|
||||
if(flips>4||bad){ printf("FAIL\n"); return 1; }
|
||||
printf("OK\n"); return 0;
|
||||
}
|
||||
Binary file not shown.
@@ -0,0 +1,64 @@
|
||||
/* o200k pre-tokenizer validation against HF-tokenizers-generated expectations.
|
||||
* Self-contained for the test-c harness: loads tests/tok_o200k_tiny.json (a
|
||||
* synthetic byte-level BPE whose Split regex is the o200k pattern — a few KB,
|
||||
* no model download) and scores tests/tok_o200k_cases.txt, whose expected ids
|
||||
* were produced by HF `tokenizers` on the same file. Guards the case-aware
|
||||
* letter matcher, contractions, digit groups, the [\r\n/]* punctuation tail,
|
||||
* whitespace branches, and added-token atomicity; round-trips every case.
|
||||
* The cl100k path is untouched by construction (dispatch requires \p{Lu} in
|
||||
* the tokenizer's own Split pattern) and stays covered by the GLM oracle. */
|
||||
#define _GNU_SOURCE
|
||||
#include "../tok.h"
|
||||
|
||||
int main(void) {
|
||||
Tok T;
|
||||
tok_load(&T, "tests/tok_o200k_tiny.json");
|
||||
if (!T.o200k) { fprintf(stderr, "test_tok_o200k: o200k pattern not detected\n"); return 1; }
|
||||
FILE *f = fopen("tests/tok_o200k_cases.txt", "rb");
|
||||
if (!f) { perror("tests/tok_o200k_cases.txt"); return 1; }
|
||||
/* fgets, not getline: MinGW's UCRT lacks getline and this must run on
|
||||
* the windows job. Case lines are short; 8 KB is generous. */
|
||||
char line[8192];
|
||||
int pass = 0, tot = 0, dpass = 0;
|
||||
while (fgets(line, sizeof(line), f)) {
|
||||
size_t nr = strlen(line);
|
||||
while (nr > 0 && (line[nr-1] == '\n' || line[nr-1] == '\r')) line[--nr] = 0;
|
||||
if (nr == 0) continue;
|
||||
char *tab = strchr(line, '\t'); if (!tab) continue;
|
||||
*tab = 0;
|
||||
const char *text = line, *idstr = tab + 1;
|
||||
char tbuf[4096]; int tn = 0;
|
||||
for (const char *q = text; *q && tn < 4095; q++) {
|
||||
if (q[0]=='\\' && q[1]=='n') { tbuf[tn++]='\n'; q++; }
|
||||
else if (q[0]=='\\' && q[1]=='t') { tbuf[tn++]='\t'; q++; }
|
||||
else if (q[0]=='\\' && q[1]=='r') { tbuf[tn++]='\r'; q++; }
|
||||
else if (q[0]=='\\' && q[1]=='\\') { tbuf[tn++]='\\'; q++; }
|
||||
else tbuf[tn++] = *q;
|
||||
}
|
||||
tbuf[tn] = 0;
|
||||
int exp[512], ne = 0;
|
||||
for (const char *q = idstr; *q; ) {
|
||||
while (*q == ',' || *q == ' ') q++;
|
||||
if (!*q) break;
|
||||
exp[ne++] = atoi(q);
|
||||
while (*q && *q != ',') q++;
|
||||
}
|
||||
int got[512]; int ng = tok_encode(&T, tbuf, tn, got, 512);
|
||||
int ok = (ng == ne);
|
||||
for (int i = 0; i < ng && ok; i++) ok = (got[i] == exp[i]);
|
||||
tot++; if (ok) pass++;
|
||||
char dec[8192]; int dn = tok_decode(&T, got, ng, dec, 8191);
|
||||
int drt = (dn == tn) && !memcmp(dec, tbuf, tn);
|
||||
if (drt) dpass++;
|
||||
if (!ok || !drt) {
|
||||
fprintf(stderr, "MISMATCH text=%s\n exp(%d):", text, ne);
|
||||
for (int i = 0; i < ne; i++) fprintf(stderr, " %d", exp[i]);
|
||||
fprintf(stderr, "\n got(%d):", ng);
|
||||
for (int i = 0; i < ng; i++) fprintf(stderr, " %d", got[i]);
|
||||
fprintf(stderr, "\n decode_ok=%d\n", drt);
|
||||
}
|
||||
}
|
||||
fclose(f);
|
||||
printf("test_tok_o200k: ENCODE %d/%d DECODE %d/%d\n", pass, tot, dpass, tot);
|
||||
return (pass == tot && dpass == tot) ? 0 : 2;
|
||||
}
|
||||
@@ -0,0 +1,40 @@
|
||||
hello world 259,32,119,111,114,263
|
||||
HelloWorld 72,101,257,111,262
|
||||
XMLHttpRequest 88,77,76,72,116,116,112,82,101,113,117,101,115,116
|
||||
helloWORLDhello 259,87,79,82,76,68,259
|
||||
dog's 100,111,103,270
|
||||
DON'T 68,79,78,39,84
|
||||
don't 100,111,110,39,116
|
||||
I'll've 73,39,257,39,118,101
|
||||
O'Brien 79,39,66,114,105,101,110
|
||||
the theatre 265,32,116,256,97,116,114,101
|
||||
12345 268,52,53
|
||||
a1b22c333d4444 97,49,98,50,50,99,51,51,51,100,52,52,52,52
|
||||
3.14 51,46,49,52
|
||||
http://x.com/a/b 104,116,116,112,58,47,47,120,46,99,111,109,47,97,47,98
|
||||
path/to/file 112,97,264,47,116,111,47,102,105,108,101
|
||||
a//b///c 97,47,47,98,47,47,47,99
|
||||
!!\r\n//x 33,33,13,10,47,47,120
|
||||
one\ntwo\r\nthree 111,110,101,10,116,119,111,13,10,264,114,101,101
|
||||
\n x 32,32,10,32,32,120
|
||||
32,32,32
|
||||
a b c 97,32,32,98,32,32,32,99
|
||||
tab\there 116,269,9,256,114,101
|
||||
Café 67,97,102,101,204,129
|
||||
naiveBayes 110,97,105,118,101,66,97,121,101,115
|
||||
Éclair 195,137,99,108,97,105,114
|
||||
北京大学 229,140,151,228,186,172,229,164,167,229,173,166
|
||||
ΑΒαβ 206,145,206,146,206,177,206,178
|
||||
Иван 208,152,208,178,208,176,208,189
|
||||
ẞßscharf 225,186,158,195,159,115,99,104,97,114,102
|
||||
i̇stanbul 105,204,135,115,116,97,110,98,117,108
|
||||
hello<|endoftext|>world 259,274,119,111,114,263
|
||||
<|message_user|>hi 275,104,105
|
||||
mixedCASEand123 109,105,120,101,100,67,65,83,69,97,110,100,268
|
||||
's 270
|
||||
's 32,39,115
|
||||
A 65
|
||||
aB 97,66
|
||||
Ab 65,98
|
||||
AB 65,66
|
||||
ab 269
|
||||
@@ -0,0 +1 @@
|
||||
{"version": "1.0", "truncation": null, "padding": null, "added_tokens": [{"id": 274, "content": "<|endoftext|>", "single_word": false, "lstrip": false, "rstrip": false, "normalized": false, "special": true}, {"id": 275, "content": "<|message_user|>", "single_word": false, "lstrip": false, "rstrip": false, "normalized": false, "special": true}], "normalizer": null, "pre_tokenizer": {"type": "Sequence", "pretokenizers": [{"type": "Split", "pattern": {"Regex": "[^\\r\\n\\p{L}\\p{N}]?[\\p{Lu}\\p{Lt}\\p{Lm}\\p{Lo}\\p{M}]*[\\p{Ll}\\p{Lm}\\p{Lo}\\p{M}]+(?i:'s|'t|'re|'ve|'m|'ll|'d)?|[^\\r\\n\\p{L}\\p{N}]?[\\p{Lu}\\p{Lt}\\p{Lm}\\p{Lo}\\p{M}]+[\\p{Ll}\\p{Lm}\\p{Lo}\\p{M}]*(?i:'s|'t|'re|'ve|'m|'ll|'d)?|\\p{N}{1,3}| ?[^\\s\\p{L}\\p{N}]+[\\r\\n/]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+"}, "behavior": "Isolated", "invert": false}, {"type": "ByteLevel", "add_prefix_space": false, "trim_offsets": true, "use_regex": false}]}, "post_processor": null, "decoder": {"type": "ByteLevel", "add_prefix_space": true, "trim_offsets": true, "use_regex": true}, "model": {"type": "BPE", "dropout": null, "unk_token": null, "continuing_subword_prefix": null, "end_of_word_suffix": null, "fuse_unk": false, "byte_fallback": false, "ignore_merges": true, "vocab": {"Ā": 0, "ā": 1, "Ă": 2, "ă": 3, "Ą": 4, "ą": 5, "Ć": 6, "ć": 7, "Ĉ": 8, "ĉ": 9, "Ċ": 10, "ċ": 11, "Č": 12, "č": 13, "Ď": 14, "ď": 15, "Đ": 16, "đ": 17, "Ē": 18, "ē": 19, "Ĕ": 20, "ĕ": 21, "Ė": 22, "ė": 23, "Ę": 24, "ę": 25, "Ě": 26, "ě": 27, "Ĝ": 28, "ĝ": 29, "Ğ": 30, "ğ": 31, "Ġ": 32, "!": 33, "\"": 34, "#": 35, "$": 36, "%": 37, "&": 38, "'": 39, "(": 40, ")": 41, "*": 42, "+": 43, ",": 44, "-": 45, ".": 46, "/": 47, "0": 48, "1": 49, "2": 50, "3": 51, "4": 52, "5": 53, "6": 54, "7": 55, "8": 56, "9": 57, ":": 58, ";": 59, "<": 60, "=": 61, ">": 62, "?": 63, "@": 64, "A": 65, "B": 66, "C": 67, "D": 68, "E": 69, "F": 70, "G": 71, "H": 72, "I": 73, "J": 74, "K": 75, "L": 76, "M": 77, "N": 78, "O": 79, "P": 80, "Q": 81, "R": 82, "S": 83, "T": 84, "U": 85, "V": 86, "W": 87, "X": 88, "Y": 89, "Z": 90, "[": 91, "\\": 92, "]": 93, "^": 94, "_": 95, "`": 96, "a": 97, "b": 98, "c": 99, "d": 100, "e": 101, "f": 102, "g": 103, "h": 104, "i": 105, "j": 106, "k": 107, "l": 108, "m": 109, "n": 110, "o": 111, "p": 112, "q": 113, "r": 114, "s": 115, "t": 116, "u": 117, "v": 118, "w": 119, "x": 120, "y": 121, "z": 122, "{": 123, "|": 124, "}": 125, "~": 126, "ġ": 127, "Ģ": 128, "ģ": 129, "Ĥ": 130, "ĥ": 131, "Ħ": 132, "ħ": 133, "Ĩ": 134, "ĩ": 135, "Ī": 136, "ī": 137, "Ĭ": 138, "ĭ": 139, "Į": 140, "į": 141, "İ": 142, "ı": 143, "IJ": 144, "ij": 145, "Ĵ": 146, "ĵ": 147, "Ķ": 148, "ķ": 149, "ĸ": 150, "Ĺ": 151, "ĺ": 152, "Ļ": 153, "ļ": 154, "Ľ": 155, "ľ": 156, "Ŀ": 157, "ŀ": 158, "Ł": 159, "ł": 160, "¡": 161, "¢": 162, "£": 163, "¤": 164, "¥": 165, "¦": 166, "§": 167, "¨": 168, "©": 169, "ª": 170, "«": 171, "¬": 172, "Ń": 173, "®": 174, "¯": 175, "°": 176, "±": 177, "²": 178, "³": 179, "´": 180, "µ": 181, "¶": 182, "·": 183, "¸": 184, "¹": 185, "º": 186, "»": 187, "¼": 188, "½": 189, "¾": 190, "¿": 191, "À": 192, "Á": 193, "Â": 194, "Ã": 195, "Ä": 196, "Å": 197, "Æ": 198, "Ç": 199, "È": 200, "É": 201, "Ê": 202, "Ë": 203, "Ì": 204, "Í": 205, "Î": 206, "Ï": 207, "Ð": 208, "Ñ": 209, "Ò": 210, "Ó": 211, "Ô": 212, "Õ": 213, "Ö": 214, "×": 215, "Ø": 216, "Ù": 217, "Ú": 218, "Û": 219, "Ü": 220, "Ý": 221, "Þ": 222, "ß": 223, "à": 224, "á": 225, "â": 226, "ã": 227, "ä": 228, "å": 229, "æ": 230, "ç": 231, "è": 232, "é": 233, "ê": 234, "ë": 235, "ì": 236, "í": 237, "î": 238, "ï": 239, "ð": 240, "ñ": 241, "ò": 242, "ó": 243, "ô": 244, "õ": 245, "ö": 246, "÷": 247, "ø": 248, "ù": 249, "ú": 250, "û": 251, "ü": 252, "ý": 253, "þ": 254, "ÿ": 255, "he": 256, "ll": 257, "hell": 258, "hello": 259, "Wo": 260, "Wor": 261, "World": 262, "ld": 263, "th": 264, "the": 265, "Ġthe": 266, "12": 267, "123": 268, "ab": 269, "'s": 270, "Ġa": 271, "./": 272, "ĊĊ": 273}, "merges": [["h", "e"], ["l", "l"], ["he", "ll"], ["hell", "o"], ["W", "o"], ["Wo", "r"], ["Wor", "ld"], ["l", "d"], ["t", "h"], ["th", "e"], ["Ġ", "the"], ["1", "2"], ["12", "3"], ["a", "b"], ["'", "s"], ["Ġ", "a"], [".", "/"], ["Ċ", "Ċ"]]}}
|
||||
Reference in New Issue
Block a user