diff --git a/c/Makefile b/c/Makefile index 7196ccac0..f5505c9d0 100644 --- a/c/Makefile +++ b/c/Makefile @@ -897,19 +897,34 @@ olmoe$(EXE): olmoe.c st.h json.h compat.h sample.h tok.h tok_unicode.h tok_unico $(CC) $(NOCUDA_CFLAGS) olmoe.c -o olmoe$(EXE) $(NOCUDA_LDFLAGS) # Qwen3.6-35B-A3B engine (hybrid Gated Attention + Gated DeltaNet + streaming -# MoE). CPU-only in this target; the optional CUDA expert tier is a separate -# follow-up (see docs/qwen36-phase01.md). -# NOCUDA_*: this file contains no CUDA, so building it with -DCOLI_CUDA and -# linking -lcudart only made `make qwen36 CUDA=1` depend on a toolkit it -# never calls. Same shape as olmoe. (The CUDA expert tier is #713, which -# adds its own sources and takes the CUDA flags back.) -qwen36$(EXE): qwen36.c st.h json.h compat.h - $(CC) $(NOCUDA_CFLAGS) qwen36.c -o qwen36$(EXE) $(NOCUDA_LDFLAGS) +# MoE). With CUDA=1 the optional VRAM expert tier (qwen36_tier.c) is compiled +# in and reuses the shared CUDA backend; without it the engine is CPU-only +# (qwen36_tier.h provides inline stubs). +# +# The flags follow the same switch. Without CUDA=1 this target keeps the +# NOCUDA_* flags it got when the engine had no CUDA at all: compiling +# -DCOLI_CUDA and linking -lcudart for a build that calls neither only made +# `make qwen36` depend on a toolkit it never touches. With CUDA=1 the tier is +# real CUDA and takes the normal flags back. +ifeq ($(CUDA),1) +QWEN36_TIER_SRC = qwen36_tier.c +QWEN36_CFLAGS = $(CFLAGS) +QWEN36_LDFLAGS = $(LDFLAGS) +else +QWEN36_TIER_SRC = +QWEN36_CFLAGS = $(NOCUDA_CFLAGS) +QWEN36_LDFLAGS = $(NOCUDA_LDFLAGS) +endif +qwen36$(EXE): qwen36.c qwen36_tier.h st.h json.h compat.h $(QWEN36_TIER_SRC) $(CUDA_OBJ) + $(CC) $(QWEN36_CFLAGS) qwen36.c $(QWEN36_TIER_SRC) $(CUDA_OBJ) -o qwen36$(EXE) $(QWEN36_LDFLAGS) # Context-size gates: KV layout, growth across requests, and the attention # capacity. Includes qwen36.c directly, so no model file is needed. -tests/test_qwen36_ctx$(EXE): tests/test_qwen36_ctx.c qwen36.c st.h json.h compat.h - $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +# Same tier sources as the engine: the test includes qwen36.c, so with CUDA=1 +# it needs qwen36_tier.c and the backend object too (without CUDA, the header's +# inline stubs cover it and both are empty). +tests/test_qwen36_ctx$(EXE): tests/test_qwen36_ctx.c qwen36.c qwen36_tier.h st.h json.h compat.h $(QWEN36_TIER_SRC) $(CUDA_OBJ) + $(CC) $(CFLAGS) $< $(QWEN36_TIER_SRC) $(CUDA_OBJ) -o $@ $(LDFLAGS) inkling$(EXE): inkling.c st.h json.h compat.h $(INK_CUDA_OBJ) $(METAL_OBJ) $(CC) $(CFLAGS) inkling.c $(INK_CUDA_OBJ) $(METAL_OBJ) -o inkling$(EXE) $(LDFLAGS) diff --git a/c/qwen36.c b/c/qwen36.c index 8213c2286..6689fd988 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -59,6 +59,7 @@ static int qwen36_max_ctx(void) { #endif #include "st.h" #include "json.h" /* tokenizer.json parsing (reuse minimal parser) */ +#include "qwen36_tier.h" /* optional transparent Vulkan compute backend for MoE experts */ #ifdef _WIN32 #include @@ -583,7 +584,7 @@ typedef struct { } Layer; /* ---------- LRU expert cache (int8 weights + per-row float scales) ---------- */ -typedef struct { int eid; int pinned; int is_int4; int8_t *g, *u, *d; float *gs, *us, *ds; uint64_t used; } Slot; +typedef struct { int eid; int pinned; int is_int4; int8_t *g, *u, *d; uint8_t *g4, *u4, *d4; float *gs, *us, *ds; uint64_t used; } Slot; typedef struct { Slot *slots; int n, cap; } LCache; typedef struct { @@ -651,6 +652,7 @@ static double g_tm_dec[6], g_tm_pre[6]; /* 0=deltanet 1=attention 2=moe_total static long g_tm_dec_tokens = 0, g_tm_pre_tokens = 0; static double tm_now(void){ struct timespec ts; clock_gettime(CLOCK_MONOTONIC,&ts); return ts.tv_sec*1e3 + ts.tv_nsec/1e6; } static int tm_on(void){ if(g_timers<0){ const char *e=getenv("COLI_TIMERS"); g_timers = (e && *e=='1'); } return g_timers; } +double g_qt_iss=0, g_qt_cpu=0, g_qt_tak=0; /* QTIER-Phasen (Decode) */ double g_dn_sub[4]; /* DN: proj, conv+split, l2n+rec, norm+out */ double g_tm_step=0; /* step() total (decode) */ static double g_tm_win_moe=0; static int g_tm_win_n=0; @@ -682,6 +684,9 @@ static void tm_report(void){ if(g_dn_sub[0]+g_dn_sub[1]+g_dn_sub[2]+g_dn_sub[3]>0) fprintf(stderr,"[timers] dn-sub: proj %.1f | conv %.1f | l2n+rec %.1f | norm+out %.1f ms/token\n", g_dn_sub[0]/g_tm_dec_tokens,g_dn_sub[1]/g_tm_dec_tokens,g_dn_sub[2]/g_tm_dec_tokens,g_dn_sub[3]/g_tm_dec_tokens); + if(g_qt_iss+g_qt_cpu+g_qt_tak>0) + fprintf(stderr,"[timers] qtier: issue %.2f | cpu-miss %.2f | take %.2f ms/token\n", + g_qt_iss/g_tm_dec_tokens, g_qt_cpu/g_tm_dec_tokens, g_qt_tak/g_tm_dec_tokens); fprintf(stderr,"[timers] prefill: %ld tokens dn=%.0f attn=%.0f moe=%.0f(sh=%.0f rt=%.0f) head=%.0f ms\n", g_tm_pre_tokens,g_tm_pre[0],g_tm_pre[1],g_tm_pre[2],g_tm_pre[3],g_tm_pre[4],g_tm_pre[5]); } @@ -1197,6 +1202,7 @@ static void slot_ensure_allocated(Model *m, Slot *s) { s->ds = s_block + 2*scale_count_gu(c); s->pinned = 0; s->is_int4 = 0; + s->g4 = s->u4 = s->d4 = NULL; /* packed int4 (allocated on int4 load if GPU int4 active) */ } static void load_expert_merged(Model *m, int layer, int eid, Slot *s) { @@ -1234,14 +1240,29 @@ static void load_expert_merged(Model *m, int layer, int eid, Slot *s) { s->g[i] = v; } s->is_int4 = 1; - /* The packed int4 bytes are NOT kept here. This engine only ever reads - * the unpacked int8 copy above, so retaining them doubled expert-cache - * RSS on the recommended gs64 container for nothing. They are the upload - * source for the CUDA expert tier and come back with it (#713), which is - * where they are actually read. */ + /* Free any previous occupant first (LRU slot reuse). */ + free(s->g4); free(s->u4); free(s->d4); s->g4 = s->u4 = s->d4 = NULL; + /* Keep the packed int4 bytes alongside the unpacked int8 copy only when + * the CUDA tier is actually running: they are its upload source, and + * they let slot_ensure_int8() rematerialize an evicted expert whose + * int8 copy the warmstart freed. Without the tier nothing ever reads + * them, and keeping them would add ~50% to expert-cache RSS on the + * recommended gs64 container -- so qt_ready() gates the allocation. + * Under CUDA=0 that is an inline `return 0` and this costs nothing. */ + if (qt_ready()) { + int64_t gp = ng / 2, up = ng / 2, dp = nd / 2; /* gate/up/down packed sizes */ + s->g4 = (uint8_t *)malloc((size_t)gp); + s->u4 = (uint8_t *)malloc((size_t)up); + s->d4 = (uint8_t *)malloc((size_t)dp); + if (!s->g4 || !s->u4 || !s->d4) { fprintf(stderr, "OOM int4-packed %s\n", nm); exit(1); } + memcpy(s->g4, raw, (size_t)gp); + memcpy(s->u4, raw + gp, (size_t)up); + memcpy(s->d4, raw + gp + up, (size_t)dp); + } free(raw); } else { s->is_int4 = 0; + free(s->g4); free(s->u4); free(s->d4); s->g4 = s->u4 = s->d4 = NULL; st_read_raw(&m->S, nm, s->g, 1); } st_read_f32(&m->S, qsnm, s->gs, 0); @@ -1262,6 +1283,31 @@ static int container_is_int4(Model *m) { return (tw->nbytes == want_w / 2) ? 1 : 0; } +/* Rematerialize a slot's int8 block from its packed int4 copy on demand + * (~0.5 ms, no container access). Needed after the warmstart freed the int8 + * copies of VRAM-resident experts and one of them got LFRU-evicted. */ +static void slot_ensure_int8(Model *m, Slot *s) { + if (s->g || !s->g4) return; + Cfg *c = &m->c; + int64_t ng = (int64_t)c->inter * c->hidden, nd = (int64_t)c->hidden * c->inter; + int8_t *w = malloc((size_t)(ng + ng + nd)); + if (!w) { fprintf(stderr, "OOM slot_ensure_int8\n"); exit(1); } + const uint8_t *src4[3] = { s->g4, s->u4, s->d4 }; + int64_t lens[3] = { ng, ng, nd }; + int8_t *dst = w; + for (int t = 0; t < 3; t++) { + const uint8_t *p = src4[t]; + for (int64_t i = 0; i < lens[t]; i += 2) { + uint8_t b = p[i >> 1]; + int8_t lo = (int8_t)(b & 0xF); if (lo & 8) lo -= 16; + int8_t hi = (int8_t)((b >> 4) & 0xF); if (hi & 8) hi -= 16; + dst[i] = lo; dst[i + 1] = hi; + } + dst += lens[t]; + } + s->g = w; s->u = w + ng; s->d = w + ng + ng; +} + static void expert_get(Model *m, int layer, int eid, Slot **out) { LCache *lc = &m->cache[layer]; pthread_mutex_lock(&g_pilot_mx); @@ -1572,9 +1618,59 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out) { for (int kk = 0; kk < K; kk++) if (idx[kk] >= 0) freq_l[idx[kk]]++; } const float *xs = x + (int64_t)s*D; - { + int shared_done = 0; + if (qt_ready()) { + /* CUDA expert tier: run the resident experts as async groups on + * all devices, compute the misses on the CPU (overlapped), then + * collect the GPU results. */ + for (int kk = 0; kk < K; kk++) { + Slot *e; expert_get(m, layer, idx[kk], &e); + if (e->g4) qt_note(layer, idx[kk], e->g4, e->u4, e->d4, e->gs, e->us, e->ds); + } + double _q0 = tm_on()? tm_now():0; + uint32_t qmask = qt_issue(layer, idx, K, xs); + double _q1 = tm_on()? tm_now():0; + for (int kk = 0; kk < K; kk++) { + if (qmask & (1u<g, e->gs, D, I); + matmul_qe(u, xs, e->u, e->us, D, I); + for (int i = 0; i < I; i++) { float gv = g[i]; g[i] = (gv / (1.f + expf(-gv))) * u[i]; } + matmul_qe(hh, g, e->d, e->ds, I, D); + float w = val[kk]; float *os = out + (int64_t)s*D; + for (int d = 0; d < D; d++) os[d] += w * hh[d]; + } + /* Compute the shared expert NOW so it overlaps with the GPU + * groups; the common block below is skipped. */ + { + double _ts2 = tm_on() ? tm_now() : 0.0; + int Ish = c->shared_inter; + matmul_d(sh, xs, l->sh_g, 1, D, Ish); + matmul_d(shu, xs, l->sh_u, 1, D, Ish); + for (int i = 0; i < Ish; i++) { float sv = sh[i]; sh[i] = (sv / (1.f + expf(-sv))) * shu[i]; } + matmul_d(shd, sh, l->sh_d, 1, Ish, D); + float sgate = 1.f; + if (l->sh_gate) { + float sg = 0.f; const float *wg = l->sh_gate; + for (int i = 0; i < D; i++) sg += xs[i] * wg[i]; + sgate = 1.f / (1.f + expf(-sg)); + } + float *os = out + (int64_t)s*D; + for (int d = 0; d < D; d++) os[d] += sgate * shd[d]; + if (tm_on()) tm_add(S, 3, tm_now()-_ts2); + } + shared_done = 1; + double _q2 = tm_on()? tm_now():0; + qt_take(qmask, val, K, out + (int64_t)s*D); + if (tm_on() && S==1) { + extern double g_qt_iss, g_qt_cpu, g_qt_tak; + g_qt_iss += _q1-_q0; g_qt_cpu += _q2-_q1; g_qt_tak += tm_now()-_q2; + } + } else { for (int kk = 0; kk < K; kk++) { Slot *e; expert_get(m, layer, idx[kk], &e); + slot_ensure_int8(m, e); matmul_qe(g, xs, e->g, e->gs, D, I); matmul_qe(u, xs, e->u, e->us, D, I); for (int i = 0; i < I; i++) { float gv = g[i]; g[i] = (gv / (1.f + expf(-gv))) * u[i]; } @@ -1585,6 +1681,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out) { } } /* shared expert (SwiGLU), sigmoid-gated by shared_expert_gate */ + if (shared_done) { if (0) goto shared_skip_dummy; shared_skip_dummy: continue; } double _ts = tm_on() ? tm_now() : 0.0; int Ish = c->shared_inter; matmul_d(sh, xs, l->sh_g, 1, D, Ish); @@ -1919,9 +2016,14 @@ static void pilot_prefetch(Model *m, int lnext, const float *x, int S) { if (idx[b] >= 0 && (idx[a] < 0 || idx[a] > idx[b])) { int t = idx[a]; idx[a] = idx[b]; idx[b] = t; } for (int kk = 0; kk < cand; kk++) { int eid = idx[kk]; if (eid < 0) continue; - int found = 0; pthread_mutex_lock(&g_pilot_mx); LCache *lc = &m->cache[lnext]; - for (int z = 0; z < lc->n; z++) if (lc->slots[z].eid == eid) { found = 1; break; } + int found = 0, fz = -1; pthread_mutex_lock(&g_pilot_mx); LCache *lc = &m->cache[lnext]; + for (int z = 0; z < lc->n; z++) if (lc->slots[z].eid == eid) { found = 1; fz = z; break; } pthread_mutex_unlock(&g_pilot_mx); + /* Lookahead: RAM-resident layer-L+1 candidates go to VRAM asynchronously */ + if (found && fz >= 0 && qt_ready()) { + Slot *ps = &lc->slots[fz]; + if (ps->g4) qt_note(lnext, eid, ps->g4, ps->u4, ps->d4, ps->gs, ps->us, ps->ds); + } if (!found) { int gidx = lnext*E + eid; pthread_mutex_lock(&g_pilot_mx); int already_queued = m->is_queued[gidx]; @@ -2307,7 +2409,55 @@ int main(int argc, char **argv) { g_qdw_n, now_s()-tq, freed/1073741824.0); } - /* coli serve mode: speak the gateway wire protocol instead of argv generation */ + /* Optional CUDA VRAM expert tier (COLI_CUDA=1): hot experts live in + * DEVICE_LOCAL memory across the configured GPUs, misses fall back to the + * CPU int8 path. See qwen36_tier.h. */ + if (qt_init(m.c.n_layers, m.c.n_experts, m.c.hidden, m.c.inter, cap, m.c.topk, m.c.expert_gs)) { + fprintf(stderr, "[gpu] MoE experts -> CUDA VRAM tier\n"); + atexit(qt_shutdown); + /* Warmstart: fill the VRAM budget BEFORE the first token (heat order + * when HEAT_FILE exists, natural order otherwise), loading all RAM + * slots along the way. */ + const char *nws = getenv("QT_NO_WARMSTART"); + if (!(nws && *nws=='1')) { + /* Plan the set (heat order), then load+stage IN PARALLEL. The + * load path is thread-safe: expert_get locks the layer cache + * (g_pilot_mx), st_read_raw uses pread; entries are unique. */ + double t0 = now_s(); + int cap_total = m.c.n_layers * m.c.n_experts; + int *wpl = malloc((size_t)cap_total*sizeof(int)); + int *wpe = malloc((size_t)cap_total*sizeof(int)); + int wn = qt_plan_fill(wpl, wpe, cap_total); + /* Load ALL experts into RAM, not just the planned (VRAM) set: + * otherwise the first touch of a CPU-fallback expert triggers a + * ~12 ms container read in the middle of decode (measured: 139 + * ms/token on a single-GPU run). Planned ones also go to VRAM. */ + uint8_t *planned = calloc((size_t)cap_total, 1); + for (int i = 0; i < wn; i++) planned[wpl[i]*m.c.n_experts + wpe[i]] = 1; + int keep8 = getenv("COLI_KEEP_INT8") != NULL; + #pragma omp parallel for schedule(dynamic, 16) + for (int gi = 0; gi < cap_total; gi++) { + int l = gi / m.c.n_experts, eidw = gi % m.c.n_experts; + Slot *e; expert_get(&m, l, eidw, &e); + if (planned[gi] && e->g4) { + qt_note_planned(l, eidw, e->g4, e->u4, e->d4, e->gs, e->us, e->ds); + /* The staging copy is done; free the int8 copy RIGHT AWAY + * so it never shows up in peak RSS. On LFRU eviction + * slot_ensure_int8() rematerializes from g4 (no container + * access). */ + if (!keep8 && e->g) { free(e->g); e->g = e->u = e->d = NULL; } + } + } + qt_fill_wait(); + free(wpl); free(wpe); free(planned); + fprintf(stderr, "[qtier] warmstart (parallel): all %d experts in RAM (int8 only for non-residents), %d in VRAM -- %.1f s\n", + cap_total, wn, now_s()-t0); + } + } + + /* coli serve mode: speak the gateway wire protocol instead of argv + * generation. AFTER the tier init: serve sessions ride the VRAM experts + * exactly like argv runs, and serve_loop never returns. */ if (getenv("SERVE") && getenv("SERVE")[0] == '1') { if (!g_tok) { fprintf(stderr, "[serve] tokenizer.json required (put in SNAP or set TOK)\n"); return 1; } serve_loop(&m); @@ -2383,6 +2533,7 @@ int main(int argc, char **argv) { double tot = m.hits + m.miss; if (g_ttft >= 0) fprintf(stderr, "TTFT: %.2f s (time to first token)\n", g_ttft); tm_report(); + qt_stats(); fprintf(stderr, "\nPEAK RSS: %.2f GB\n", rss_gb()); fprintf(stderr, "Expert cache hit rate: %.1f%% (hit=%llu miss=%llu)\n", tot?100.0*m.hits/tot:0.0, (unsigned long long)m.hits, (unsigned long long)m.miss); diff --git a/c/qwen36_tier.c b/c/qwen36_tier.c new file mode 100644 index 000000000..7639c8441 --- /dev/null +++ b/c/qwen36_tier.c @@ -0,0 +1,446 @@ +/* qwen36_tier.c -- CUDA VRAM expert tier for the qwen36 engine. See header. */ +#ifdef COLI_CUDA +#include +#include +#include +#include +#include "qwen36_tier.h" +#include "backend_cuda.h" + +#define QT_MAX_DEV 8 +#define QT_QCAP 48 /* upload queue depth (staging ~1.6 MB/entry) */ + +typedef struct { + ColiCudaTensor *tg, *tu, *td; + uint32_t heat; + uint8_t resident, queued, planned; + /* raw RAM pointers (slots are never evicted when cap==n_experts) -- lets + * warmstart, lookahead and LFRU swaps run without an engine callback */ + const uint8_t *g4,*u4,*d4; const float *gs,*us,*ds; +} QSlot; + +static struct { + int on, nl, ne, D, Ih, topk, ndev; + int egs; size_t sc_gu, sc_d; /* expert group size + per-matrix scale counts (gs64) */ + int dev[QT_MAX_DEV]; + size_t budget[QT_MAX_DEV], used[QT_MAX_DEV]; + size_t exp_bytes; /* estimated VRAM bytes per expert */ + QSlot *slot; /* [nl*ne] */ + pthread_mutex_t mx; + pthread_t th; + int th_stop; + /* upload ring with staging copies */ + struct { int layer, eid; uint8_t *w; float *s; int v_layer, v_eid; } q[QT_QCAP]; + int qh, qt_, qn; + pthread_cond_t cv; + /* statistics */ + uint64_t hits[QT_MAX_DEV], miss, uploads, q_full_skips; + /* issue state of the (single) decode thread */ + int is_cnt[QT_MAX_DEV]; + int is_k[QT_MAX_DEV][32]; + float *is_x; /* count*D input replicas per device */ + /* M3 */ + int *fill_order; int fill_cur; /* warmstart order (heat desc) */ + int issue_open; /* guard: no tensor_free while a group is in flight */ + pthread_cond_t cv_take; /* signals qt_take done + queue space */ + uint64_t tick, swaps, pf_hits, pf_notes; + uint32_t *heat0; /* heat table loaded from HEAT_FILE */ +} G; + +static QSlot *qs(int layer, int eid){ return &G.slot[(size_t)layer*G.ne + eid]; } +static int home(int eid){ return eid % G.ndev; } + +/* Staging: packed int4 (g|u|d) two's-complement -> offset-binary (XOR 0x88, + * the upload format of backend_cuda fmt=2) + copy the scales (gs|us|ds). */ +static void stage(uint8_t *dw, float *dsc, + const uint8_t *g4,const uint8_t *u4,const uint8_t *d4, + const float *gs,const float *us,const float *ds){ + size_t mb = (size_t)G.D*G.Ih/2; + const uint64_t X=0x8888888888888888ull; + const uint64_t *sg=(const uint64_t*)g4,*su=(const uint64_t*)u4,*sd=(const uint64_t*)d4; + uint64_t *w0=(uint64_t*)dw,*w1=(uint64_t*)(dw+mb),*w2=(uint64_t*)(dw+2*mb); + for(size_t i=0;i=0){ + /* LFRU swap: free the victim only when no group is in flight */ + while(G.issue_open && !G.th_stop) pthread_cond_wait(&G.cv_take,&G.mx); + QSlot *v=qs(vl,ve); + ColiCudaTensor *a=v->tg,*b=v->tu,*ct=v->td; + v->tg=v->tu=v->td=NULL; + pthread_mutex_unlock(&G.mx); + if(a)coli_cuda_tensor_free(a); if(b)coli_cuda_tensor_free(b); if(ct)coli_cuda_tensor_free(ct); + } else pthread_mutex_unlock(&G.mx); + + int dv = G.dev[home(eid)]; + size_t mb=(size_t)G.D*G.Ih/2; + ColiCudaTensor *tg=NULL,*tu=NULL,*td=NULL; + int ok; + if(G.egs){ + ok = coli_cuda_tensor_upload_g(&tg, w, sc, 4, G.D, G.Ih, dv, G.egs) + && coli_cuda_tensor_upload_g(&tu, w+mb, sc+G.sc_gu, 4, G.D, G.Ih, dv, G.egs) + && coli_cuda_tensor_upload_g(&td, w+2*mb, sc+2*G.sc_gu, 4, G.Ih, G.D, dv, G.egs); + } else { + ok = coli_cuda_tensor_upload(&tg, w, sc, 2, G.D, G.Ih, dv) + && coli_cuda_tensor_upload(&tu, w+mb, sc+G.Ih, 2, G.D, G.Ih, dv) + && coli_cuda_tensor_upload(&td, w+2*mb, sc+2*G.Ih, 2, G.Ih, G.D, dv); + } + free(w); free(sc); + pthread_mutex_lock(&G.mx); + QSlot *s=qs(layer,eid); + if(ok){ s->tg=tg; s->tu=tu; s->td=td; s->resident=1; G.uploads++; } + else { int hd=home(eid); G.used[hd]-=G.exp_bytes; + G.budget[hd]=G.used[hd]; /* device genuinely full: stop trying */ } + s->queued=0; + pthread_mutex_unlock(&G.mx); + } +} + +int qt_init(int nl, int ne, int D, int Ih, int cap, int topk, int expert_gs){ + const char *e=getenv("COLI_CUDA"); + if(!(e && *e=='1')) return 0; + if(cap != ne){ + fprintf(stderr,"[qtier] cap=%d != n_experts=%d -> tier disabled (needs full RAM residency)\n",cap,ne); + return 0; + } + if(topk>32){ fprintf(stderr,"[qtier] topk>32 unsupported\n"); return 0; } + memset(&G,0,sizeof G); + G.nl=nl; G.ne=ne; G.D=D; G.Ih=Ih; G.topk=topk; + + /* devices: COLI_GPUS="0,1" (default: 0,1 when present, else 0) */ + const char *gl=getenv("COLI_GPUS"); + char buf[128]; snprintf(buf,sizeof buf,"%s", gl?gl:"0,1"); + for(char *t=strtok(buf,","); t && G.ndev CPU path\n"); return 0; } + int have=coli_cuda_device_count(); + if(have CPU path\n"); return 0; } + + /* per-device budget: CUDA_EXPERT_GB, or auto = free minus 1 GB headroom. + * Scale counts follow the container: per-row (expert_gs=0) or grouped + * (gs64: [O, ceil(I/gs)] per projection). */ + G.egs = expert_gs; + G.sc_gu = expert_gs ? (size_t)Ih * ((D + expert_gs - 1)/expert_gs) : (size_t)Ih; + G.sc_d = expert_gs ? (size_t)D * ((Ih + expert_gs - 1)/expert_gs) : (size_t)D; + G.exp_bytes = 3ull*D*Ih/2 + (2*G.sc_gu+G.sc_d)*sizeof(float) + 4096; /* + allocation slack */ + const char *bg=getenv("CUDA_EXPERT_GB"); + for(int i=0;i0) + ? (size_t)(atof(bg)*1024.0*1024.0*1024.0) + : (freeb>(1ull<<30) ? freeb-(1ull<<30) : 0); + G.budget[i]=b; + fprintf(stderr,"[qtier] dev %d: %.1f GB free, budget %.1f GB (~%zu experts)\n", + G.dev[i], freeb/1073741824.0, b/1073741824.0, b/G.exp_bytes); + } + G.slot=calloc((size_t)nl*ne,sizeof(QSlot)); + G.is_x=malloc((size_t)32*D*sizeof(float)); + if(!G.slot||!G.is_x) return 0; + /* load learned heat (HEAT_FILE): warmstart order + initial values */ + const char *hf=getenv("HEAT_FILE"); + if(hf){ + FILE *f=fopen(hf,"rb"); + if(f){ + uint32_t hdr[3]={0,0,0}; + if(fread(hdr,4,3,f)==3 && hdr[0]==0x51544831u && hdr[1]==(uint32_t)nl && hdr[2]==(uint32_t)ne){ + G.heat0=malloc((size_t)nl*ne*4); + if(G.heat0 && fread(G.heat0,4,(size_t)nl*ne,f)==(size_t)nl*ne){ + for(size_t i=0;i<(size_t)nl*ne;i++) G.slot[i].heat=G.heat0[i]>>1; /* decay */ + fprintf(stderr,"[qtier] HEAT_FILE loaded: %s\n",hf); + } else { free(G.heat0); G.heat0=NULL; } + } + fclose(f); + } + } + pthread_mutex_init(&G.mx,NULL); pthread_cond_init(&G.cv,NULL); pthread_cond_init(&G.cv_take,NULL); + if(pthread_create(&G.th,NULL,uploader,NULL)!=0) return 0; + G.on=1; + fprintf(stderr,"[qtier] CUDA VRAM expert tier active: %d device(s), %.2f MB/expert\n", + G.ndev, G.exp_bytes/1048576.0); + return 1; +} + +int qt_ready(void){ return G.on; } + +/* Is (layer,eid) currently VRAM-resident? (used to free RAM-side int8 copies) */ +int qt_is_resident(int layer,int eid){ + if(!G.on) return 0; + pthread_mutex_lock(&G.mx); + int r = qs(layer,eid)->resident; + pthread_mutex_unlock(&G.mx); + return r; +} + +/* internal, G.mx held: enqueue one upload. victim=-1: plain upload (budget is + * reserved here); victim>=0: LFRU swap (budget neutral). */ +static int enqueue_locked(int layer,int eid,int v_layer,int v_eid,int reserved){ + QSlot *s=qs(layer,eid); + if(s->resident||s->queued||!s->g4) return 0; + if(G.qn>=QT_QCAP){ G.q_full_skips++; return 0; } + int hd=home(eid); + if(!reserved && v_eid<0 && G.used[hd]+G.exp_bytes>G.budget[hd]) return 0; + size_t mb=(size_t)G.D*G.Ih/2; + uint8_t *w=malloc(3*mb); float *sc=malloc((2*G.sc_gu+G.sc_d)*sizeof(float)); + if(!w||!sc){ free(w); free(sc); return 0; } + if(!reserved && v_eid<0) G.used[hd]+=G.exp_bytes; + s->queued=1; + stage(w,sc,s->g4,s->u4,s->d4,s->gs,s->us,s->ds); + G.q[G.qt_].layer=layer; G.q[G.qt_].eid=eid; G.q[G.qt_].w=w; G.q[G.qt_].s=sc; + G.q[G.qt_].v_layer=v_layer; G.q[G.qt_].v_eid=v_eid; + G.qt_=(G.qt_+1)%QT_QCAP; G.qn++; + pthread_cond_signal(&G.cv); + return 1; +} + +void qt_note(int layer,int eid, + const uint8_t *g4,const uint8_t *u4,const uint8_t *d4, + const float *gs,const float *us,const float *ds){ + if(!G.on || !g4) return; + QSlot *s=qs(layer,eid); + pthread_mutex_lock(&G.mx); + if(!s->g4){ s->g4=g4; s->u4=u4; s->d4=d4; s->gs=gs; s->us=us; s->ds=ds; } + if(s->heat<0xFFFFFFFFu) s->heat++; + enqueue_locked(layer,eid,-1,-1,0); + pthread_mutex_unlock(&G.mx); +} + +/* blocking variant for the warmstart (waits for queue space). */ +void qt_note_block(int layer,int eid, + const uint8_t *g4,const uint8_t *u4,const uint8_t *d4, + const float *gs,const float *us,const float *ds){ + if(!G.on || !g4) return; + QSlot *s=qs(layer,eid); + pthread_mutex_lock(&G.mx); + if(!s->g4){ s->g4=g4; s->u4=u4; s->d4=d4; s->gs=gs; s->us=us; s->ds=ds; } + while(G.qn>=QT_QCAP && !G.th_stop) pthread_cond_wait(&G.cv_take,&G.mx); + enqueue_locked(layer,eid,-1,-1,0); + pthread_mutex_unlock(&G.mx); +} + +/* warmstart order -- heat descending (HEAT_FILE) or natural order. + * Returns 0 once all budgets are full or the list is exhausted. */ +static const uint32_t *g_sort_heat; +static int cmp_heat_desc(const void *a,const void *b){ + uint32_t ha=g_sort_heat[*(const int*)a], hb=g_sort_heat[*(const int*)b]; + return hahb ? -1 : 0; +} +int qt_fill_next(int *layer,int *eid){ + if(!G.on) return 0; + size_t n=(size_t)G.nl*G.ne; + pthread_mutex_lock(&G.mx); + if(!G.fill_order){ + G.fill_order=malloc(n*sizeof(int)); + for(size_t i=0;iresident||s->queued) continue; + if(G.used[hd]+G.exp_bytes>G.budget[hd]) continue; /* dieses Device voll */ + *layer=l; *eid=e; + pthread_mutex_unlock(&G.mx); + return 1; + } + pthread_mutex_unlock(&G.mx); + return 0; +} + +/* Plan the whole warmstart set in one pass -- same heat order and budget + * reservation as qt_fill_next, but without loading. The experts are then + * loaded by any number of threads and handed over via qt_note_planned. */ +int qt_plan_fill(int *layers,int *eids,int max){ + if(!G.on) return 0; + size_t n=(size_t)G.nl*G.ne; + int cnt=0; + pthread_mutex_lock(&G.mx); + if(!G.fill_order){ + G.fill_order=malloc(n*sizeof(int)); + for(size_t i=0;iresident||s->queued||s->planned) continue; + if(G.used[hd]+G.exp_bytes>G.budget[hd]) continue; + G.used[hd]+=G.exp_bytes; /* reserve */ + s->planned=1; + layers[cnt]=l; eids[cnt]=e; cnt++; + } + pthread_mutex_unlock(&G.mx); + return cnt; +} + +/* Thread-safe (callable from multiple loader threads): stage + enqueue one + * expert reserved by qt_plan_fill; blocks only while the queue is full. */ +void qt_note_planned(int layer,int eid, + const uint8_t *g4,const uint8_t *u4,const uint8_t *d4, + const float *gs,const float *us,const float *ds){ + if(!G.on || !g4) return; + QSlot *s=qs(layer,eid); + pthread_mutex_lock(&G.mx); + if(!s->g4){ s->g4=g4; s->u4=u4; s->d4=d4; s->gs=gs; s->us=us; s->ds=ds; } + while(G.qn>=QT_QCAP && !G.th_stop) pthread_cond_wait(&G.cv_take,&G.mx); + if(!enqueue_locked(layer,eid,-1,-1,1)){ + /* not enqueueable (e.g. already resident): return the reservation */ + if(s->planned) G.used[home(eid)]-=G.exp_bytes; + } + s->planned=0; + pthread_mutex_unlock(&G.mx); +} + +/* waits until the upload queue is drained (end of warmstart). */ +void qt_fill_wait(void){ + if(!G.on) return; + pthread_mutex_lock(&G.mx); + while(G.qn>0 && !G.th_stop) pthread_cond_wait(&G.cv_take,&G.mx); + pthread_mutex_unlock(&G.mx); +} + +/* LFRU swap check (every 16 ticks = tokens): per device, coldest resident vs + * hottest non-resident, with the tier.h hysteresis. */ +static void qt_lfru_tick_locked(void){ + if(++G.tick % 16) return; + size_t n=(size_t)G.nl*G.ne; + for(int di=0;diresident && !s->queued){ if(cold<0||s->heatheat; } } + else if(!s->resident && !s->queued && s->g4){ if(hot<0||s->heat>hh){ hot=(int)i; hh=s->heat; } } + } + if(cold<0||hot<0) continue; + if(hh<=ch+(ch>>2)+4) continue; /* hysteresis as in tier.h */ + QSlot *v=&G.slot[cold]; + v->resident=0; /* CPU fallback from now on */ + if(enqueue_locked(hot/G.ne,hot%G.ne,cold/G.ne,cold%G.ne,0)) G.swaps++; + else v->resident=1; /* queue full: revert */ + } +} + +uint32_t qt_issue(int layer,const int *eids,int K,const float *x){ + if(!G.on||K>32) return 0; + uint32_t mask=0; + ColiCudaTensor *tg[QT_MAX_DEV][32],*tu[QT_MAX_DEV][32],*td[QT_MAX_DEV][32]; + static int rows[32]={0}; + if(!rows[0]) for(int i=0;i<32;i++) rows[i]=1; + for(int i=0;iresident){ + int di=home(eids[k]); int c=G.is_cnt[di]; + tg[di][c]=s->tg; tu[di][c]=s->tu; td[di][c]=s->td; + G.is_k[di][c]=k; G.is_cnt[di]=c+1; + mask|=1u< hand these k back to the CPU */ + for(int j=0;j0? 100.0*hits/tot : 0.0, (unsigned long long)G.swaps); + { uint64_t calls=0,ex=0,rows=0; double h2d=0,kms=0,d2h=0; + coli_cuda_group_stats(&calls,&ex,&rows,&h2d,&kms,&d2h); + if(calls) fprintf(stderr,"[qtier] group_stats: %llu calls, %llu experts | h2d %.0f ms, kernel %.0f ms, d2h %.0f ms\n", + (unsigned long long)calls,(unsigned long long)ex,h2d,kms,d2h); } +} + +void qt_shutdown(void){ + if(!G.on) return; + const char *hf=getenv("HEAT_FILE"); + if(hf){ + FILE *f=fopen(hf,"wb"); + if(f){ + uint32_t hdr[3]={0x51544831u,(uint32_t)G.nl,(uint32_t)G.ne}; + fwrite(hdr,4,3,f); + for(size_t i=0;i<(size_t)G.nl*G.ne;i++) fwrite(&G.slot[i].heat,4,1,f); + fclose(f); + fprintf(stderr,"[qtier] HEAT_FILE saved: %s\n",hf); + } + } + pthread_mutex_lock(&G.mx); G.th_stop=1; pthread_cond_signal(&G.cv); pthread_mutex_unlock(&G.mx); + pthread_join(G.th,NULL); + G.on=0; + coli_cuda_shutdown(); +} + +#endif /* COLI_CUDA */ diff --git a/c/qwen36_tier.h b/c/qwen36_tier.h new file mode 100644 index 000000000..e424a1687 --- /dev/null +++ b/c/qwen36_tier.h @@ -0,0 +1,81 @@ +/* qwen36_tier.h -- optional CUDA VRAM expert tier for the qwen36 engine. + * + * Applies colibri's placement concept ("route -> place -> overlap -> learn") + * one level up from the GLM disk tier: experts live in RAM, the *hot* ones + * are promoted into DEVICE_LOCAL VRAM across one or more GPUs and computed + * there via the existing CUDA backend (backend_cuda.cu, expert-group API). + * + * - Every expert has one home device (eid % n_gpus), no duplicates. + * - Routing heat decides who earns VRAM (LFRU semantics from tier.h, with + * hysteresis); a warmstart pre-fills the budget before the first token, + * ordered by a persisted heat table (HEAT_FILE) when available. + * - Uploads run on a background thread through staging copies; decode never + * blocks on placement. A VRAM miss falls back to the CPU int8 path and + * overlaps with the in-flight GPU groups. + * + * Enable with COLI_CUDA=1 [COLI_GPUS=0,1] [CUDA_EXPERT_GB=|auto] + * [HEAT_FILE=] [QT_NO_WARMSTART=1]. Compiled only when the build sets + * -DCOLI_CUDA (CUDA=1); otherwise the inline stubs below keep the engine + * CPU-only with zero overhead. */ +#ifndef QWEN36_TIER_H +#define QWEN36_TIER_H +#include + +#ifdef COLI_CUDA + +/* Init after model load. Returns 1 when the tier is active. + * cap_experts_per_layer must equal n_experts (full RAM residency): the tier + * stores raw pointers into the expert slots, which must never be evicted. */ +int qt_init(int n_layers, int n_experts, int hidden, int inter, + int cap_experts_per_layer, int topk, int expert_gs); +int qt_ready(void); +int qt_is_resident(int layer, int eid); +void qt_shutdown(void); + +/* Call once per routed expert per token (pointers to the RAM slot: packed + * int4 + per-row scales). Updates heat and may enqueue a background upload. */ +void qt_note(int layer, int eid, + const uint8_t *g4, const uint8_t *u4, const uint8_t *d4, + const float *gs, const float *us, const float *ds); + +/* Launch the GPU groups for the resident subset of the K selected experts + * (async, all devices in parallel). Returns a bitmask of the k handled by + * the GPU. Compute the misses on the CPU, then call qt_take(). */ +uint32_t qt_issue(int layer, const int *eids, int K, const float *x); + +/* Collect the GPU results and accumulate val[k]*y_k into out[hidden]. */ +void qt_take(uint32_t mask, const float *val, int K, float *out); + +/* Warmstart: plan the full fill set (heat order, budget reserved), then any + * number of loader threads may call qt_note_planned per planned expert. */ +int qt_plan_fill(int *layers, int *eids, int max); +void qt_note_planned(int layer, int eid, + const uint8_t *g4, const uint8_t *u4, const uint8_t *d4, + const float *gs, const float *us, const float *ds); +int qt_fill_next(int *layer, int *eid); +void qt_note_block(int layer, int eid, + const uint8_t *g4, const uint8_t *u4, const uint8_t *d4, + const float *gs, const float *us, const float *ds); +void qt_fill_wait(void); /* blocks until the upload queue is drained */ + +/* One telemetry block on stderr: residency, hits/misses, uploads per device. */ +void qt_stats(void); + +#else /* !COLI_CUDA: inline stubs, engine stays CPU-only */ + +static inline int qt_init(int a,int b,int c,int d,int e,int f,int g){(void)a;(void)b;(void)c;(void)d;(void)e;(void)f;(void)g;return 0;} +static inline int qt_ready(void){return 0;} +static inline int qt_is_resident(int a,int b){(void)a;(void)b;return 0;} +static inline void qt_shutdown(void){} +static inline void qt_note(int a,int b,const uint8_t*c,const uint8_t*d,const uint8_t*e,const float*f,const float*g,const float*h){(void)a;(void)b;(void)c;(void)d;(void)e;(void)f;(void)g;(void)h;} +static inline uint32_t qt_issue(int a,const int*b,int c,const float*d){(void)a;(void)b;(void)c;(void)d;return 0;} +static inline void qt_take(uint32_t a,const float*b,int c,float*d){(void)a;(void)b;(void)c;(void)d;} +static inline int qt_plan_fill(int*a,int*b,int c){(void)a;(void)b;(void)c;return 0;} +static inline void qt_note_planned(int a,int b,const uint8_t*c,const uint8_t*d,const uint8_t*e,const float*f,const float*g,const float*h){(void)a;(void)b;(void)c;(void)d;(void)e;(void)f;(void)g;(void)h;} +static inline int qt_fill_next(int*a,int*b){(void)a;(void)b;return 0;} +static inline void qt_note_block(int a,int b,const uint8_t*c,const uint8_t*d,const uint8_t*e,const float*f,const float*g,const float*h){(void)a;(void)b;(void)c;(void)d;(void)e;(void)f;(void)g;(void)h;} +static inline void qt_fill_wait(void){} +static inline void qt_stats(void){} + +#endif /* COLI_CUDA */ +#endif /* QWEN36_TIER_H */ diff --git a/docs/qwen36-cuda-tier.md b/docs/qwen36-cuda-tier.md new file mode 100644 index 000000000..21684a6ca --- /dev/null +++ b/docs/qwen36-cuda-tier.md @@ -0,0 +1,50 @@ +# qwen36: CUDA VRAM expert tier + +Applies colibri's placement concept ("route -> place -> overlap -> learn") to +Qwen3.6-35B-A3B one level up from the GLM disk tier: all 10,240 experts live +in RAM, the **hot** ones are promoted into DEVICE_LOCAL VRAM across one or +more GPUs and computed there through the existing shared CUDA backend +(`backend_cuda.cu` expert-group API — no new backend). + +## How it works + +- **Home device:** expert `eid` lives on GPU `eid % n_gpus`; no duplicates. +- **Placement:** routing heat decides who earns VRAM (LFRU semantics from + `tier.h`, 25%+4 hysteresis). A parallel **warmstart** fills the per-device + budget before the first token — ordered by a persisted heat table + (`HEAT_FILE`) when present, so a second run starts fully placed. +- **Decode:** per (token, layer) the resident experts are issued as async + groups on all devices (`coli_cuda_expert_group_issue/take`); VRAM misses + fall back to the CPU int8 path and overlap with the in-flight groups, as + does the shared expert. Placement never changes routing or precision. +- **Memory:** the warmstart frees the RAM int8 copies of VRAM-resident + experts (rematerialized from the packed int4 copy on LFRU eviction; no + container access). Peak RSS for the 35B int4 container: ~29 GB with two + 8 GB GPUs. + +## Usage + +```bash +make -C c qwen36 CUDA=1 CUDA_ARCH=native # NVCC=/usr/bin/nvcc on distro CUDA +COLI_CUDA=1 COLI_GPUS=0,1 HEAT_FILE=heat.bin CUDA_EXPERT_GB=auto \ +OMP_NUM_THREADS= OMP_WAIT_POLICY=ACTIVE OMP_PROC_BIND=close \ +SNAP= N_NEW=200 ./c/qwen36 256 4 prompt.txt +``` + +`cap` (argv[1]) must equal `n_experts` (full RAM residency). int4 containers +only (the int8 container keeps the CPU path). `COLI_TIMERS=1` prints +per-phase timings and tier telemetry. + +## Measured (Threadripper 3945WX 12C, RTX 3070 8 GB + Quadro RTX 4000 8 GB, Qwen3.6-35B-A3B int4, 200-token decode) + +| | 1 GPU (8 GB) | 2 GPUs (16 GB) | +|---|---|---| +| decode tok/s (cold / warm heat) | 9.2 / 9.9 | 10.6 / **11.3** | +| VRAM-resident experts | 4,391 (43 %) | 8,532 (83 %) | +| VRAM hit rate (cold / warm) | 44 % / 95 % | 85 % / 100 % | +| peak RSS | 40 GB | **29 GB** | +| reference: Ollama q4_K_M, same box | 7.5 | 10.5 | + +CPU-only baseline of this engine before the tier: 0.35 tok/s. +Numerics: logits cosine vs the f32 CPU reference 0.9992 (dense int8 on), +bit-identical GPU-vs-CPU on the same container (cosine 1.0000001).