From 63966ba3f8866417ee567b63824deed8165f7965 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Tue, 14 Jul 2026 01:23:18 +0200 Subject: [PATCH 1/3] Dual-SSD streaming: COLI_MODEL_MIRROR reads experts from two model copies MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A second (read-only) copy of the model on another drive is registered as a per-shard read replica; expert loads are split between the drives by a deterministic (layer,eid) hash weighted by COLI_DISK_WEIGHTS=, or, by default, by a startup bandwidth probe using the engine's own access pattern (parallel ~19 MB reads, O_DIRECT). Cold decode is disk-bound (~11 GB/token), so two NVMe queues add up. - st.h: mirror accepted per file only if size + safetensors header are byte-identical to the primary (identical data_offsets by construction, so every pread is valid on either copy); partial mirrors work (smaller second SSD holding only some shards); the mirror is never written — .coli_usage / .coli_kv stay on the primary. - glm.c: routing covers the coalesced slab pread, O_DIRECT, mmap/Metal and scale reads, plus the OMP-parallel pin/autopin warmup (streams from both drives). Deterministic routing keeps readahead/PILOT WILLNEED on the same drive as the demand read and avoids caching an expert twice. A mirror read error falls back to the primary (one warning, no crash). Per-drive bytes are reported in a MIRROR: stats line. - tests: test_st_mirror covers validation, read equality on both replicas, and the rejection paths (divergent header, size mismatch, missing file). Measured on 2x NVMe (GLM-5.2 int4, greedy, DRAFT=0, DIRECT=1, cold-ish cache): 0.42 -> 0.57 tok/s (+36%), expert-disk service 15.2s -> 10.3s, byte-identical output; probe chose a 48/52 split. --- README.md | 16 ++++ c/.gitignore | 1 + c/Makefile | 5 +- c/coli | 3 + c/colibri.c | 164 ++++++++++++++++++++++++++++++++++++--- c/st.h | 87 ++++++++++++++++++++- c/tests/test_st_mirror.c | 109 ++++++++++++++++++++++++++ 7 files changed, 370 insertions(+), 15 deletions(-) create mode 100644 c/tests/test_st_mirror.c diff --git a/README.md b/README.md index d6f9bcb..9e33501 100644 --- a/README.md +++ b/README.md @@ -82,6 +82,22 @@ precision are the same whether an expert answered from VRAM or from disk. VRAM / RAM / NVMe three-tier expert residency

+### Dual-SSD: two copies of the model, twice the read bandwidth + +Decode is disk-bound on most machines, and expert reads are read-only — so if you have a **second SSD**, put a full copy of the model on it and let the engine stream from both drives at once: + +```bash +COLI_MODEL=/fast/glm52_i4 COLI_MODEL_MIRROR=/second/glm52_i4 ./coli chat +COLI_DISK_WEIGHTS=9,3 ... # optional: primary,mirror bandwidth ratio (else measured at startup) +``` + +Each expert is routed to one drive by a deterministic hash, weighted by the two drives' measured (or declared) bandwidth, so readahead/PILOT prefetch and the demand read always hit the same drive and nothing is cached twice. The aggregate bandwidth is the sum of both drives — a 9 GB/s + 3 GB/s pair reads experts ~33% faster than the fast drive alone, and the OMP-parallel pin/warmup load streams from both. Details worth knowing: + +- the mirror is **validated at startup** (per-file size + safetensors header must be byte-identical to the primary); divergent or missing files silently stay on the primary, so a **partial mirror is fine** — a smaller second SSD holding only some shards still helps; +- the mirror is **never written**: `.coli_usage`, `.coli_kv` and all sidecars stay on the primary; +- a read error on the mirror falls back to the primary (one warning, no crash), so unplugging the second drive mid-run degrades instead of killing the server; +- routing never changes tokens — both copies are byte-identical, and the per-run `MIRROR:` stats line shows GB served per drive. + The same engine spans the whole range: on a 25 GB laptop everything streams from disk (slow but correct); on a large host the entire expert set becomes resident (`CUDA_EXPERT_GB=auto PIN_GB=all`) and disk drops out of the decode path diff --git a/c/.gitignore b/c/.gitignore index d8e1f88..2ea32c9 100644 --- a/c/.gitignore +++ b/c/.gitignore @@ -3,3 +3,4 @@ glm_tiny/ olmoe_hf/ olmoe_i4/ .build-config +tests/test_st_mirror diff --git a/c/Makefile b/c/Makefile index af17357..bb49499 100644 --- a/c/Makefile +++ b/c/Makefile @@ -166,7 +166,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_mirror$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -300,6 +300,9 @@ tests/test_st_pread$(EXE): tests/test_st_pread.c st.h json.h compat.h tests/test_st$(EXE): tests/test_st.c st.h json.h compat.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +tests/test_st_mirror$(EXE): tests/test_st_mirror.c st.h json.h compat.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + tests/test_tier$(EXE): tests/test_tier.c tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) diff --git a/c/coli b/c/coli index d6728fe..6378ce5 100755 --- a/c/coli +++ b/c/coli @@ -15,6 +15,9 @@ Run GLM-5.2 (744B) locally on CPU with roughly 15-26 GB of RAM. Configuration through environment variables or flags (also valid after the subcommand): COLI_MODEL= model directory (default /home/vincenzo/glm52_i4) + COLI_MODEL_MIRROR= second copy of the model on another drive: expert reads + are split across both SSDs (COLI_DISK_WEIGHTS=9,3 sets the + primary,mirror bandwidth ratio; default: measured at startup) --ram N RAM budget in GB (automatically sizes the expert cache) --repin N adapt RAM/VRAM experts every N tokens --topp P adaptive expert top-p --topk N fixed top-k diff --git a/c/colibri.c b/c/colibri.c index 03eb878..e5e08f3 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -994,6 +994,57 @@ static void *map_of_fd(int fd){ return base; } +/* ==================== DUAL-SSD: two model copies, two drives ==================== + * COLI_MODEL_MIRROR= registers a SECOND (read-only) copy of the model on + * another drive; expert reads are split between the two according to + * COLI_DISK_WEIGHTS=, (relative bandwidth; without the env it + * is measured at startup with the engine's own access pattern). Cold decode is + * disk-bound (~11 GB/token): two NVMe drives reading in parallel add up. */ +static const char *g_mirror_dir=NULL; /* COLI_MODEL_MIRROR / SNAP_MIRROR */ +static int g_mirror=0; /* 1 = mirror active (at least one shard accepted) */ +static int g_mir_share=64; /* expert share routed to the mirror, out of 256 */ +static _Atomic int64_t g_mir_bytes[2]; /* bytes served per drive: [0] primary, [1] mirror */ +static _Atomic int64_t g_mir_nread[2]; + +/* replica of one expert: DETERMINISTIC hash of (layer,eid). Determinism is a + * requirement, not a style choice: the readahead/PILOT WILLNEED and the demand + * pread must hit the same fd/page-cache, and in buffered mode an expert must + * never be cached twice (one copy per drive). */ +static inline int expert_route(int layer,int eid){ + if(!g_mirror) return 0; + uint32_t h=(uint32_t)layer*2654435761u ^ (uint32_t)eid*0x9E3779B9u; + h^=h>>16; h*=0x45d9f3bu; h^=h>>16; + return (int)(h&255) < g_mir_share; +} + +/* buffered fd of the replica, falling back to the primary if the file is not mirrored */ +static inline int rep_bfd(shards *S,int fd,int rep){ + int r=st_fd_rep(S,fd,rep); return r<0?fd:r; +} + +static int pread_full(int fd, void *buf, int64_t n, int64_t off, const char *tag); + +/* pread on the chosen replica with fallback to the primary on error/short-read: + * an unreadable sector (or an unmount) of the mirror must never kill the process + * when the primary can serve the same bytes. Delegates to pread_full so the + * mirror path inherits the short-read/EINTR loop and honest reporting (#236). + * Accounts bytes per drive. Returns 0 = ok, -1 = real error/EOF (like pread_full). */ +static ssize_t mir_pread(shards *S,int fd,int rep,void *buf,int64_t n,int64_t off,const char *tag){ + int rfd = st_fd_rep(S,fd,rep); + int used = rep && rfd>=0; + if(rfd<0) rfd=fd; + int rc=pread_full(rfd,buf,n,off,tag); + if(rc && used){ + static _Atomic int warned; + if(!atomic_exchange(&warned,1)) + fprintf(stderr,"[MIRROR] read error on the mirror copy — falling back to the primary drive\n"); + used=0; rc=pread_full(fd,buf,n,off,tag); + } + if(!rc){ atomic_fetch_add_explicit(&g_mir_bytes[used],n,memory_order_relaxed); + atomic_fetch_add_explicit(&g_mir_nread[used],1,memory_order_relaxed); } + return rc; +} + /* carica un expert nello slot. Container pre-quantizzato: le 3 matrici sono contigue nel * file -> UNA pread coalescente da ~19 MB dentro `slab` (+ le scale in fslab); i QT sono * viste dentro lo slab (zero copie). Fallback per modelli non quantizzati (oracolo tiny). @@ -1061,10 +1112,12 @@ static int expert_load_impl(Model *m, int layer, int eid, ESlot *s, int fatal){ else { __atomic_add_fetch(&m->ld_main,1,__ATOMIC_RELAXED); __atomic_add_fetch(&m->bytes_main,(uint64_t)tb,__ATOMIC_RELAXED); } } + int rep=expert_route(layer,eid); /* DUAL-SSD: this expert's replica */ + if(rep && st_fd_rep(&m->S,tw[0]->fd,1)<0) rep=0; /* shard not in the mirror (partial) */ if(g_mmap){ void *bw[3],*bq[3]; int okm=1; for(int k=0;k<3;k++){ - bw[k]=map_of_fd(tw[k]->fd); bq[k]=map_of_fd(tq[k]->fd); + bw[k]=map_of_fd(rep_bfd(&m->S,tw[k]->fd,rep)); bq[k]=map_of_fd(rep_bfd(&m->S,tq[k]->fd,rep)); if(!bw[k]||!bq[k]||((tw[k]->off)&3)||((tq[k]->off)&3)) okm=0; } if(okm){ @@ -1099,7 +1152,9 @@ static int expert_load_impl(Model *m, int layer, int eid, ESlot *s, int fatal){ * the final resident set only, after GPU release has already nulled out the * pointers for anything that isn't genuinely RAM-tier. */ atomic_fetch_add_explicit(&g_prof_io,(int64_t)(n+nq),memory_order_relaxed); + atomic_fetch_add_explicit(&g_mir_bytes[rep],tw[k]->nbytes+tq[k]->nbytes,memory_order_relaxed); } + atomic_fetch_add_explicit(&g_mir_nread[rep],1,memory_order_relaxed); s->eid=eid; return 0; } } @@ -1163,7 +1218,7 @@ static int expert_load_impl(Model *m, int layer, int eid, ESlot *s, int fatal){ int64_t pos[3]; int done=0; if(contig){ int64_t off0=tw[ord[0]]->off; - int dfd = g_direct ? st_direct_fd(&m->S, tw[ord[0]]->fd) : -1; + int dfd = g_direct ? st_direct_fd_rep(&m->S, tw[ord[0]]->fd, rep) : -1; if(dfd>=0){ /* O_DIRECT: offset/len allineati a 4K */ int64_t base=off0 & ~4095LL, need=(off0-base)+wtot; int64_t len=(need+4095)&~4095LL; @@ -1171,28 +1226,31 @@ static int expert_load_impl(Model *m, int layer, int eid, ESlot *s, int fatal){ if(r>=need){ pos[ord[0]]=off0-base; pos[ord[1]]=pos[ord[0]]+tw[ord[0]]->nbytes; pos[ord[2]]=pos[ord[1]]+tw[ord[1]]->nbytes; done=1; + atomic_fetch_add_explicit(&g_mir_bytes[rep],(int64_t)r,memory_order_relaxed); + atomic_fetch_add_explicit(&g_mir_nread[rep],1,memory_order_relaxed); } } if(!done){ /* fallback bufferizzato */ - if(pread_full(tw[ord[0]]->fd, s->slab, wtot, off0, "pread expert")){ if(fatal) exit(1); return -1; } + if(mir_pread(&m->S, tw[ord[0]]->fd, rep, s->slab, wtot, off0, "pread expert")){ if(fatal) exit(1); return -1; } pos[ord[0]]=0; pos[ord[1]]=tw[ord[0]]->nbytes; pos[ord[2]]=tw[ord[0]]->nbytes+tw[ord[1]]->nbytes; done=1; } } if(!done){ /* non contigui: 3 pread bufferizzate */ int64_t o=0; for(int a=0;a<3;a++){ int k=ord[a]; - if(pread_full(tw[k]->fd, s->slab+o, tw[k]->nbytes, tw[k]->off, "pread expert")){ if(fatal) exit(1); return -1; } + if(mir_pread(&m->S, tw[k]->fd, rep, s->slab+o, tw[k]->nbytes, tw[k]->off, "pread expert")){ if(fatal) exit(1); return -1; } pos[k]=o; o+=tw[k]->nbytes; } } float *fp[3]; int64_t fo=0; /* scale (piccole) */ for(int k=0;k<3;k++){ - if(pread_full(tq[k]->fd, (char*)(s->fslab+fo), tq[k]->nbytes, tq[k]->off, "pread qs")){ if(fatal) exit(1); return -1; } + if(mir_pread(&m->S, tq[k]->fd, rep, (char*)(s->fslab+fo), tq[k]->nbytes, tq[k]->off, "pread qs")){ if(fatal) exit(1); return -1; } fp[k]=s->fslab+fo; fo+=tq[k]->nbytes/4; } atomic_fetch_add_explicit(&g_prof_io,wtot+fo*4,memory_order_relaxed); if(g_drop){ /* scarta subito le pagine: evita che la page - * cache in pressione strangoli il throughput */ - posix_fadvise(tw[ord[0]]->fd, tw[ord[0]]->off, wtot, POSIX_FADV_DONTNEED); - for(int k=0;k<3;k++) posix_fadvise(tq[k]->fd, tq[k]->off, tq[k]->nbytes, POSIX_FADV_DONTNEED); + * cache in pressione strangoli il throughput. + * The drop targets the fd of the replica READ. */ + posix_fadvise(rep_bfd(&m->S,tw[ord[0]]->fd,rep), tw[ord[0]]->off, wtot, POSIX_FADV_DONTNEED); + for(int k=0;k<3;k++) posix_fadvise(rep_bfd(&m->S,tq[k]->fd,rep), tq[k]->off, tq[k]->nbytes, POSIX_FADV_DONTNEED); } QT *qt[3]={&s->g,&s->u,&s->d}; int OO[3]={I,I,D}, II[3]={D,D,I}; for(int k=0;k<3;k++){ @@ -1541,13 +1599,14 @@ static void expert_host_ensure(Model *m, int layer, ESlot *s){ #endif /* prefetch asincrono dei pesi di un expert (e delle sue scale .qs): avvia il readahead - * cosi' le letture sincrone successive trovano la page-cache calda. */ + * cosi' le letture sincrone successive trovano la page-cache calda. The readahead + * targets the SAME replica that will serve the pread (expert_route is deterministic). */ static void expert_prefetch(Model *m, int layer, int eid){ - char nm[300]; + char nm[300]; int rep=expert_route(layer,eid); const char *suf[3]={"gate_proj.weight","up_proj.weight","down_proj.weight"}; for(int k=0;k<3;k++){ - snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.%d.%s",layer,eid,suf[k]); st_prefetch(&m->S,nm); - char qs[320]; snprintf(qs,sizeof(qs),"%s.qs",nm); st_prefetch(&m->S,qs); + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.experts.%d.%s",layer,eid,suf[k]); st_prefetch_rep(&m->S,nm,rep); + char qs[320]; snprintf(qs,sizeof(qs),"%s.qs",nm); st_prefetch_rep(&m->S,qs,rep); } } @@ -3794,6 +3853,14 @@ static void profile_print(Model *m, double elapsed){ (unsigned long long)m->cpu_expert_rows,m->t_egpu,m->t_route,m->t_p2p,(unsigned long long)m->n_p2p, elapsed-m->t_ewait-m->t_emm-m->t_attn-m->t_head-m->t_route-m->t_p2p>0? elapsed-m->t_ewait-m->t_emm-m->t_attn-m->t_head-m->t_route-m->t_p2p:0); + if(g_mirror){ + double b0=atomic_load_explicit(&g_mir_bytes[0],memory_order_relaxed)/1e9; + double b1=atomic_load_explicit(&g_mir_bytes[1],memory_order_relaxed)/1e9; + printf("MIRROR: primary %.2f GB (%lld reads) | mirror %.2f GB (%lld reads) — %.0f%% of expert bytes from the mirror\n", + b0,(long long)atomic_load_explicit(&g_mir_nread[0],memory_order_relaxed), + b1,(long long)atomic_load_explicit(&g_mir_nread[1],memory_order_relaxed), + b0+b1>0?100.0*b1/(b0+b1):0.0); + } #ifdef COLI_METAL if(g_metal_enabled){ uint64_t ok=0,fb=0,ex=0; double su=0,gp=0,sc=0; coli_metal_moe_counts(&ok,&fb,&ex); coli_metal_moe_times(&su,&gp,&sc); @@ -3899,6 +3966,8 @@ static void run_replay(Model *m, const int *full, int nfull, int np){ m->hits=m->miss=m->ereq=m->gpu_expert_calls=0; m->hit_pin=m->hit_ecache=0; profile_reset(m); ProfBase pb; prof_base(m,&pb); + atomic_store(&g_mir_bytes[0],0); atomic_store(&g_mir_bytes[1],0); + atomic_store(&g_mir_nread[0],0); atomic_store(&g_mir_nread[1],0); double t0=now_s(); int steps=0; for(int i=np-1;infd;i++){ + if(rep && S->mfds[i]<0) continue; + int64_t sz=lseek(rep?S->mfds[i]:S->fds[i],0,SEEK_END); + if(sz>bsz){ bsz=sz; big=i; } + } + const int64_t blk=19ll<<20; const int NB=8; + if(big<0 || bszmdfds[big] : S->dfds[big]; + int fd = dfd>=0? dfd : (rep? S->mfds[big] : S->fds[big]); + if(dfd<0) fprintf(stderr,"[MIRROR] no O_DIRECT on %s: the probe may read the page cache — " + "set COLI_DISK_WEIGHTS for an accurate split\n", rep?"the mirror":"the primary"); + double t0=now_s(); int64_t tot=0; + #pragma omp parallel for schedule(dynamic,1) reduction(+:tot) + for(int i=0;i0) tot+=r; + compat_aligned_free(buf); + } + } + double dt=now_s()-t0; + return (dt>0 && tot>0)? tot/1e9/dt : 0; +} + +/* DUAL-SSD setup: register the mirror copy, derive the read split from + * COLI_DISK_WEIGHTS=, or from a startup bandwidth probe. + * Runs after model_init (needs the shard index) and BEFORE any pin/autopin + * load, so the OMP-parallel pin warmup already streams from both drives. */ +static void mirror_setup(Model *m){ + if(!g_mirror_dir) return; + const char *snap=getenv("SNAP"); + if(snap && !strcmp(snap,g_mirror_dir)){ + fprintf(stderr,"[MIRROR] mirror dir equals the model dir — ignored\n"); return; + } + int nf=st_mirror_init(&m->S,g_mirror_dir); + if(nf<=0){ + fprintf(stderr,"[MIRROR] %s: no usable shard (missing or divergent copy) — " + "running on the primary drive only\n",g_mirror_dir); + return; + } + g_mirror=1; + double wp=0,wm=0; const char *w=getenv("COLI_DISK_WEIGHTS"); const char *how="COLI_DISK_WEIGHTS"; + if(w && (sscanf(w," %lf , %lf",&wp,&wm)!=2 || wp<=0 || wm<=0)){ + fprintf(stderr,"[MIRROR] invalid COLI_DISK_WEIGHTS '%s' (want e.g. 9,3) — probing instead\n",w); + wp=wm=0; + } + if(wp<=0||wm<=0){ + wp=mirror_probe_bw(&m->S,0); wm=mirror_probe_bw(&m->S,1); how="measured"; + if(wp>0&&wm>0) fprintf(stderr,"[MIRROR] probe: primary %.2f GB/s, mirror %.2f GB/s\n",wp,wm); + else { wp=wm=1; how="fallback 1:1 (probe failed)"; } + } + int share=(int)(256.0*wm/(wp+wm)+0.5); + g_mir_share = share<1?1 : share>255?255 : share; + fprintf(stderr,"[MIRROR] %s: %d/%d shards | read split primary %.0f%% / mirror %.0f%% (%s)\n", + g_mirror_dir,nf,m->S.nfd,100.0*(256-g_mir_share)/256,100.0*g_mir_share/256,how); +} + typedef struct { int l,e; uint32_t c; } PinRec; static int pin_rec_cmp(const void *a,const void *b){ const PinRec *x=a,*y=b; return x->cc?1:x->c>y->c?-1:0; @@ -5205,6 +5339,9 @@ int main(int argc, char **argv){ fprintf(stderr,"URING=1 is supported only on Linux\n"); return 2; #endif } + g_mirror_dir = getenv("COLI_MODEL_MIRROR"); /* DUAL-SSD: second model copy */ + if(!g_mirror_dir||!*g_mirror_dir) g_mirror_dir = getenv("SNAP_MIRROR"); + if(g_mirror_dir&&!*g_mirror_dir) g_mirror_dir = NULL; g_idot = getenv("IDOT")?atoi(getenv("IDOT")):1; /* 0 = kernel f32 esatti (A/B) */ g_spec_pin = getenv("SPEC_PIN")?atoi(getenv("SPEC_PIN")):1; /* #163: 0 = gate S-dipendenti storici / legacy S-dependent gates */ if(getenv("ROUTE_TRACE")&&*getenv("ROUTE_TRACE")){ @@ -5327,6 +5464,9 @@ int main(int argc, char **argv){ " Keep it on a native Linux fs (ext4/xfs/zfs) for memory efficiency and speed.\n", snap); } #endif + /* DUAL-SSD: register the mirror copy BEFORE any pin/autopin load, so the + * OMP-parallel pin warmup already streams from both drives. */ + mirror_setup(&m); /* HOT-STORE: PIN= [PIN_GB=g] -> top expert per frequenza fissi in RAM. * Va PRIMA di cap_for_ram: i pinnati contano nel residente. */ if(getenv("PIN")){ diff --git a/c/st.h b/c/st.h index a2c2c3f..8ed0388 100644 --- a/c/st.h +++ b/c/st.h @@ -40,6 +40,9 @@ typedef struct { int dfds[512]; /* gemelli O_DIRECT (aperti pigramente): -2 = non ancora provato */ char *paths[512]; int nfd; + int mfds[512]; /* MIRROR: fds of the second model copy (dual-SSD), -1 = absent */ + int mdfds[512]; /* O_DIRECT twins of the second copy, -1 = absent */ + int nmirror; /* files accepted into the mirror (0 = mirror inactive) */ int *hidx; /* hash map nome->indice (open addressing): con ~120k tensori * (GLM: 256 expert x 78 layer x 3 x 2) la scansione lineare * costava decine di secondi/token (misurato sul primo run reale) */ @@ -99,10 +102,80 @@ static int st_open_fd(shards *S, const char *path) { /* fd gemello O_DIRECT dello stesso file (bypassa la page cache: il buffered read su * ext4-in-VHDX si strozza a ~0.8 GB/s, O_DIRECT arriva a 2.3+; misurato). -1 se non disponibile. */ -static int st_direct_fd(shards *S, int fd) { - for (int i = 0; i < S->nfd; i++) if (S->fds[i] == fd) return S->dfds[i]; +static int st_fidx(shards *S, int fd) { + for (int i = 0; i < S->nfd; i++) if (S->fds[i] == fd) return i; return -1; } +static int st_direct_fd(shards *S, int fd) { + int i = st_fidx(S, fd); return i < 0 ? -1 : S->dfds[i]; +} + +/* ---- MIRROR (dual-SSD): second read-only copy of the model on another drive ---- + * st_fd_rep/st_direct_fd_rep: fd of replica `rep` (0 = primary, 1 = mirror) for + * the SAME file identified by its primary fd. -1 if that replica is absent. */ +static int st_fd_rep(shards *S, int fd, int rep) { + if (!rep) return fd; + if (!S->nmirror) return -1; + int i = st_fidx(S, fd); return i < 0 ? -1 : S->mfds[i]; +} +static int st_direct_fd_rep(shards *S, int fd, int rep) { + if (!rep) return st_direct_fd(S, fd); + if (!S->nmirror) return -1; + int i = st_fidx(S, fd); return i < 0 ? -1 : S->mdfds[i]; +} + +/* Registers / as a read replica of every already-indexed shard. + * A file is accepted ONLY if its size and safetensors header are byte-identical + * to the primary: the data_offsets then match by construction, so every pread + * is valid on either copy. Missing or divergent files simply stay on the + * primary (the mirror may be partial, e.g. a smaller SSD holding only the + * expert shards). Returns the number of accepted files. The mirror is NEVER + * written to: .coli_usage/.coli_kv keep deriving from the primary alone. */ +static int st_mirror_init(shards *S, const char *dir) { + if (S->nmirror) for (int i = 0; i < S->nfd; i++) { /* re-init: drop the old replica */ + if (S->mfds[i] >= 0) close(S->mfds[i]); + if (S->mdfds[i] >= 0) close(S->mdfds[i]); + } + for (int i = 0; i < ST_MAX_SHARDS; i++) { S->mfds[i] = -1; S->mdfds[i] = -1; } + S->nmirror = 0; + for (int i = 0; i < S->nfd; i++) { + const char *base = strrchr(S->paths[i], '/'); +#ifdef _WIN32 + const char *b2 = strrchr(S->paths[i], '\\'); + if (b2 && (!base || b2 > base)) base = b2; +#endif + base = base ? base + 1 : S->paths[i]; + char mp[2048]; snprintf(mp, sizeof(mp), "%s/%s", dir, base); + int mfd = open(mp, COMPAT_O_RDONLY); + if (mfd < 0) continue; /* partial mirror: this shard stays on the primary */ + int64_t sza = lseek(S->fds[i], 0, SEEK_END), szb = lseek(mfd, 0, SEEK_END); + if (sza != szb) { + fprintf(stderr, "[MIRROR] %s: size differs from the primary copy — file skipped\n", mp); + close(mfd); continue; + } + uint64_t ha = 0, hb = 0; int ok = 1; /* identical header => identical data_offsets */ + if (pread(S->fds[i], &ha, 8, 0) != 8 || pread(mfd, &hb, 8, 0) != 8 || + ha != hb || ha == 0 || ha > (uint64_t)256 << 20 || (int64_t)(8 + ha) > sza) ok = 0; + if (ok) { + char *ba = malloc(ha), *bb = malloc(ha); + if (!ba || !bb || pread(S->fds[i], ba, ha, 8) != (ssize_t)ha || + pread(mfd, bb, ha, 8) != (ssize_t)ha || memcmp(ba, bb, ha)) ok = 0; + free(ba); free(bb); + } + if (!ok) { + fprintf(stderr, "[MIRROR] %s: header differs from the primary copy — file skipped\n", mp); + close(mfd); continue; + } + S->mfds[i] = mfd; +#ifdef O_DIRECT + S->mdfds[i] = open(mp, COMPAT_O_RDONLY | O_DIRECT); +#elif defined(__APPLE__) + S->mdfds[i] = compat_open_direct(mp); +#endif + S->nmirror++; + } + return S->nmirror; +} /* indicizza tutti i model-*.safetensors in snap_dir */ /* pread completo: chunk-loop (una singola pread si ferma a ~2^31 byte su Linux @@ -256,6 +329,16 @@ static void st_prefetch(shards *S, const char *name) { if (t) posix_fadvise(t->fd, t->off, t->nbytes, POSIX_FADV_WILLNEED); } +/* like st_prefetch, but on replica `rep`'s drive: the WILLNEED must warm the + * page cache of the SAME fd the later demand pread will hit. */ +static void st_prefetch_rep(shards *S, const char *name, int rep) { + st_tensor *t = st_find(S, name); + if (!t) return; + int fd = st_fd_rep(S, t->fd, rep); + if (fd < 0) fd = t->fd; + posix_fadvise(fd, t->off, t->nbytes, POSIX_FADV_WILLNEED); +} + /* legge un tensore in un buffer float32 fornito dal chiamante (numel float). * drop=1 -> consiglia al kernel di scartare le pagine (per gli expert in streaming). */ static int64_t st_read_f32(shards *S, const char *name, float *out, int drop) { diff --git a/c/tests/test_st_mirror.c b/c/tests/test_st_mirror.c new file mode 100644 index 0000000..769c99d --- /dev/null +++ b/c/tests/test_st_mirror.c @@ -0,0 +1,109 @@ +/* Dual-SSD mirror (st_mirror_init & friends): a second read-only model copy is + * accepted only when byte-identical in size and safetensors header, reads on + * either replica return the same bytes, and divergent/missing copies degrade + * to the primary instead of being trusted. Fixture dirs are created in the + * current working directory and removed on exit. */ +#include +#include +#include + +#include "../st.h" + +#ifdef _WIN32 +#include +#define MKDIR(p) _mkdir(p) +#else +#include +#define MKDIR(p) mkdir(p, 0777) +#endif + +#define CHECK(condition) do { \ + if (!(condition)) { \ + fprintf(stderr, "%s:%d: check failed: %s\n", __FILE__, __LINE__, #condition); \ + return 1; \ + } \ +} while (0) + +#define DIR_A "tmp_mirror_a" /* primary */ +#define DIR_B "tmp_mirror_b" /* identical copy */ +#define DIR_C "tmp_mirror_c" /* same size, divergent header */ +#define DIR_D "tmp_mirror_d" /* different size */ +#define DIR_E "tmp_mirror_e" /* empty (missing file) */ + +/* one-tensor safetensors file; flip lets us corrupt one header byte and + * pad lets us grow the payload, both without changing anything else */ +static int write_model(const char *dir, int flip, int pad) { + char hdr[128], path[256]; + int hl = snprintf(hdr, sizeof(hdr), + "{\"t0\":{\"dtype\":\"F32\",\"shape\":[8],\"data_offsets\":[0,32]}}"); + if (flip) hdr[hl - 2] = ' '; /* inside the JSON header, same length */ + snprintf(path, sizeof(path), "%s/model.safetensors", dir); + FILE *f = fopen(path, "wb"); + if (!f) { perror(path); return -1; } + uint64_t n = (uint64_t)hl; + fwrite(&n, 8, 1, f); + fwrite(hdr, 1, (size_t)hl, f); + float data[8] = {1, -2, 3.5f, 0, 42, -0.5f, 7, 8}; + fwrite(data, 4, 8, f); + for (int i = 0; i < pad; i++) fputc(0, f); + fclose(f); + return 0; +} + +static void cleanup(void) { + const char *dirs[] = {DIR_A, DIR_B, DIR_C, DIR_D, DIR_E}; + for (int i = 0; i < 5; i++) { + char path[256]; + snprintf(path, sizeof(path), "%s/model.safetensors", dirs[i]); + remove(path); + remove(dirs[i]); + } +} + +int main(void) { + cleanup(); + CHECK(MKDIR(DIR_A) == 0 && MKDIR(DIR_B) == 0 && MKDIR(DIR_C) == 0 && + MKDIR(DIR_D) == 0 && MKDIR(DIR_E) == 0); + CHECK(write_model(DIR_A, 0, 0) == 0); + CHECK(write_model(DIR_B, 0, 0) == 0); + CHECK(write_model(DIR_C, 1, 0) == 0); + CHECK(write_model(DIR_D, 0, 64) == 0); + + shards S; + st_init(&S, DIR_A); + CHECK(S.n == 1 && S.nfd == 1); + st_tensor *t = st_find(&S, "t0"); + CHECK(t != NULL); + + /* without a mirror: replica 0 is the identity, replica 1 is absent */ + CHECK(st_fd_rep(&S, t->fd, 0) == t->fd); + CHECK(st_fd_rep(&S, t->fd, 1) == -1); + + /* identical copy: accepted, and both replicas serve the same bytes */ + CHECK(st_mirror_init(&S, DIR_B) == 1); + int mfd = st_fd_rep(&S, t->fd, 1); + CHECK(mfd >= 0 && mfd != t->fd); + float a[8], b[8]; + CHECK(pread(t->fd, a, t->nbytes, t->off) == t->nbytes); + CHECK(pread(mfd, b, t->nbytes, t->off) == t->nbytes); + CHECK(memcmp(a, b, sizeof(a)) == 0); + st_prefetch_rep(&S, "t0", 1); /* smoke: WILLNEED on the mirror fd */ + st_prefetch_rep(&S, "t0", 0); + + /* divergent header, same size: rejected */ + CHECK(st_mirror_init(&S, DIR_C) == 0); + CHECK(st_fd_rep(&S, t->fd, 1) == -1); + + /* different size: rejected */ + CHECK(st_mirror_init(&S, DIR_D) == 0); + + /* missing file: rejected (partial mirror with zero shards) */ + CHECK(st_mirror_init(&S, DIR_E) == 0); + + /* unknown fd never maps to a replica */ + CHECK(st_fd_rep(&S, 987654, 1) == -1); + + cleanup(); + puts("safetensors mirror tests: ok"); + return 0; +} From e9225d2ffc67c177dfac03c42bf5e32480abe77d Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Sun, 19 Jul 2026 15:02:14 +0200 Subject: [PATCH 2/3] docs: dual-SSD streaming env vars in ENVIRONMENT.md --- docs/ENVIRONMENT.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 9f0d749..394a58f 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -78,6 +78,15 @@ Format: `VAR` — default — effect. --- +## Dual-SSD streaming + +| Variable | Default | Effect | +|---|---|---| +| `COLI_MODEL_MIRROR` | unset | Path to a second, byte-identical (read-only) copy of the model on another drive; expert reads are split across both. Partial mirrors work (only the shards present are used). | +| `COLI_DISK_WEIGHTS` | unset (startup bandwidth probe) | Split ratio `,` (e.g. `1,1` for 50/50, `9,3` for a fast+slow pair). Unset = probe both drives with the engine's own access pattern at startup. | + +Per-drive byte counts are reported in a `MIRROR:` stats line. Combine with `DIRECT=1` so the two copies never compete for page cache. + ## CUDA (NVIDIA) | Variable | Default | Effect | From 507c8a1808dd3505f86231527181c75433f4dec2 Mon Sep 17 00:00:00 2001 From: JustVugg Date: Mon, 20 Jul 2026 21:58:15 +0200 Subject: [PATCH 3/3] fix(ci): remove accidentally-committed test_st_pread binary MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit #421 committed the compiled tests/test_st_pread (mode 644, non-executable). On a fresh CI checkout make sees it as up-to-date and skips rebuilding, so run_tests.py hits 'Permission denied' exec'ing a non-executable file — the linux/macos 'make check' failure. Remove the binary and gitignore it next to the existing test_st_mirror rule. Co-Authored-By: Claude Opus 4.8 --- c/.gitignore | 1 + c/tests/test_st_pread | Bin 30536 -> 0 bytes 2 files changed, 1 insertion(+) delete mode 100644 c/tests/test_st_pread diff --git a/c/.gitignore b/c/.gitignore index 2ea32c9..51695f9 100644 --- a/c/.gitignore +++ b/c/.gitignore @@ -4,3 +4,4 @@ olmoe_hf/ olmoe_i4/ .build-config tests/test_st_mirror +tests/test_st_pread diff --git a/c/tests/test_st_pread b/c/tests/test_st_pread deleted file mode 100644 index 8cf27103acc8a9513c7adc9fd46169663577a4dd..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 30536 zcmeHw3wTu3x%Qq(2$4(nP*EdR*~VZ)4KV?v2@=f23GBfML@r8I9Fmy{Nl7x(xp1*n zg2^cR9|x$XdhGedL+_sYXgO`6h=nHH)z(IP!{TWp72N}(s4a3)=Dgq9Yt8H#GyVIV z=YO96dHyF0X1@Ks>-)a-t#4iS+G|htUf1B7W3^Z$k@BS35^;eU22#f}Hu(7gNS#zI zjm7sw=^Cj3d_KdB{5pdm*GNaKnY57Od7z|MOb0{I&l|8LsPK>|>E%kxZAOlu66Z;8 z1Z64og`}|W6dhM$@PeB4Sb0rcK8r5p6%y2}N3m(U$Ixr*$fQs6dLPN-2DqOTF49}W z_117bK~GafkOUR+Bp>Nh&FyB>N>V}+G+VEY>t$1ip(Y7RWl&mlJna3Sl-I=dwkm!VtldRnaFYHm?c%D!4*FU~s^odKq=bc`$ytrWG zO?!ejoMiDNze$H=D3Lurgcwi9)P^L(;$c}&xRtOfd|vgv(pmFi?DGT;m9B>3=Z46b zodd6i>`?MA=D_dHVdoX#$`Iu~m%|Rkhb)((GL%1Zj&f(@z`u~gPD>8_+d1&Y9Qfrq z%C+UdJ96NT9QNPHffwemvp5GH&tc~TALPKt z=D@ph;1xN_U6wxM8Y8{;0p%*LCNh7g@c~JTDPZdC1V0+ zrrEQSGTg1r?Hy8U(C3pN?TrK^PlrF`^R@?x4|>`d4g0-Pt3TlDl0s`kVP7Yh;7X}; zr5DEmDIAP+d78sM6ygc54fxzG?hd~PT&J&76bWPaFQI^`Qqb4j!32D3+QSfQ4Fub} z!mZ5DPM+n)nZtt$GzY`(&gS+mDF7`meDb1lRM47ccWZlBb4UAq%yXm4Btx!y3@fdy zE)v`b`KX}wkl!R>x61!)f|>2{-tE_tc7M>7(@F^fZ`0oLX^H8m>s=)2b!6b8({^ z_4Tc2N6>?d8*lCKclj1Kw{%dbS9JQjc!1o7=#YZwAXc3J@^OyM$8?dO`d5%TN@Mwu zv`8b7LmKo?(6kIm5Jyt^QX%^$oz)l%bl#a$1&7LN-@rlXZZ!v zeD=JtR<$VdD>7*l%#puNK1S!&C>>2 z!kXr3<1F#vr>XiY#2-uZv{9CLEzQ#gSz>RRr;V}1PtrVXh$X(0=4m4=@o1W-4Y0(6 zX`VK|62UZ28(xW)G*27ViG^vNHog+|DLy%=s8Nzy<%tf7_21;Ahi=Q{AIRcYW$|5E zd}|heZx(+?7C%3WpOeMU&f;fe@fBJ8^;!HiS^VW$d{Gu(ki}n|ml=;US^Rrh{M%Xl zpR@QsWbvMi7T=Y{w`TG8X7P7q z@$<9zIa&PdEPh6cS7W>T4kVMh1GCXDX`a>Z!jT%U?ZK|BUV-DpN24%r$9C%^c6CZ@ znj!WeiH%I5lrNM&zyLQ)QocN~8`jzQaK#sw7OFjO4ce2zvrAO%xcyH19rh*m#rB1Z z)b;Puy)w0DOsVt_j_cslfph*5wb%>xsf1wTqjpZ! zwhSU1uK4zC=Ma@EK1n7IeCi*u06W%@Qx8S!4knXyM@^02bSEqv_;g*ABuQu>+u?y#O%>2r|b z(qNS->iwF1oEcFFvbUtr<||Td~G?w{VfFy=t`HEY-G4rrj4@$1T(d9L9;1@JtJ>?7Rv?UgSP!Nd}ye83w3RtOMT(Es_mt?K)9-}=Y&;z>=Q`X z<4Y_x%(}esb}S#5b87qva=tM>eWXi!)&>o7RohaCBd7KTb9(aLSdTv8h?mbo<05qW^Qcv{?}6JfMyXToVgb_5=x+i?7a=ncgE0D%J_c3W zNs&@x+gTUZ-gnmQcgnZ!S0{g@#ur&st=~2IO+5@{=lcE3R?YXR+{h!YnBjEH@VMuI z#Z~hxRb17M$k9s>JT*SzwmE2?7z+~|rb$DMvFNz8?N1Yjnz$@yk+%(8Ri9lu*g#R< zjzt)TprNpZx?`i~?S`sTCqKu|t4@^GSo5yj8;9Jy*e_1n&@WxuMg0W~8}ypLG9``4Dy}vua{c%9{7cRk9r<~ zu-yAL9N;c#2VL4RGzLlxqr~T_M12|+6^l|a2hP3UhsLCFj^0rREEl0p5^Uf-i-`49kC3C^!+|5#R@o3eC2!n{MKjD%%c480 z+U@}q=g_vF1vz;+I$1B7mz+KOEe$ot<&7;QAKOCmnA8q_fw)`C=Ed>{jR{1dqASi) zV}xHqmNG9keNaWq_AbNt>6wl47qfwMZ)NnTyivumE4Hdk?f+{YTG^?6d~zJE|IyBc zu9}l_?@^4D+hVgSP*^Jp(|?N@oZ8tLo2#JRQE(aDr!mUaQ~kuykEQj?N&jE6^=Fke zRAFFzzYr!cfsKbyqxOlaZI=KzwWRQ8KkWTg?!C^icRh``xrVL3VhoWmX2wSC9rO`0 zrN-u<>f@`?mgF33fxB@5(o>x$M;GEijiNU=3ArU-j@Hq^x=M?@aR$qKP41mS&`ji& z5`yZgFefMx`xZ5t2L~D%8y)5u&2KA1v)!m zhWMHqo2M{FjnC>o`2~u1c)Yq31DDb|^;_!47uA|0U%yK2=`U2}AOBgcITvnFW0w7I zUMLR~qBYn4Mvc{~_0)DzYVoQ*HGY|n^|u;})6m^Rr%L@xh>vmRAB!V4epOGr5JU6k zIGt0_JLv=@_s&PBirXzNWB$~lB%;RiotPx@)%f_45TsHt8B%EV_s@_nJ4X2wdi0Lw zNhn>j8q;RuKBiAnPkbBN8k)t;g)4k48Ek`SjL#wnC)s*C@0$!^}bFxuK zBTLopDq|xOBW%-Cbf%0M<1og?q5jH845ArC-%H2q*^ElFGOVz(IVWSx(O^d_h7#i;|;90X`72$HS!_gXL<1Ih86KA3K{maQw1H5TgUHLNJVh5^Ff@Rc|oX zmvYnsOKRLVpdGo}ewY0o`!f5zcK0%8yyfgZzP{zN0QykXYl*Ws*$7;_Sd~otF962* z8*4_S`_o~3|k?K%2 zERnUnPR~Jm&q0eTzLgf$^R#asr*n6VEuXY+Mu~yNM_>_WCEGj6WWCn=eL@o@ILq+0 zRALqD3bos75h~M^^m}YEWK2olgB>Gk_u(kX)+dcxKcajeRMGfHp}{TKmxz!XlbXIF zrCf_6{pr(GuUb14*<}KghKatb(L6V&WXItMViTBaaauAoYJ+G;!zh<$x_UuM>5UJV z((zd-{~`EC7{{6&oij_BXwSs@t1&}&1Mei_5GeRA}RK=sE_bQDXs{qlx) zU?8@;1444|Q#jWDgo3Kh>>^jN=G0#R8)HUx(KO`3!c|#gnMS(GPIB6Y5NiEGRaFF%Q3o+L7&v6fC-a>qlgOzsd2D1|Duj`2Y9Kkd2R> zh32IZj|*Q&j&(s2fxR2V2rHGBxP;BXwQYvVoya7W8bjquL!~UGayzN~0RtDyR<@Kj zE>NCdDjMM~h_b7XJp|Ed<-Jpo23_?YHAd|r>M#>Wq=~L1S+i+V!6upi)lIe>b)w~ z*gbEZZd%-co`pqz?Ru1uyzExk{##lBcsmd^KB9$4&Dro}*s8vsjh{7?SJnUwf}6MYB{LBdJ!wC*oePZ8wLuKgT{O+^l3l4yLK#erN|9S z)SlYcAc-0Acz6tA^#?Hef1E^Pc;V~85y!z>mO)oC>nBztll!GcUz0aJ2NU)?8?{6B zJDl1X?4)R)+7~;uUtw?M_lv5I(v|SZ@rECA?}JpHWuIDe^8T-=+KHVwdE$~~%;z22 z`|9LBIpfo>QET?9@`3@qj@mrhxAtZI4!9Jn?n&l{59{{@Y0;aj8j<|3!ohm#bV z|H@+CmYBA0f{-ywjm8oDryc*hmx)v6zJ10$f!a^8N#fLAuj*TZxzR9l9=p*-mpJ{d zFmez_hX1bE!^PO4aB3f52^j}1K-UVrC<$d@*`n5b7@nZUaKS!67T+a{;pccx zn~hPqHt|s zoveC(Tel%R-WFeE)lO=^mv`8(e7h5R;e{%8TH-5VZG5S`qq461;ysar-D^-Q6w+ih z;D7~qWZiT4zWafvhbLCe?CocJfWS?od6aRL@Zuk6*wRqcJf@dT=fHKt9~5~}tmtV#3}3A+9*hL4BTm{B*=itIa>q15&J71UVO&Wawk zZx8ldq{KuTCqHkp%0I8m$3_>f%L@)VN=Hep+TzlJUD!atk)yOwpQ~dj8yjEh)DCF< z4e{+#8F^8kZ(n~giA5zjw|ZgsefjsuybtvVd(cDba+Z z?)yeWh2={cn%5#~&$> z(2dAxM7@A2#w)XFER1U}sh*GePkgTDrPJ%NvXi6@i>KS<-p|8h$~qPQnzjGLn4V*| z^t{x){=%nBpd11h-~OE2{aN5FIm&Gmt?c21t**O+5tbd$HAxfw^Q6#xC{=CYr{sqqMB5 zPu_7uUGJ&zJ^BLlOINJSsU19dMQZ-qhr#lB3_~VUT_^82U_aw%uqctY>bgJ07@vtI z?>Tih7-jX9m{lq;pF#@5E&oTVWnea)(lGuo$P`)#K64yhF6Jmz_Fz7j*vxa_!Z?fy z>Ku9_{628O2KEpw=r9xxA-4`!*YMOd42w`Grub@CdJNj#`|L4SseeuKcnNr5wVaKZw!4NoL5!Z`CmoVfN z@|fzh9lVBCCE5XMS7TDFbjD_!9M8)&CNCr73VFv%XB_iU>O0-%XQrn%vO@Pe8)E)S zT>98+{tM?}xYQV%Tjq={P@J9vcB2A)sy4o#N@~#jmCUHzdjR9FF;+q6FxmxVlO@mk z5wcJwylq6G$GXp#%h7Qx!fiYpb(Vw;Lqf&SONX3#kp9)TlF37~uBKK!R7Za~DkZLi zmn2IP*^N9I64c0PJ7T`IyA#*J8WcyQGg%>HvI;>NQeqDc{2}_rkCB<9Z|whTs&Ak# zk;nSAFrC#Ku1-9IOxC|M`h(p2FNCN1#GA90$pbdl-LO?tS?gU zcngbjr^jE3QxJ6yTG1Iho#i_92HZ}tYj+wQBD|Bj-(|-9meC)=Pnr8)dE`Nnd-BQd z^B4tTnj2~Hm|p){jyhxOu#@U2Eymoh_P?E{YB>Gszk~P4LDj&~;3JVjeK%k=wz#y| zg{rV|i2DVJTC^uKq;MPhW2BBznz$LIrG5j_BYyRP+6|}S2n>Yj1coCp z9D(5o3`bx%0>csb|0)6&wrwOylR`6ge#6v+o<}P-$5k_Jt zYV!xf3O%t#X$~t53+DVIHRV|MHx==m9k0^b-r>6iroL7ZE}8kY5-&YU)PFo8`U$Sl#ySo${Drg z_5G}Fdr0Z>hm~e!-l7Hb6@SZDe4a2Xkm6{RqkLgf=nG0)a|r%+;Hg8MKCi;6d!xB( zH?m5PV(~ZYnPrr_S&PuuGQ}SWEB;m%N2SFd>7p1=j;Vdjc9?m^zseVE?eMP-DR^Ab zjOnaRg1**BsJ&~2Qr6}Vg{k`fpyKfdgONa(nP(PJ+}-WVl-gRQEYcNfU(tnHco6!@ zW$nrrl`knbDAVwpRw?x)r83qo&#SM-! zdj3!(G+DX+dgV@n7A>|fUa}|!L!O=|V?2GS3{NmBSI|L3$19nvTwANmU*d96;Yb;B ziy|J6FBF0htQbP*!Eh(+W)EYUi7G_8Rw^BSe*iUT^LHUvbW%JaNny`qnyL6#N@1z9 zJ;cIfM3qL#DC3b%=ya=7Pm-Dxy0O#m^>tL3I-0b)1?_s{q>%BjC8Nu7-8S?-$#`Vc z5Pee!dFjzkoxX50>ke*-^&iRHHzd};QvIB&8Pc3-QzbM~BoIKDL~k=3rs{`o%m_k- zhdSLPF9pNrS*g(Dw0z*Ge$+OiC9e;)@YbW}MwQW5t2c#smD$v=>6Vdr@TFj~cO zm(lC+WT#Q%iYHz_W%|wa&Zax>^>u{qTfbq`<6FM>%uk;0JM_x$-}vi^56&hF#*V+L zWOC(~W;eJN-PPRYUv+0_pm3UO0x#hPBxvQY*>8rWI?w^|r+<@72C&Jzd?1-D z1MM43Cf5Mp_9nJHLAPS_Zz=YQH)8`e0J??pr8X*4QjpjTu?E8F_@` zcE3~yTi27IV?#-b*eCqF*Hrw%cCA5^mLFCsr_ei@JY7b9jb(?&C_ZM4E_JP+0%$ww&< zjX-HV1?rg5+q&~_A3M@fG|-)&S6~@cROu)xvzvVwEqP(W3wuvsJp7Gz%3*I4+1p6= z;`#NddfBO7j-r6|l~G02sA+b6M(6zzMO5GRnVg2vFzdkQS170Q75o~2D&ib7N^;o> ztgEDXa^9+?1vc5bYK^VPx+-8BYhBf58)IFyyu+4fUDcG=J!%w$8`1LW=t`*4&6zdm zTkSZO3}upI#+aUAf!QLfBsQFeBQP9+;RuKbi02xL=NMvYVu_x;NNJ2gBuOz2=qDMJ z^6L$XpWd(}o|kwP=jmQCCH#hkCGp%uI#*E=`G3Ed^b^00=5QpsflX-x7r>u8uylmy z8&?l_b{;<`(VR-D*dUVm8F>V8HL0!35D=%bLYm`|==?&7e$7P*H~m>!$}2ALZeEb^ z`v{k#-^)-E<>S|SED1T>o@DhyI2nXHoQ#*aJ>2FK`AiSLjb!|jyrTGP7{-fwaP$Ay zh)6JZ2<#ZdT zeVh((YLxQNe~agc&b{^4nM&D`mPl7PqEt<*m{w6aWk!UtQy-W*y`pk@Mb%`U@sF%q zkeM%ikXk1TS?b616xK(0M^ktP{fYcg^49^k;%^+6;|F<^*jf;@3LnCostJ{(_YERB zK?$#;UvMI^H5=%c@gdw+ODKL~%EzS-#}>>UbpZzB%NS9eI3oG2lMI}#eQ>xMpP~GG zP{{N5CKy`y9Sg??vhaW9xR?)x{IfV9f6BN&6h}(W<71I7m&E+WKzt42!{Wf>Bk*@| zV8JXU=1YN(hQXoy&sDDpkhfs=TD(+PFla72lslCDGG^xrDSIBMVimkp67K;Jr8t;; zacVvl_+7v=o^#4G=ogbi;or!Ce6g&KtlL;e?4^j()62I6I~KkCAychn-Q@ zq3vIl1FryXgMC|;{hM;gt2yv{2)B+)oliv-T5`zWp99~V1OIjoe0vUjMh^UN4*YfC zjXrq|sy`d>Ui(=}>mQ4f&yP`eUr2@KG%D zS#RdB^Y^eJA8-#|nmHKIa9pYg{Nz%OS7kz~|<`znlZ_$br*y`G<<< z13Bcs0eqbGv*UUT!#~^4|1*c3mvhJu0#`&3P4s*YJ7aV7!;&2MG~gvByW(&@$y>3g zRRm(t#X0Qwa^P!n;G2OD`Tl@Ko?y7D!jJvZ=9YGMxOs(?v2}*sPOBrr>!<}wCzaN{*@W$(Rujuf%GLU(g%P}B#@(RuT-C8qa?f=wsJFY^3+BvO)UepS*k12yfJWBK2kKI93cz#zGu}Ip z&cWLN=*0{vfL;~g&VDgMDtFp+WM;mjAeF~7|Iyn7=uHe6()1#U3;->Ow?Je-ZXfNT zOU7FkGIIDE6*4f^8dTS?`PCCCk?H)^5-EUurq@Me0Pup|Ns)mWuXo6RXnQ>aaMMjN zYESnn^h@*xbUOZ?iwrprdvQbtgsw$5(cpG4OgBjIGKj4F^lL6s8Pm{q=`vgKehj0~ z?AJE9y&=E5t+~tFfmb>>7oZexdzU+c7jUHR1*L2n&5t3*->H$xpu+f@FH#r@8#hKW zAc)bcH!?8udqPqfMg>|zAzpF*u8&l{VIS}1;4jh0$Y8Gu$v})(l4L^Uat30&a3ce9 zyYbcwG3d-~h`Gsl4TkCEAt}wzMqh~27#Q+38E9B^wv+{`9w2rHD+6RN{750wrlFP1 zuOdn1u*RZEg}o{z1BPAv&ww&soRZ3!1{dgcCn$vdN6GO^g@0>!@JQK%l);6iayS58O7!D8X_4Op6!A?;^VM}@DA&%$lE zNtfab8^&~OXozDJ)!f+*V}?5Vr-V~TXD6oj%roWxjHCT9EJ1KBXryASvnUBF_QMR+ zkQK*7ufvD-poD%k*B7*sgiI-O6(~t`d280seve$KNtX_>49Q4z-EP(w`vZb18PcEC zzZ3a*`h;KjFZTHaE#vxPpD)}0X5`aWnb1GN>o4dw?lAdCaTfYwKaTbXNJZ$2eMLdV zJ|aX|A{*xPRWP(iC-lX>oS-AmHAU(EFlF0C)@6d>QM zpW4S*ZfJ;dvh@W&2FhehlrOHsZOaV_CFQZ;4Jy9J@$xGLz^Txml7*A6l*}}%iCKQ# zaz;}3X|4uK_j*LV#r%5-FyXtQVSFEMGw64-`mZQg$d1FoQ2HzR{T$i{HK#0A`ftF! BPSgMZ