From fc52d73ed8c0cb44ba11b1c347db8bd43b16572c Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Thu, 30 Jul 2026 09:26:31 +0200 Subject: [PATCH 01/17] feat(qwen36): add Qwen3.6-35B-A3B engine (CPU): hybrid Gated Attention + Gated DeltaNet + streaming MoE Self-contained engine for Qwen3.6-35B-A3B (35B total / 3B active, 256 experts per layer), split out of PR #602 per maintainer guidance (engine first, GPU backends as follow-ups; no shared GLM/Inkling files touched): - c/qwen36.c: 40-layer hybrid forward (10x Gated Attention with partial RoPE + per-head q/k RMSNorm; 30x Gated DeltaNet recurrent linear attention with carried conv ring + per-head state), streaming MoE with per-layer LRU expert cache, shared expert, router lookahead prefetch (PILOT), resident-expert pinning (COLIBRI_RESIDENT), per-phase timers (COLI_TIMERS=1). - CPU performance work: optional per-row int8 quantization of the large dense matrices (COLI_DENSE_I8, on by default, f32 originals freed), hand-written AVX2/FMA int8 GEMV, DeltaNet recurrence parallelized per value head. Decode on a 12-core Zen2 + int4 container: ~2.5 tok/s CPU-only (vs 0.35 tok/s scalar baseline). - c/tools/convert_qwen36.py: HF -> container converter (true int4 packing, per-row scales, --selftest). Containers now run out of the box: config.json is written flat (VL checkpoints nest dims under text_config) and tokenizer.json is bundled. - tools/make_qwen36_tiny.py + make_qwen36_oracle.py: weight-free tiny model and oracle for engine validation without the 70 GB checkpoint. - docs/qwen36-phase01.md, qwen36-phase02.md: design notes. Numerics: logits cosine vs f32 reference 0.9992 with COLI_DENSE_I8=1, bit-identical with COLI_DENSE_I8=0. Co-authored-by: minne100 --- c/Makefile | 6 + c/qwen36.c | 2009 +++++++++++++++++++++++++++++++++ c/tools/convert_qwen36.py | 572 ++++++++++ c/tools/make_qwen36_oracle.py | 139 +++ c/tools/make_qwen36_tiny.py | 158 +++ docs/qwen36-phase01.md | 97 ++ docs/qwen36-phase02.md | 103 ++ 7 files changed, 3084 insertions(+) create mode 100644 c/qwen36.c create mode 100644 c/tools/convert_qwen36.py create mode 100644 c/tools/make_qwen36_oracle.py create mode 100644 c/tools/make_qwen36_tiny.py create mode 100644 docs/qwen36-phase01.md create mode 100644 docs/qwen36-phase02.md diff --git a/c/Makefile b/c/Makefile index ee941025f..c6bfa7d65 100644 --- a/c/Makefile +++ b/c/Makefile @@ -876,6 +876,12 @@ fp8-bench: fp8_bench$(EXE) olmoe$(EXE): olmoe.c st.h json.h compat.h sample.h tok.h tok_unicode.h tok_unicode_o200k.h omp_tune.h route_trace.h serve_codec.h $(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). +qwen36$(EXE): qwen36.c st.h json.h compat.h + $(CC) $(CFLAGS) qwen36.c -o qwen36$(EXE) $(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 new file mode 100644 index 000000000..fa9b0702f --- /dev/null +++ b/c/qwen36.c @@ -0,0 +1,2009 @@ +/* Qwen3.6-35B-A3B inference engine in pure C, Phase 2: Gated Attention + Gated + * DeltaNet (recurrent linear attention) + streaming MoE. + * + * The full model is a hybrid: 10 x (3 x Gated DeltaNet -> MoE, 1 x Gated + * Attention -> MoE). Phase 1 implemented ONLY the 25% attention layers and + * treated the DeltaNet layers as identity; Phase 2 implements BOTH: + * - Gated Attention (GQA, per-head q/k RMSNorm, partial RoPE, output gate). + * - Gated DeltaNet: causal depthwise conv1d + recurrent gated-delta-rule with a + * carried conv ring + state S[h]=[kdim,vdim], then per-head Gated RMSNorm. + * Every layer (attention or DeltaNet) carries its own MoE/MLP block. + * + * DENSE (embed, attn/dn q/k/v/o & projections, q/k norms, RMSNorm, router gate, + * shared expert, lm_head, final norm) resident in RAM (float32). Expert weights + * read from disk on-demand via pread + posix_fadvise(DONTNEED), cached LRU + * per-layer, with a PILOT prefetch thread -- the same mechanism that fits + * GLM-5.2 in 15 GB. + * + * Env vars (inherited from olmoe.c): PILOT, HOT, WARMUP, WIDE, SMOOTH, CONF_LIMIT. + * Plus: SNAP=, and argv: qwen36 [ref.json] [PPL=1]. + */ +#define _GNU_SOURCE +#include +#include +#include +#include +#include +#include +#include +#if defined(__AVX2__) +#include +#endif +#if defined(__APPLE__) || defined(__linux__) || defined(__FreeBSD__) +#include +#include +#endif +#include "st.h" +#include "json.h" /* tokenizer.json parsing (reuse minimal parser) */ + +#ifdef _WIN32 +#include +#define sleep_ms(ms) Sleep(ms) +#else +#include +#define sleep_ms(ms) usleep((ms) * 1000) +#endif + +/* ---------- tokenizer (optional, for human-readable output) ---------- */ +static char **g_tok = NULL; /* id -> piece string (strdup'd) */ +static int g_tok_n = 0; + +static int hexnib(char c){ + if (c>='0'&&c<='9') return c-'0'; + if (c>='a'&&c<='f') return c-'a'+10; + if (c>='A'&&c<='F') return c-'A'+10; + return 0; +} + +/* ===== text -> ids : BPE encoder (mirrors HF/Qwen tokenizer.json) ===== + * Builds piece->id (reverse vocab) + pair->rank (merges) maps, plus the + * GPT-2 byte-to-unicode mapping. Encode = special-token split + GPT-2 regex + * pre-tokenize + per-piece ByteLevel map + BPE merges. */ +typedef struct { char **keys; int *vals; int *used; int cap; } SMap; +static unsigned shash(const char *s){ unsigned h=2166136261u; while(*s){ h^=(unsigned char)*s++; h*=16777619u; } return h; } +static void smap_init(SMap *m,int cap){ m->cap=cap; m->keys=calloc(cap,sizeof(char*)); m->vals=malloc(cap*sizeof(int)); m->used=calloc(cap,sizeof(int)); } +static void smap_put(SMap *m,const char *k,int v){ if(!k)return; unsigned h=shash(k)&(m->cap-1); while(m->used[h]){ if(m->keys[h]&&strcmp(m->keys[h],k)==0){m->vals[h]=v;return;} h=(h+1)&(m->cap-1);} m->used[h]=1; m->keys[h]=(char*)k; m->vals[h]=v; } +static int smap_get(SMap *m,const char *k){ if(!m||!m->cap||!k)return -1; unsigned h=shash(k)&(m->cap-1); while(m->used[h]){ if(m->keys[h]&&strcmp(m->keys[h],k)==0)return m->vals[h]; h=(h+1)&(m->cap-1);} return -1; } + +static SMap g_rev; /* piece string -> id (encode) */ +static SMap g_merge; /* "a\x1F b" pair -> rank (encode) */ +static char byte_sym_utf8[256][8]; /* byte -> UTF-8 of mapped codepoint */ +static short g_unmap[512]; /* mapped codepoint -> original byte (-1 = unused) */ +static int g_nspecial = 0; +static char **g_sp_str = NULL; static int *g_sp_id = NULL; static int *g_sp_len = NULL; + +static const char *jstr(jval *o,const char *k){ jval *v=json_get(o,k); return (v&&v->t==J_STR)?v->str:NULL; } +static double jnum(jval *o,const char *k){ jval *v=json_get(o,k); return (v&&v->t==J_NUM)?v->num:0; } + +enum { U_W=0, U_L=1, U_M=2, U_N=3, U_P=4, U_O=5 }; +static int uclass(unsigned cp){ + if (cp==0x20||cp==0x09||cp==0x0A||cp==0x0D||cp==0x0B||cp==0x0C) return U_W; + if (cp==0x00A0||cp==0x2000||cp==0x2001||cp==0x2002||cp==0x2003||cp==0x2004||cp==0x2005||cp==0x2006||cp==0x2007||cp==0x2008||cp==0x2009||cp==0x200A||cp==0x2028||cp==0x2029||cp==0x202F||cp==0x205F||cp==0x3000||cp==0xFEFF) return U_W; + if (cp>=0x30&&cp<=0x39) return U_N; + if (cp>=0xFF10&&cp<=0xFF19) return U_N; + if (cp>=0x0660&&cp<=0x0669) return U_N; + if ((cp>=0x41&&cp<=0x5A)||(cp>=0x61&&cp<=0x7A)) return U_L; + if (cp>=0x00C0&&cp<=0x024F) return U_L; + if (cp>=0x0400&&cp<=0x04FF) return U_L; + if (cp>=0x0600&&cp<=0x06FF) return U_L; + if (cp>=0x1F00&&cp<=0x1FFF) return U_L; + if (cp>=0x3040&&cp<=0x30FF) return U_L; + if (cp>=0x3400&&cp<=0x4DBF) return U_L; + if (cp>=0x4E00&&cp<=0x9FFF) return U_L; + if (cp>=0xAC00&&cp<=0xD7A3) return U_L; + if (cp>=0x300&&cp<=0x36F) return U_M; + if (cp>=0x1AB0&&cp<=0x1AFF) return U_M; + if (cp>=0x1DC0&&cp<=0x1DFF) return U_M; + if (cp>=0x20D0&&cp<=0x20FF) return U_M; + if (cp>=0xFE20&&cp<=0xFE2F) return U_M; + if (cp>=0x21&&cp<=0x2F) return U_P; + if (cp>=0x3A&&cp<=0x40) return U_P; + if (cp>=0x5B&&cp<=0x60) return U_P; + if (cp>=0x7B&&cp<=0x7E) return U_P; + if (cp>=0x3000&&cp<=0x303F) return U_P; + if (cp>=0xFF01&&cp<=0xFF0F) return U_P; + if (cp>=0xFF1A&&cp<=0xFF20) return U_P; + if (cp>=0xFF3B&&cp<=0xFF40) return U_P; + if (cp>=0xFF5B&&cp<=0xFF65) return U_P; + if (cp>=0x2010&&cp<=0x2027) return U_P; + if (cp>=0x2030&&cp<=0x205E) return U_P; + return U_O; +} +static int utf8_decode(const char *s,int i,int n,int *adv){ + unsigned char c=(unsigned char)s[i]; int cp,a; + if(c<0x80){cp=c;a=1;} + else if((c>>5)==6){cp=c&0x1F;a=2;} + else if((c>>4)==14){cp=c&0x0F;a=3;} + else if((c>>3)==30){cp=c&0x07;a=4;} + else {cp=c;a=1;} + for(int k=1;k>6); byte_sym_utf8[b][k++]=0x80|(c&0x3F); } + else { byte_sym_utf8[b][k++]=0xE0|(c>>12); byte_sym_utf8[b][k++]=0x80|((c>>6)&0x3F); byte_sym_utf8[b][k++]=0x80|(c&0x3F); } + byte_sym_utf8[b][k]=0; + g_unmap[cp]=(short)b; /* reverse: mapped codepoint -> original byte */ + } +} +static void push_id(int **ids,int *n,int *cap,int v){ if(*n==*cap){*cap*=2; *ids=realloc(*ids,*cap*sizeof(int));} (*ids)[(*n)++]=v; } + +static int try_special(const char *s,int i,int n,int *id_out){ + int best_len=0,best_id=-1; + for(int k=0;kn) continue; + if(memcmp(s+i,g_sp_str[k],L)==0){ if(L>best_len){best_len=L;best_id=g_sp_id[k];} } + } + *id_out=best_id; return best_len; +} +/* Pre-tokenize splitter, mirrors the HF/Qwen regex alternation: + * (?i:'s|'t|'re|'ve|'m|'ll|'d) | [^\r\n\p{L}\p{N}]?[\p{L}\p{M}]+ | \p{N} + * | ?[^\s\p{L}\p{M}\p{N}]+[\r\n]* | \s*[\r\n]+ | \s+(?!\S) | \s+ + * Returns the byte index just past the piece starting at i. */ +static int pretok_end(const char *s,int i,int n){ + if (s[i]=='\''){ + const char *cands[]={"ll","ve","re","s","t","m","d"}; int clen[]={2,2,2,1,1,1,1}; + int best=0; + for(int c=0;c<7;c++){ int L=clen[c]; if(i+1+L>n) continue; int ok=1; for(int k=0;kbest)best=L; } + if(best>0) return i+1+best; + } + int adv; unsigned c0=utf8_decode(s,i,n,&adv); + { /* rule2: optional non-(cr/lf/letter/number) prefix then letter/mark run */ + int k=i; unsigned c=c0; int prefix=0; + if(k1){ + int best=-1,besti=-1; + for(int k=0;k=0 && (best<0||r0){ push_id(&ids,&n,&cap,sid); i+=L; continue; } + int j=pretok_end(text,i,tlen); if(j<=i) j=i+utf8_adv(text,i); + bpe_piece(text+i,j-i,&ids,&n,&cap); + i=j; + } + *out_ids=ids; *out_n=n; +} + +/* Load Qwen tokenizer.json and build an id->piece table. Only needs the + * "model.vocab" map (piece string -> id); merges are irrelevant for decoding. */ +static void load_tokenizer(const char *path){ + FILE *f = fopen(path, "rb"); + if (!f) { fprintf(stderr, "[tok] cannot open %s\n", path); return; } + fseek(f,0,SEEK_END); long n = ftell(f); fseek(f,0,SEEK_SET); + char *buf = malloc(n+1); + if (fread(buf,1,(size_t)n,f) != (size_t)n) { /* ignore short read */ } + buf[n] = 0; fclose(f); + char *arena = NULL; + jval *root = json_parse(buf, &arena); + jval *model = json_get(root, "model"); if (!model) model = root; + jval *vocab = json_get(model, "vocab"); + if (!vocab) vocab = json_get(model, "tokens"); + if (!vocab) { fprintf(stderr, "[tok] no model.vocab/tokens in %s\n", path); free(buf); return; } + int mx = 0; + if (vocab->t == J_OBJ){ + for (int i=0;ilen;i++){ int id=(int)vocab->kids[i]->num; if(id>mx)mx=id; } + } else { + mx = vocab->len - 1; + } + g_tok = calloc((size_t)mx+1, sizeof(char*)); + if (vocab->t == J_OBJ){ + for (int i=0;ilen;i++){ int id=(int)vocab->kids[i]->num; if(id>=0 && id<=mx) g_tok[id]=strdup(vocab->keys[i]); } + } else { + for (int i=0;ilen;i++){ if(vocab->kids[i] && vocab->kids[i]->t==J_STR) g_tok[i]=strdup(vocab->kids[i]->str); } + } + g_tok_n = mx+1; + + /* ---- encoder tables (text -> ids) ---- */ + smap_init(&g_rev, 1<<19); + for (int i=0;it==J_ARR){ + for (int r=0;rlen;r++){ + const char *e = merges->kids[r]->str; if(!e) continue; + const char *sp = strchr(e, ' '); if(!sp) continue; + int la=(int)(sp-e), lb=(int)strlen(sp+1); + char *key=malloc(la+1+lb+1); + memcpy(key,e,la); key[la]=0x1F; memcpy(key+la+1,sp+1,lb); key[la+1+lb]=0; + smap_put(&g_merge, key, r); + } + } + jval *adds = json_get(root, "added_tokens"); + if (adds && adds->t==J_ARR && g_nspecial==0){ + g_nspecial = adds->len; + g_sp_str = malloc(g_nspecial*sizeof(char*)); + g_sp_id = malloc(g_nspecial*sizeof(int)); + g_sp_len = malloc(g_nspecial*sizeof(int)); + for (int k=0;klen;k++){ + jval *t = adds->kids[k]; + const char *c = jstr(t,"content"); + g_sp_str[k] = c?strdup(c):strdup(""); + g_sp_id[k] = (int)jnum(t,"id"); + g_sp_len[k] = (int)strlen(g_sp_str[k]); + } + } + build_byte_sym(); + + fprintf(stderr, "[tok] loaded %d pieces (max id %d) from %s\n", vocab->len, mx, path); + free(buf); +} + +/* Decode token ids to text using g_tok, writing to stdout. Handles Qwen's + * byte-representation markers (Ġ=space, Ċ=newline, ▁=space) and <0xXX> byte + * fallback. Only active when a tokenizer was loaded. */ +/* ---- streaming / incremental decode support ---- */ +static int g_stream = 0; /* 1 = emit tokens as they are generated */ +static unsigned char g_sbuf[16]; /* carries a partial UTF-8 char across tokens */ +static int g_sbn = 0; + +/* ---- OpenAI-compatible output + timing ---- */ +static int g_openai = 0; /* 1 = emit OpenAI Chat Completions format (SSE/JSON) */ +static double g_gen_t0 = 0; /* generate() start (monotonic seconds) */ +static double g_ttft = -1; /* time to first token (s); -1 = unset */ +static long g_oa_created = 0; /* unix timestamp for OpenAI "created" */ +static char g_oa_id[64]; /* OpenAI-style id, e.g. chatcmpl-... */ +static const char *g_model = "qwen3.6-35b-a3b-colibri"; +static double now_s(void); /* forward decl; defined later near model code */ + +/* Output sink for server mode: when g_sock_out >= 0, SSE/JSON bytes are routed + * to the live socket via g_sock_send instead of stdout. Lets qwen36_serve.c + * reuse all emit logic without any change to the CLI path. */ +static long long g_sock_out = -1; +static void (*g_sock_send)(long long fd, const char *buf, int n) = NULL; + +/* JSON-escape a byte string into out (no surrounding quotes). Returns length. */ +static int json_escape(const unsigned char *s, int n, char *out, int outsz){ + int o = 0; + for (int i=0;i *bn) break; + if (*outn + need <= 255){ for (int x=0;x` lines). */ +static void sse_chunk(const char *json){ + char hdr[8]; int hl = snprintf(hdr, sizeof hdr, "data: "); + if (g_sock_out >= 0 && g_sock_send){ + g_sock_send(g_sock_out, hdr, hl); + g_sock_send(g_sock_out, json, (int)strlen(json)); + g_sock_send(g_sock_out, "\n\n", 2); + } else { + fwrite(hdr, 1, (size_t)hl, stdout); + fwrite(json, 1, (size_t)strlen(json), stdout); + fwrite("\n\n", 1, 2, stdout); + fflush(stdout); + } +} + +/* Decode a single token id into its raw (unmapped) bytes. + * The vocab stores byte-level BPE pieces: each piece is UTF-8 of the + * GPT-2 byte_to_unicode-mapped codepoints. We reverse that mapping so the + * output is the original text bytes (correct for CJK / non-ASCII too). + * <0xXX> byte-fallback tokens emit the raw byte directly. */ +static void decode_id_to_bytes(int id, unsigned char *out, int *outn){ + *outn = 0; + if (!g_tok || id<0 || id>=g_tok_n) return; + const unsigned char *pc = (const unsigned char*)g_tok[id]; + /* byte-fallback token: <0xXX> -> raw byte */ + if (pc[0]=='<' && pc[1]=='0' && pc[2]=='x' && pc[5]=='>'){ + out[(*outn)++] = (unsigned char)(hexnib((char)pc[3])*16 + hexnib((char)pc[4])); + return; + } + int i = 0; + while (pc[i]){ + int cp, extra; + if (pc[i] < 0x80){ cp = pc[i]; extra = 0; } + else if ((pc[i] & 0xE0) == 0xC0){ cp = pc[i] & 0x1F; extra = 1; } + else if ((pc[i] & 0xF0) == 0xE0){ cp = pc[i] & 0x0F; extra = 2; } + else if ((pc[i] & 0xF8) == 0xF0){ cp = pc[i] & 0x07; extra = 3; } + else { i++; continue; } /* stray lead byte, skip */ + int ok = 1; + for (int e=0; e= 0) out[(*outn)++] = (unsigned char)g_unmap[cp]; /* reverse byte_to_unicode */ + else out[(*outn)++] = (unsigned char)cp; + if (*outn >= 255) break; + } +} + +/* Decode a range of token ids into a NUL-terminated text buffer (non-streaming). */ +static int decode_range(const int *arr, int from, int to, char *ob, int obsz){ + unsigned char sb[16]; int sbn = 0; int o = 0; + for (int i=from;i *bn) break; + fwrite(buf+j, 1, (size_t)need, stdout); + memmove(buf+j, buf+j+need, *bn-j-need); + *bn -= need; + } + } +} + +static void print_decoded(const int *arr, int from, int to){ + unsigned char buf[16]; int bn = 0; + for (int i=from;i 0){ + char esc[1024]; json_escape(chunk, cn, esc, sizeof esc); + char jb[1200]; + snprintf(jb, sizeof jb, + "{\"id\":\"%s\",\"object\":\"chat.completion.chunk\",\"created\":%ld,\"model\":\"%s\"," + "\"choices\":[{\"index\":0,\"delta\":{\"content\":\"%s\"},\"finish_reason\":null}]}", + g_oa_id, g_oa_created, g_model, esc); + sse_chunk(jb); + } + return; + } + /* default raw-text streaming */ + if (!g_tok){ printf("%d ", id); fflush(stdout); return; } + unsigned char tmp[256]; int tn = 0; + decode_id_to_bytes(id, tmp, &tn); + out_bytes(g_sbuf, &g_sbn, tmp, tn); + fflush(stdout); /* make streaming visible immediately even when piped */ +} +static void stream_flush(void){ if (g_sbn){ fwrite(g_sbuf, 1, (size_t)g_sbn, stdout); g_sbn = 0; } } + +/* Emit the final OpenAI Chat Completions response for a finished generation. + * Streaming: flushes any trailing partial UTF-8 as a last content chunk, then + * sends the termination chunk (finish_reason + usage + timings) and "data: [DONE]". + * Non-streaming: sends a single chat.completion JSON object. + * When g_sock_out >= 0 the bytes go to the live socket; otherwise to stdout. */ +static void emit_openai_result(const int *out, int np, int n_new, int stream){ + double total = now_s() - g_gen_t0; + if (g_ttft < 0) g_ttft = total; /* non-streaming: all tokens arrive at once */ + double gen_t = total - g_ttft; + double tps = (gen_t > 1e-6 && n_new > 1) ? n_new / gen_t : (total > 0 ? n_new / total : 0.0); + if (stream){ + if (g_sbn > 0){ + unsigned char chunk[16]; int cn = 0; + for (int k=0;k 0){ + char esc[256]; json_escape(chunk, cn, esc, sizeof esc); + char jb[400]; + snprintf(jb, sizeof jb, + "{\"id\":\"%s\",\"object\":\"chat.completion.chunk\",\"created\":%ld,\"model\":\"%s\"," + "\"choices\":[{\"index\":0,\"delta\":{\"content\":\"%s\"},\"finish_reason\":null}]}", + g_oa_id, g_oa_created, g_model, esc); + sse_chunk(jb); + } + } + char jb[700]; + snprintf(jb, sizeof jb, + "{\"id\":\"%s\",\"object\":\"chat.completion.chunk\",\"created\":%ld,\"model\":\"%s\"," + "\"choices\":[{\"index\":0,\"delta\":{},\"finish_reason\":\"stop\"}]," + "\"usage\":{\"prompt_tokens\":%d,\"completion_tokens\":%d,\"total_tokens\":%d}," + "\"timings\":{\"ttft_s\":%.3f,\"tokens_per_sec\":%.3f,\"total_s\":%.3f}}", + g_oa_id, g_oa_created, g_model, np, n_new, np+n_new, g_ttft, tps, total); + sse_chunk(jb); + char done[16]; int dl = snprintf(done, sizeof done, "data: [DONE]\n\n"); + if (g_sock_out >= 0 && g_sock_send) g_sock_send(g_sock_out, done, dl); + else { fwrite(done, 1, (size_t)dl, stdout); fflush(stdout); } + } else { + char text[1<<16]; decode_range(out, np, np+n_new, text, sizeof text); + char esc[1<<16]; json_escape((const unsigned char*)text, (int)strlen(text), esc, sizeof esc); + char buf[1<<20]; + int bl = snprintf(buf, sizeof buf, + "{\"id\":\"%s\",\"object\":\"chat.completion\",\"created\":%ld,\"model\":\"%s\"," + "\"choices\":[{\"index\":0,\"message\":{\"role\":\"assistant\",\"content\":\"%s\"},\"finish_reason\":\"stop\"}]," + "\"usage\":{\"prompt_tokens\":%d,\"completion_tokens\":%d,\"total_tokens\":%d}," + "\"timings\":{\"ttft_s\":%.3f,\"tokens_per_sec\":%.3f,\"total_s\":%.3f}}\n", + g_oa_id, g_oa_created, g_model, esc, np, n_new, np+n_new, g_ttft, tps, total); + if (g_sock_out >= 0 && g_sock_send) g_sock_send(g_sock_out, buf, bl); + else { fwrite(buf, 1, (size_t)bl, stdout); fflush(stdout); } + } +} + +/* ---------- config ---------- */ +typedef struct { + int hidden, n_layers, n_active; + int q_heads, kv_heads, head_dim; /* k/v head dim == attention head dim */ + int q_head_dim; /* q per-head total = head_dim*2 when attn_output_gate */ + int k_head_dim, v_head_dim, o_in; /* o_in = q_heads*head_dim (o_proj input) */ + int rope_dim, rotary_dim; /* rotary_dim = actual rotated dims (head_dim*partial_rotary_factor) */ + int n_experts, topk, inter, shared_inter, vocab; + int n_group, topk_group; + float theta, eps, partial_rotary_factor; + int norm_topk, has_qk_norm, has_bias, attn_output_gate; + uint8_t *is_attn; /* [n_layers] 1 if Gated Attention layer, 0 if DeltaNet */ + /* Gated DeltaNet (linear_attention) dims, read from qwen36_meta.json. */ + int dn_vheads, dn_kheads, dn_kdim, dn_vdim, dn_convk, dn_conv_dim; +} Cfg; + +/* ---------- per-layer dense weights ---------- */ +typedef struct { + float *in_ln, *post_ln, *q, *k, *v, *o, *qn, *kn, *gate, *gate_bias; + float *sh_g, *sh_u, *sh_d, *sh_gate; /* shared expert (dense f32) + shared_expert_gate */ + /* Gated DeltaNet (linear_attention) dense weights (f16->f32 via st_read_f32). */ + float *dn_qkv, *dn_z, *dn_b, *dn_a; /* in_proj_qkv/z/b/a */ + float *dn_conv; /* conv1d.weight [conv_dim, convk] (groups=conv_dim) */ + float *dn_dtbias, *dn_alog; /* dt_bias[vh], A_log[vh] */ + float *dn_norm; /* RMSNormGated weight [vdim] */ + float *dn_out; /* out_proj [hidden, value_dim] */ +} Layer; + +/* ---------- LRU expert cache (int8 weights + per-row float scales) ---------- */ +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 { + Cfg c; + shards S; + int quant_bits; + float *embed, *lm_head, *final_norm; + Layer *L; + LCache *cache; /* [n_layers] */ + int *active_of; /* [n_layers] original->active idx (Phase 2: identity for all layers) */ + float **DN_rec; /* [n_layers] recurrent state S[h]=[kdim,vdim] for DeltaNet layers (NULL for attn) */ + float **DN_conv; /* [n_layers] conv ring [conv_dim, convk-1] for DeltaNet layers (NULL for attn) */ + uint64_t clock, hits, miss; + float **K, **V; int kv_len, max_t, kv_cap; + double dense_load_s; + uint32_t *freq; + int freq_token_count, hot_pinned, hot_n, warmup_tokens, token_count; + float *momentum_logits; + float pilot_smooth, pilot_conf_limit; + uint8_t *is_pinned; + uint8_t *is_queued; + uint8_t *seen; /* prefill-collected experts (COLIBRI_RESIDENT) */ + int resident_mode; /* 0 off; 1 pin this-prompt experts (CPU no-evict -> GPU resident) */ + int resident_collecting; /* prefill in progress, collecting routed experts */ + int first_step; /* the first step() call is the prefill */ +} Model; + +static pthread_mutex_t g_pilot_mx = PTHREAD_MUTEX_INITIALIZER; +static struct { int l, e; } pilot_q[4096]; +static volatile unsigned pilot_r = 0, pilot_w = 0; +static Model *pilot_m = NULL; +static int g_pilot = 0; +static int g_wide = 1; + +static void pilot_prefetch(Model *m, int lnext, const float *x, int S); +static void *pilot_worker(void *arg); +static void ensure_pilot_worker_started(Model *m); +static void slot_ensure_allocated(Model *m, Slot *s); + +static void ensure_pilot_worker_started(Model *m) { + if (!pilot_m) { + pilot_m = m; + pthread_t t; + if (pthread_create(&t, NULL, pilot_worker, NULL) != 0) { + fprintf(stderr, "Error: Failed to create pilot prefetch worker thread\n"); + exit(1); + } + pthread_detach(t); + } +} + +/* ---------- utility ---------- */ +static double now_s(void) { struct timespec t; clock_gettime(CLOCK_MONOTONIC, &t); return t.tv_sec + t.tv_nsec*1e-9; } +#if defined(__APPLE__) +static double rss_gb(void) { struct rusage r; getrusage(RUSAGE_SELF, &r); return r.ru_maxrss / (1024.0*1024.0*1024.0); } +#else +static double rss_gb(void) { struct rusage r; getrusage(RUSAGE_SELF, &r); return r.ru_maxrss / (1024.0*1024.0); } +#endif + +/* ---- M-PROF (R2): per-phase wall-clock accumulators, COLI_TIMERS=1 ---- */ +static int g_timers = -1; +static double g_tm_dec[6], g_tm_pre[6]; /* 0=deltanet 1=attention 2=moe_total 3=shared 4=router 5=lm_head */ +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_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; +static void tm_add(int S, int idx, double ms){ + if(S==1){ + g_tm_dec[idx]+=ms; + if(idx==2) g_tm_win_moe+=ms; + if(idx==5 && ++g_tm_win_n==32){ + fprintf(stderr,"[timers] window: moe %.0f ms/token (last 32)\n", g_tm_win_moe/32.0); + g_tm_win_moe=0; g_tm_win_n=0; + } + } else g_tm_pre[idx]+=ms; +} +static void tm_report(void){ + if(!tm_on()) return; + static const char *nm[6]={"deltanet","attention","moe_total","(shared)","(router)","lm_head"}; + fprintf(stderr,"[timers] decode: %ld tokens (shared/router are subsets of moe_total)\n", g_tm_dec_tokens); + double sum=0; + for(int i=0;i<6;i++){ + fprintf(stderr,"[timers] %-10s %9.1f ms %8.2f ms/token\n", + nm[i], g_tm_dec[i], g_tm_dec_tokens? g_tm_dec[i]/g_tm_dec_tokens:0.0); + if(i!=3&&i!=4) sum+=g_tm_dec[i]; + } + fprintf(stderr,"[timers] %-10s %9.1f ms %8.2f ms/token\n","TOTAL",sum,g_tm_dec_tokens?sum/g_tm_dec_tokens:0.0); + if(g_tm_step>0) + fprintf(stderr,"[timers] step() total: %.1f ms/token (outside the phases: %.1f)\n", + g_tm_step/g_tm_dec_tokens, + (g_tm_step-(g_tm_dec[0]+g_tm_dec[1]+g_tm_dec[2]+g_tm_dec[5]))/g_tm_dec_tokens); + 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); + 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]); +} +static float *falloc(int64_t n) { float *p = malloc(n*sizeof(float)); if(!p){fprintf(stderr,"OOM %ld\n",(long)n);exit(1);} return p; } + +/* y[S,O] = x[S,I] @ W^T, W is [O,I] row-major */ +static void matmul(float *y, const float *x, const float *W, int S, int I, int O) { + #pragma omp parallel for schedule(static) + for (int o = 0; o < O; o++) { + const float *w = W + (int64_t)o * I; + for (int s = 0; s < S; s++) { + const float *xs = x + (int64_t)s * I; + float acc = 0.f; + for (int i = 0; i < I; i++) acc += xs[i] * w[i]; + y[(int64_t)s * O + o] = acc; + } + } +} + +/* y[1,O] = x[1,I] @ W^T with W quantized: q[O,I] int8 + scale per row. */ +#if defined(__ARM_NEON) +#include +static inline int32_t dot_i8_16(const int8_t *a, const int8_t *b) { + int32x4_t acc = vdupq_n_s32(0); + int8x16_t va = vld1q_s8(a), vb = vld1q_s8(b); +#if defined(__ARM_FEATURE_DOTPROD) + acc = vdotq_s32(acc, va, vb); +#else + acc = vpadalq_s16(acc, vmull_s8(vget_low_s8(va), vget_low_s8(vb))); + acc = vpadalq_s16(acc, vmull_s8(vget_high_s8(va), vget_high_s8(vb))); +#endif + return vaddvq_s32(acc); +} +#endif +static void matmul_q(float *y, const float *x, const int8_t *q, const float *scale, int I, int O) { +#if defined(__ARM_NEON) + static int idot = -1; + if (idot < 0) { const char *e = getenv("IDOT"); idot = !(e && *e == '0'); } + if (idot && I % 16 == 0 && I <= 4096) { + int nb = I / 16; int8_t xi[4096]; float xs[256]; + for (int b = 0; b < nb; b++) { + const float *xb = x + b*16; + float am = 0.f; for (int i = 0; i < 16; i++) { float a = fabsf(xb[i]); if (a > am) am = a; } + float s = am/127.f; if (s < 1e-12f) s = 1e-12f; + xs[b] = s; float inv = 1.f/s; + for (int i = 0; i < 16; i++) xi[b*16+i] = (int8_t)lrintf(xb[i]*inv); + } + #pragma omp parallel for schedule(static) + for (int o = 0; o < O; o++) { + const int8_t *w = q + (int64_t)o * I; + float acc = 0.f; + for (int b = 0; b < nb; b++) acc += xs[b]*(float)dot_i8_16(xi+b*16, w+b*16); + y[o] = acc * scale[o]; + } + return; + } +#endif +#if defined(__AVX2__) && defined(__FMA__) + /* Hand-vectorized int8->f32 GEMV (gcc does not auto-vectorize the + * convert+accumulate chain). 32 weights per iteration, FMA accumulate. */ + #pragma omp parallel for schedule(static) if(O >= 256) + for (int o = 0; o < O; o++) { + const int8_t *w = q + (int64_t)o * I; + __m256 a0 = _mm256_setzero_ps(), a1 = _mm256_setzero_ps(); + __m256 a2 = _mm256_setzero_ps(), a3 = _mm256_setzero_ps(); + int i = 0; + for (; i + 32 <= I; i += 32) { + __m128i b0 = _mm_loadu_si128((const __m128i*)(w + i)); + __m128i b1 = _mm_loadu_si128((const __m128i*)(w + i + 16)); + a0 = _mm256_fmadd_ps(_mm256_loadu_ps(x+i), _mm256_cvtepi32_ps(_mm256_cvtepi8_epi32(b0)), a0); + a1 = _mm256_fmadd_ps(_mm256_loadu_ps(x+i+8), _mm256_cvtepi32_ps(_mm256_cvtepi8_epi32(_mm_srli_si128(b0,8))), a1); + a2 = _mm256_fmadd_ps(_mm256_loadu_ps(x+i+16), _mm256_cvtepi32_ps(_mm256_cvtepi8_epi32(b1)), a2); + a3 = _mm256_fmadd_ps(_mm256_loadu_ps(x+i+24), _mm256_cvtepi32_ps(_mm256_cvtepi8_epi32(_mm_srli_si128(b1,8))), a3); + } + a0 = _mm256_add_ps(_mm256_add_ps(a0,a1), _mm256_add_ps(a2,a3)); + __m128 s = _mm_add_ps(_mm256_castps256_ps128(a0), _mm256_extractf128_ps(a0,1)); + s = _mm_add_ps(s, _mm_movehl_ps(s,s)); + s = _mm_add_ss(s, _mm_shuffle_ps(s,s,1)); + float acc = _mm_cvtss_f32(s); + for (; i < I; i++) acc += x[i] * (float)w[i]; + y[o] = acc * scale[o]; + } +#else + #pragma omp parallel for schedule(static) + for (int o = 0; o < O; o++) { + const int8_t *w = q + (int64_t)o * I; + float acc = 0.f; + for (int i = 0; i < I; i++) acc += x[i] * (float)w[i]; + y[o] = acc * scale[o]; + } +#endif +} + +/* ---- Dense int8: per-row quantized copies of the large f32 matrices. + * matmul_d dispatches via pointer lookup to matmul_q; COLI_DENSE_I8=0 falls + * back to f32 (reference path for parity tests). ~4x less memory traffic. */ +#define QDW_MAX 1024 +static struct { const float *w; int8_t *q; float *sc; int I, O; } g_qdw[QDW_MAX]; +static int g_qdw_n = 0; +static int dense_i8_on(void){ static int v=-1; if(v<0){ const char *e=getenv("COLI_DENSE_I8"); v=!(e&&*e=='0'); } return v; } +static void qdw_register(const float *W, int I, int O){ + if (!W || !dense_i8_on() || g_qdw_n >= QDW_MAX) return; + int8_t *q = malloc((size_t)O*I); float *sc = malloc((size_t)O*sizeof(float)); + if (!q || !sc) { free(q); free(sc); return; } + #pragma omp parallel for schedule(static) + for (int o = 0; o < O; o++) { + const float *r = W + (int64_t)o*I; float am = 0.f; + for (int i = 0; i < I; i++) { float a = fabsf(r[i]); if (a > am) am = a; } + float s = am > 1e-12f ? am/127.f : 1.f; sc[o] = s; float inv = 1.f/s; + int8_t *d = q + (int64_t)o*I; + for (int i = 0; i < I; i++) { int v = (int)lrintf(r[i]*inv); if (v>127) v=127; if (v<-127) v=-127; d[i] = (int8_t)v; } + } + g_qdw[g_qdw_n].w=W; g_qdw[g_qdw_n].q=q; g_qdw[g_qdw_n].sc=sc; g_qdw[g_qdw_n].I=I; g_qdw[g_qdw_n].O=O; g_qdw_n++; +} +static void matmul_d(float *y, const float *x, const float *W, int S, int I, int O){ + for (int i = 0; i < g_qdw_n; i++) if (g_qdw[i].w == W && g_qdw[i].I == I) { + for (int s = 0; s < S; s++) matmul_q(y+(int64_t)s*O, x+(int64_t)s*I, g_qdw[i].q, g_qdw[i].sc, I, O); + return; + } + matmul(y, x, W, S, I, O); +} + +/* rmsnorm over a row of length D (in-place capable: out may == x). + * Qwen3_5MoeRMSNorm: out = (x * rsqrt(mean(x^2)+eps)) * (1.0 + weight). */ +static void rmsnorm_row(float *out, const float *x, const float *w, int D, float eps) { + double ms = 0; for (int i = 0; i < D; i++) ms += (double)x[i]*x[i]; + float r = 1.f / sqrtf((float)(ms / D) + eps); + for (int i = 0; i < D; i++) out[i] = x[i] * r * (1.0f + w[i]); +} + +static void softmax_row(float *x, int n) { + float m = -1e30f; for (int i = 0; i < n; i++) if (x[i] > m) m = x[i]; + float s = 0; for (int i = 0; i < n; i++) { x[i] = expf(x[i]-m); s += x[i]; } + for (int i = 0; i < n; i++) x[i] /= s; +} + +/* softplus(z) = log(1+exp(z)), stable for large z (HF GatedDeltaNet g_rule). */ +static float softplus_f(float z) { return z > 20.f ? z : log1pf(expf(z)); } + +/* ---------- loading ---------- */ +static double req_num(jval *r, const char *k){ + jval *v=json_get(r,k); + if(!v||v->t!=J_NUM){ fprintf(stderr,"config.json: missing or non-numeric \"%s\"\n",k); exit(1); } + return v->num; +} +static void load_cfg(Cfg *c, const char *snap) { + char path[2048]; snprintf(path, sizeof(path), "%s/config.json", snap); + FILE *f = fopen(path, "rb"); if(!f){perror(path);exit(1);} + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + if(n<0 || n>(256L<<20)){ fprintf(stderr,"%s: config.json missing or larger than 256 MB\n",path); exit(1); } + char *buf = malloc((size_t)n+1); if(!buf){ fprintf(stderr,"OOM reading %s\n",path); exit(1); } + if(fread(buf,1,(size_t)n,f)!=(size_t)n){ fprintf(stderr,"%s: short read\n",path); exit(1); } buf[n]=0; fclose(f); + char *arena=NULL; jval *r = json_parse(buf, &arena); + c->hidden = (int)req_num(r,"hidden_size"); + c->n_layers = (int)req_num(r,"num_hidden_layers"); + c->vocab = (int)req_num(r,"vocab_size"); + c->eps = (float)req_num(r,"rms_norm_eps"); + jval *th = json_get(r,"rope_theta"); c->theta = th ? (float)th->num : 10000.f; + free(buf); free(arena); + /* Phase-1 defaults; overridden by qwen36_meta.json in load_meta. + * Clamped so a missing meta file can never produce a divide-by-zero. */ + c->q_heads = (c->hidden >= 16) ? (c->hidden / 16) : 1; + if (c->q_heads < 1) c->q_heads = 1; + c->kv_heads = c->q_heads / 8; if (c->kv_heads < 1) c->kv_heads = 1; + c->head_dim = c->hidden / c->q_heads; if (c->head_dim < 1) c->head_dim = 1; + c->k_head_dim = c->head_dim; c->v_head_dim = c->head_dim; + c->q_head_dim = c->head_dim * 2; /* includes attn_output_gate */ + c->o_in = c->q_heads * c->head_dim; + c->rotary_dim = (c->head_dim >= 4) ? (c->head_dim / 4) : 2; + if (c->rotary_dim % 2 != 0) c->rotary_dim += 1; + c->rope_dim = c->head_dim; + c->partial_rotary_factor = 0.25f; + c->n_experts = 256; c->topk = 8; c->inter = 512; c->shared_inter = 512; + c->n_group = 1; c->topk_group = 1; c->norm_topk = 1; c->has_qk_norm = 1; c->has_bias = 0; + c->attn_output_gate = 1; c->n_active = 0; + c->is_attn = calloc(c->n_layers, sizeof(uint8_t)); + for (int i = 0; i < c->n_layers; i++) c->is_attn[i] = (i % 4 == 3) ? 1 : 0; +} + +/* Read qwen36_meta.json (emitted FLAT by convert_qwen36.py) to override the + * Phase-1 defaults with the real model dimensions. The converter derives the + * head dims from the actual weight shapes, so these are authoritative. Falls + * back silently to the i%4==3 pattern and defaults if the file is absent. */ +static void load_meta(Cfg *c, const char *snap) { + char path[2048]; snprintf(path, sizeof(path), "%s/qwen36_meta.json", snap); + FILE *f = fopen(path, "rb"); if (!f) { printf("[meta] %s not found; using i%%4==3 + defaults\n", path); return; } + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + char *buf = malloc((size_t)n+1); if(!buf){fclose(f);return;} + if(fread(buf,1,(size_t)n,f)!=(size_t)n){ free(buf); fclose(f); return; } buf[n]=0; fclose(f); + char *arena=NULL; jval *r = json_parse(buf, &arena); + if (r && r->t == J_OBJ) { + jval *v; + #define G(name,field) if((v=json_get(r,name))&&v->t==J_NUM) c->field=(int)v->num + G("hidden", hidden); G("n_layers", n_layers); G("n_active", n_active); + G("q_heads", q_heads); G("kv_heads", kv_heads); G("head_dim", head_dim); + G("q_head_dim", q_head_dim); G("k_head_dim", k_head_dim); G("v_head_dim", v_head_dim); + G("o_in", o_in); G("rope_dim", rope_dim); G("qk_rope_head_dim", rope_dim); + G("num_experts", n_experts); G("topk", topk); + G("moe_inter", inter); G("shared_inter", shared_inter); + G("n_group", n_group); G("topk_group", topk_group); + G("dn_vheads", dn_vheads); G("dn_kheads", dn_kheads); G("dn_kdim", dn_kdim); + G("dn_vdim", dn_vdim); G("dn_convk", dn_convk); G("dn_conv_dim", dn_conv_dim); + #undef G + if((v=json_get(r,"partial_rotary_factor"))&&v->t==J_NUM) c->partial_rotary_factor=(float)v->num; + if((v=json_get(r,"rope_theta"))&&v->t==J_NUM) c->theta=(float)v->num; + if((v=json_get(r,"rms_eps"))&&v->t==J_NUM) c->eps=(float)v->num; + if((v=json_get(r,"attn_output_gate"))&&v->t==J_BOOL) c->attn_output_gate=v->boolean; + if((v=json_get(r,"norm_topk_prob"))&&v->t==J_BOOL) c->norm_topk=v->boolean; + if((v=json_get(r,"has_bias"))&&v->t==J_BOOL) c->has_bias=v->boolean; + if((v=json_get(r,"has_qk_norm"))&&v->t==J_BOOL) c->has_qk_norm=v->boolean; + /* derive rotary_dim from head_dim * partial_rotary_factor (HF formula) */ + if (c->partial_rotary_factor > 0.f) + c->rotary_dim = (int)(c->head_dim * c->partial_rotary_factor + 0.5f); + else + c->rotary_dim = c->head_dim; + if (c->rotary_dim < 2) c->rotary_dim = 2; + if (c->rotary_dim % 2 != 0) c->rotary_dim += 1; + if (c->rotary_dim > c->head_dim) c->rotary_dim = c->head_dim; + /* rebuild is_attn from explicit layer_types if present */ + jval *lt = json_get(r,"layer_types"); + if (lt && lt->t==J_ARR) { + for (int i=0;in_layers;i++) c->is_attn[i]=0; + c->n_active=0; + for (int i=0;ilen && in_layers;i++){ + const char *s = (lt->kids[i]->t==J_STR)? lt->kids[i]->str : ""; + if (s && strcmp(s,"full_attention")==0) { c->is_attn[i]=1; c->n_active++; } + } + } + } + free(buf); free(arena); + fprintf(stderr, "[meta] loaded: q_heads=%d kv_heads=%d head_dim=%d q_head_dim=%d k_head_dim=%d v_head_dim=%d " + "o_in=%d rotary_dim=%d n_experts=%d topk=%d inter=%d shared_inter=%d n_group=%d topk_group=%d " + "attn_output_gate=%d n_active=%d\n", + c->q_heads, c->kv_heads, c->head_dim, c->q_head_dim, c->k_head_dim, c->v_head_dim, + c->o_in, c->rotary_dim, c->n_experts, c->topk, c->inter, c->shared_inter, + c->n_group, c->topk_group, c->attn_output_gate, c->n_active); + if (c->dn_vheads > 0) + fprintf(stderr, "[meta] DeltaNet: vheads=%d kheads=%d kdim=%d vdim=%d convk=%d conv_dim=%d\n", + c->dn_vheads, c->dn_kheads, c->dn_kdim, c->dn_vdim, c->dn_convk, c->dn_conv_dim); +} + +static float *load_t(Model *m, const char *name) { + int64_t n = st_numel(&m->S, name); + if (n < 0) { fprintf(stderr, "missing %s\n", name); exit(1); } + float *p = falloc(n); + st_read_f32(&m->S, name, p, 0); + return p; +} + +static void model_init(Model *m, const char *snap, int cap, int bits) { + memset(m, 0, sizeof(*m)); + m->quant_bits = bits; + load_cfg(&m->c, snap); + load_meta(&m->c, snap); + if (m->c.rotary_dim > m->c.head_dim || m->c.rotary_dim % 2 != 0) { + fprintf(stderr, "rotary_dim %d invalid for head_dim %d\n", m->c.rotary_dim, m->c.head_dim); exit(1); + } + st_init(&m->S, snap); + Cfg *c = &m->c; + double t0 = now_s(); + m->embed = load_t(m, "model.embed_tokens.weight"); + m->lm_head = load_t(m, "lm_head.weight"); + m->final_norm = load_t(m, "model.norm.weight"); + m->L = calloc(c->n_layers, sizeof(Layer)); + /* Phase 2: the converter stores EVERY layer (Gated-Attention + Gated DeltaNet) + * under its OWN original index model.layers.{i}. So active_of is the identity + * map; experts and dense weights are read from model.layers.{i} for all i. */ + m->active_of = malloc((size_t)c->n_layers * sizeof(int)); + for (int i = 0; i < c->n_layers; i++) m->active_of[i] = i; + char nm[256]; + for (int i = 0; i < c->n_layers; i++) { + int ai = m->active_of[i]; /* == i for Phase 2 */ + Layer *l = &m->L[i]; + /* input/post layernorms + MoE exist for every layer */ + #define LD(field, suffix) snprintf(nm,sizeof(nm),"model.layers.%d." suffix,ai); l->field = load_t(m,nm) + LD(in_ln, "input_layernorm.weight"); + LD(post_ln,"post_attention_layernorm.weight"); + LD(gate, "mlp.gate.weight"); + #undef LD + /* q/k norms are per-head [head_dim]; only on attention layers, load if present */ + if (c->has_qk_norm) { + snprintf(nm,sizeof(nm),"model.layers.%d.self_attn.q_norm.weight", ai); + l->qn = st_has(&m->S, nm) ? load_t(m, nm) : NULL; + snprintf(nm,sizeof(nm),"model.layers.%d.self_attn.k_norm.weight", ai); + l->kn = st_has(&m->S, nm) ? load_t(m, nm) : NULL; + } else { l->qn = NULL; l->kn = NULL; } + /* router correction bias (optional) */ + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.gate.e_score_correction_bias", ai); + if (st_has(&m->S, nm)) { l->gate_bias = falloc(c->n_experts); st_read_f32(&m->S, nm, l->gate_bias, 0); } + else l->gate_bias = NULL; + /* shared expert (dense f32) */ + #define LD2(field, suffix) snprintf(nm,sizeof(nm),"model.layers.%d.mlp.shared_expert." suffix,ai); l->field = load_t(m,nm) + LD2(sh_g, "gate_proj.weight"); LD2(sh_u, "up_proj.weight"); LD2(sh_d, "down_proj.weight"); + #undef LD2 + /* shared_expert_gate: Linear(hidden -> 1), sigmoid-gated shared expert */ + snprintf(nm,sizeof(nm),"model.layers.%d.mlp.shared_expert_gate.weight", ai); + l->sh_gate = st_has(&m->S, nm) ? load_t(m, nm) : NULL; + if (c->is_attn[i]) { + /* Gated Attention (full_attention) layer */ + #define LD3(field, suffix) snprintf(nm,sizeof(nm),"model.layers.%d.self_attn." suffix,ai); l->field = load_t(m,nm) + LD3(q, "q_proj.weight"); LD3(k, "k_proj.weight"); + LD3(v, "v_proj.weight"); LD3(o, "o_proj.weight"); + #undef LD3 + l->dn_qkv=l->dn_z=l->dn_b=l->dn_a=l->dn_conv=NULL; + l->dn_dtbias=l->dn_alog=l->dn_norm=l->dn_out=NULL; + } else { + /* Gated DeltaNet (linear_attention) layer */ + l->q=l->k=l->v=l->o=NULL; + #define LD4(field, suffix) snprintf(nm,sizeof(nm),"model.layers.%d.linear_attn." suffix,ai); l->field = load_t(m,nm) + LD4(dn_qkv, "in_proj_qkv.weight"); LD4(dn_z, "in_proj_z.weight"); + LD4(dn_b, "in_proj_b.weight"); LD4(dn_a, "in_proj_a.weight"); + LD4(dn_conv,"conv1d.weight"); LD4(dn_dtbias, "dt_bias"); + LD4(dn_alog,"A_log"); LD4(dn_norm, "norm.weight"); + LD4(dn_out, "out_proj.weight"); + #undef LD4 + } + } + m->cache = calloc(c->n_layers, sizeof(LCache)); + for (int i = 0; i < c->n_layers; i++) { + m->cache[i].cap = cap; + m->cache[i].slots = calloc(cap, sizeof(Slot)); + } + /* per-layer DeltaNet recurrent + conv state (only for linear_attention layers) */ + m->DN_rec = calloc(c->n_layers, sizeof(float*)); + m->DN_conv = calloc(c->n_layers, sizeof(float*)); + for (int i = 0; i < c->n_layers; i++) { + if (c->is_attn[i]) { m->DN_rec[i] = NULL; m->DN_conv[i] = NULL; continue; } + if (c->dn_vheads <= 0) { fprintf(stderr, "layer %d is DeltaNet but dn dims missing from meta\n", i); exit(1); } + m->DN_rec[i] = calloc((size_t)c->dn_vheads * c->dn_kdim * c->dn_vdim, sizeof(float)); + m->DN_conv[i] = calloc((size_t)c->dn_conv_dim * (c->dn_convk - 1), sizeof(float)); + } + m->freq = calloc((size_t)c->n_layers * c->n_experts, sizeof(uint32_t)); + m->hot_pinned = 0; m->freq_token_count = 0; + m->hot_n = getenv("HOT") ? atoi(getenv("HOT")) : 0; + m->warmup_tokens = getenv("WARMUP") ? atoi(getenv("WARMUP")) : 5; + m->token_count = 0; + m->momentum_logits = calloc((size_t)c->n_layers * c->n_experts, sizeof(float)); + float sv = getenv("SMOOTH") ? (float)atof(getenv("SMOOTH")) : 0.3f; + if (sv < 0.f) sv = 0.f; if (sv > 0.95f) sv = 0.95f; + m->pilot_smooth = sv; + m->is_pinned = calloc((size_t)c->n_layers * c->n_experts, sizeof(uint8_t)); + m->seen = calloc((size_t)c->n_layers * c->n_experts, 1); + m->resident_mode = getenv("COLIBRI_RESIDENT") ? atoi(getenv("COLIBRI_RESIDENT")) : 0; + m->resident_collecting = 0; + m->first_step = 1; + m->is_queued = calloc((size_t)c->n_layers * c->n_experts, sizeof(uint8_t)); + float cl = getenv("CONF_LIMIT") ? (float)atof(getenv("CONF_LIMIT")) : 0.92f; + if (cl < 0.1f) cl = 0.1f; if (cl > 1.0f) cl = 1.0f; + m->pilot_conf_limit = cl; + m->dense_load_s = now_s() - t0; +} + +static void slot_ensure_allocated(Model *m, Slot *s) { + if (s->g) return; + Cfg *c = &m->c; + int64_t ng = (int64_t)c->inter * c->hidden; + int64_t nd = (int64_t)c->hidden * c->inter; + int8_t *w_block = malloc(ng + ng + nd); + if (!w_block) { fprintf(stderr, "Error: OOM allocating slot weights\n"); exit(1); } + s->g = w_block; + s->u = w_block + ng; + s->d = w_block + ng + ng; + float *s_block = falloc(c->inter + c->inter + c->hidden); + s->gs = s_block; + s->us = s_block + c->inter; + s->ds = s_block + c->inter + c->inter; + 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) { + char nm[256], qsnm[256]; + int la = m->active_of[layer]; /* container stores experts under active index */ + snprintf(nm, sizeof(nm), "model.layers.%d.mlp.experts.%d.merged_weight", la, eid); + snprintf(qsnm, sizeof(qsnm), "model.layers.%d.mlp.experts.%d.qs", la, eid); + Cfg *cc = &m->c; + int64_t ng = (int64_t)cc->inter * cc->hidden, nd = (int64_t)cc->hidden * cc->inter; + int64_t want_w = ng + ng + nd; + int64_t want_s = (int64_t)cc->inter + cc->inter + cc->hidden; + st_tensor *tw = st_find(&m->S, nm), *ts = st_find(&m->S, qsnm); + if (!tw || (tw->nbytes != want_w && tw->nbytes != want_w / 2)) { + fprintf(stderr, "%s: expert weight is %lld bytes — expected %lld (int8) or %lld (int4)\n", + nm, (long long)(tw ? tw->nbytes : -1), (long long)want_w, (long long)(want_w / 2)); exit(1); } + if (!ts || ts->numel != want_s) { + fprintf(stderr, "%s: scale array is %lld elems — expected %lld (refusing)\n", + qsnm, (long long)(ts ? ts->numel : -1), (long long)want_s); exit(1); } + /* int4 detection by ON-DISK SIZE (robust against a mislabeled meta.ebits, e.g. the + i8 container whose meta says ebits=4 but stores int8). True int4 packed uint8 is + exactly N/2 bytes (N = 3*inter*hidden, always even). Unpack in-place to int8 so the + rest of the MoE path (matmul_q) is unchanged. Nibble convention (must match + c/tools/convert_qwen36.py pack_int4): LOW nibble = element 2k, HIGH nibble = 2k+1; + each nibble is signed 4-bit (sign-extend if bit3 set). */ + if (tw->nbytes == want_w / 2) { + static int noted = 0; + if (!noted) { fprintf(stderr, "[qwen36] int4 packed weights detected — unpacking to int8 in slot\n"); noted = 1; } + uint8_t *raw = (uint8_t *)malloc((size_t)(want_w / 2)); + if (!raw) { fprintf(stderr, "OOM reading int4 expert %s\n", nm); exit(1); } + st_read_raw(&m->S, nm, raw, 1); + for (int64_t i = 0; i < want_w; i++) { + uint8_t byte = raw[i >> 1]; + int8_t v = (int8_t)((i & 1) ? ((byte >> 4) & 0xF) : (byte & 0xF)); + if (v & 8) v -= 16; /* sign-extend signed 4-bit */ + s->g[i] = v; + } + s->is_int4 = 1; + /* Keep the packed bytes for the int4 GPU shader. The unpacked int8 copy + * above is retained for the CPU fallback / int8-GPU path. Only allocate + * when a GPU int4 backend is actually active, to avoid doubling expert + * memory on CPU-only or int8-GPU runs. 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: they are + * the upload source for the (optional) CUDA expert tier and allow int8 + * rematerialisation without touching the container again. */ + { + 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); +} + +/* Robust int4 detection by on-disk size of one expert tensor (ignores a possibly + * mislabeled meta.ebits — cf. load_expert_merged). Returns 1 if the container + * stores true int4 packed weights, 0 otherwise. Used to pick the Vulkan + * pipeline at init time. */ +static int container_is_int4(Model *m) { + Cfg *cc = &m->c; + int64_t ng = (int64_t)cc->inter * cc->hidden, nd = (int64_t)cc->hidden * cc->inter; + int64_t want_w = ng + ng + nd; + char nm[256]; + snprintf(nm, sizeof(nm), "model.layers.0.mlp.experts.0.merged_weight"); + st_tensor *tw = st_find(&m->S, nm); + if (!tw) return 0; + return (tw->nbytes == want_w / 2) ? 1 : 0; +} + +static void expert_get(Model *m, int layer, int eid, Slot **out) { + LCache *lc = &m->cache[layer]; + pthread_mutex_lock(&g_pilot_mx); + for (int i = 0; i < lc->n; i++) if (lc->slots[i].eid == eid) { + m->hits++; lc->slots[i].used = ++m->clock; *out = &lc->slots[i]; + pthread_mutex_unlock(&g_pilot_mx); return; + } + m->miss++; + Cfg *c = &m->c; Slot *s; + if (lc->n < lc->cap) { s = &lc->slots[lc->n++]; slot_ensure_allocated(m, s); } + else { + int lru = -1; + for (int i = 0; i < lc->n; i++) { + if (lc->slots[i].pinned || lc->slots[i].eid < 0) continue; + if (lru < 0 || lc->slots[i].used < lc->slots[lru].used) lru = i; + } + if (lru < 0) { + for (int i = 0; i < lc->n; i++) { if (lc->slots[i].eid < 0) continue; if (lru < 0 || lc->slots[i].used < lc->slots[lru].used) lru = i; } + } + if (lru < 0) lru = 0; + s = &lc->slots[lru]; s->pinned = 0; + } + s->eid = -1; s->used = ++m->clock; + pthread_mutex_unlock(&g_pilot_mx); + load_expert_merged(m, layer, eid, s); + pthread_mutex_lock(&g_pilot_mx); + s->eid = eid; s->pinned = m->is_pinned[layer * c->n_experts + eid]; s->used = ++m->clock; + *out = s; pthread_mutex_unlock(&g_pilot_mx); +} + +static void pin_hot_experts(Model *m) { + Cfg *c = &m->c; + if (m->hot_n <= 0 || m->hot_pinned) return; + m->hot_pinned = 1; + int is_dynamic = (m->hot_n >= 100); + double thresh = is_dynamic ? (double)m->hot_n / 1000.0 : 0.0; + int pinned_total = 0; + for (int l = 0; l < c->n_layers; l++) { + uint32_t *freq_l = m->freq + (int64_t)l * c->n_experts; + uint64_t layer_total = 0; + for (int e = 0; e < c->n_experts; e++) layer_total += freq_l[e]; + if (layer_total == 0) continue; + int max_pin = m->cache[l].cap - 8; if (max_pin < 4) max_pin = 4; + int hn = is_dynamic ? max_pin : (m->hot_n < c->n_experts ? m->hot_n : c->n_experts); + if (hn > 256) hn = 256; + int hot_eids[256], actual_hn = 0; + for (int k = 0; k < hn; k++) { + int best = -1; uint32_t bv = 0; + for (int e = 0; e < c->n_experts; e++) { + int already = 0; + for (int j = 0; j < k; j++) if (hot_eids[j] == e) { already = 1; break; } + if (!already && freq_l[e] > bv) { bv = freq_l[e]; best = e; } + } + if (best < 0 || bv == 0) break; + if (is_dynamic && bv < thresh * layer_total) break; + hot_eids[k] = best; actual_hn++; + } + for (int k = 0; k < actual_hn; k++) { + int eid = hot_eids[k]; + m->is_pinned[l * c->n_experts + eid] = 1; + LCache *lc = &m->cache[l]; + int found = 0; + pthread_mutex_lock(&g_pilot_mx); + for (int i = 0; i < lc->n; i++) if (lc->slots[i].eid == eid) { lc->slots[i].pinned = 1; found = 1; break; } + pthread_mutex_unlock(&g_pilot_mx); + if (!found && g_pilot > 0) { + ensure_pilot_worker_started(m); + unsigned w = __atomic_load_n(&pilot_w, __ATOMIC_RELAXED); + unsigned r = __atomic_load_n(&pilot_r, __ATOMIC_ACQUIRE); + int gidx = l * c->n_experts + eid; + pthread_mutex_lock(&g_pilot_mx); + int already = m->is_queued[gidx]; + if (!already && w - r < 4096) { + pilot_q[w & 4095].l = l; pilot_q[w & 4095].e = eid; m->is_queued[gidx] = 1; + __atomic_store_n(&pilot_w, w + 1, __ATOMIC_RELEASE); + } + pthread_mutex_unlock(&g_pilot_mx); + } + pinned_total++; + } + } + fprintf(stderr, "[HOT] Pinned %d experts (top-%d/layer) after %d warmup tokens\n", pinned_total, m->hot_n, m->freq_token_count); +} + +/* COLIBRI_RESIDENT: after prefill (mode 1) and continuously through decode (mode 2), + * pin every expert this prompt routed to, so the CPU LRU never evicts their RAM + * slots. `quiet` suppresses the log line when nothing new was pinned + * (used for the per-token mid-decode calls). A per-layer pin budget = cap prevents + * pinning more experts than fit in the cache (which would deadlock the LRU). */ +static int apply_resident(Model *m, int quiet) { + Cfg *c = &m->c; + int newly = 0, over = 0; + for (int l = 0; l < c->n_layers; l++) { + uint8_t *row = m->seen + (int64_t)l * c->n_experts; + int cap = m->cache[l].cap; + int already = 0; + for (int e = 0; e < c->n_experts; e++) if (m->is_pinned[l * c->n_experts + e]) already++; + int budget = cap - already; /* free pin slots in this layer */ + int seen = 0; + for (int e = 0; e < c->n_experts; e++) { + if (!row[e]) continue; + seen++; + if (m->is_pinned[l * c->n_experts + e]) continue; /* already pinned */ + if (budget <= 0) { over++; continue; } /* layer full, skip */ + m->is_pinned[l * c->n_experts + e] = 1; + newly++; budget--; + } + if (seen > cap) over += seen - cap; + LCache *lc = &m->cache[l]; + for (int i = 0; i < lc->n; i++) + if (lc->slots[i].eid >= 0 && row[lc->slots[i].eid]) + lc->slots[i].pinned = 1; + } + if (!quiet || newly > 0) + fprintf(stderr, "[RESIDENT] Pinned %d new experts (CPU no-evict -> GPU resident)%s\n", + newly, over > 0 ? " | WARN: exceed per-layer cap, raise cap for full coverage" : ""); + return newly; +} + +/* ---------- RoPE: applied to the FIRST rope_dim dims of each head (Qwen3 partial rope) ---------- */ +static void rope_head_partial(float *x, int pos, int rope_dim, int head_dim, float theta) { + int h = rope_dim / 2; + for (int j = 0; j < h; j++) { + float inv = powf(theta, -2.0f * j / rope_dim); + float ang = pos * inv, cs = cosf(ang), sn = sinf(ang); + float a = x[j], b = x[j+h]; + x[j] = a*cs - b*sn; + x[j+h] = b*cs + a*sn; + } +} + +/* Gated Attention (GQA) matching HF Qwen3_5MoeAttention: + * - q_proj outputs query(head_dim) ++ attn_output_gate(head_dim); k/v are head_dim. + * - per-head q/k RMSNorm (weight [head_dim], 1.0+weight). + * - partial RoPE on the first rotary_dim dims of each head (text: mRoPE == standard). + * - scale = head_dim^-0.5; GQA repeat_kv. + * - attn_out = attn_out * sigmoid(gate), then o_proj (input dim = q_heads*head_dim). */ +static void attention(Model *m, Layer *l, int layer, float *x, int S, int pos_base, float *out) { + Cfg *c = &m->c; + int H = c->q_heads, KV = c->kv_heads, hd = c->head_dim, D = c->hidden; + int kvd = c->k_head_dim; + int qdim = c->q_head_dim; /* per-head q total (query+gate) */ + int q_out = H * qdim; /* q_proj output dim */ + int kv_out = KV * kvd; /* k/v_proj output dim */ + int q_per_kv = H / KV; + int rotary = c->rotary_dim; + /* HF always chunks q_proj output into query(head_dim) ++ gate(head_dim), + * regardless of the attn_output_gate config flag -- so split whenever the + * q per-head dim exceeds the (k/v) head dim. */ + int gate_dim = (qdim > hd) ? (qdim - hd) : 0; + float *q = falloc((int64_t)S*q_out); + float *k = falloc((int64_t)S*kv_out); + float *vv= falloc((int64_t)S*kv_out); + matmul_d(q, x, l->q, S, D, q_out); + matmul_d(k, x, l->k, S, D, kv_out); + matmul_d(vv, x, l->v, S, D, kv_out); + /* split q into query (first hd) and gate (next gate_dim), both per head */ + float *query = falloc((int64_t)S*H*hd); + float *gate = falloc((int64_t)S*H*gate_dim); + for (int s = 0; s < S; s++) { + for (int hh = 0; hh < H; hh++) { + const float *qs = q + (int64_t)s*q_out + hh*qdim; + memcpy(query + ((int64_t)s*H + hh)*hd, qs, hd*sizeof(float)); + if (gate_dim) memcpy(gate + ((int64_t)s*H + hh)*gate_dim, qs + hd, gate_dim*sizeof(float)); + } + } + for (int s = 0; s < S; s++) { + for (int hh = 0; hh < H; hh++) { + float *qh = query + ((int64_t)s*H + hh)*hd; + if (l->qn) rmsnorm_row(qh, qh, l->qn, hd, c->eps); + rope_head_partial(qh, pos_base + s, rotary, hd, c->theta); + } + for (int kvh = 0; kvh < KV; kvh++) { + float *kh = k + (int64_t)s*KV*kvd + kvh*kvd; + if (l->kn) rmsnorm_row(kh, kh, l->kn, kvd, c->eps); + rope_head_partial(kh, pos_base + s, rotary, kvd, c->theta); + } + } + for (int s = 0; s < S; s++) for (int kvh = 0; kvh < KV; kvh++) { + int t = pos_base + s; + memcpy(m->K[layer] + ((int64_t)kvh*m->max_t + t)*kvd, k + (int64_t)s*KV*kvd + kvh*kvd, kvd*sizeof(float)); + memcpy(m->V[layer] + ((int64_t)kvh*m->max_t + t)*kvd, vv + (int64_t)s*KV*kvd + kvh*kvd, kvd*sizeof(float)); + } + float scale = 1.f / sqrtf((float)hd); + float *ctx = falloc((int64_t)S*H*hd); + #pragma omp parallel for collapse(2) schedule(static) + for (int hh = 0; hh < H; hh++) { + for (int s = 0; s < S; s++) { + int kvh = hh / q_per_kv; + int qpos = pos_base + s; + const float *qv = query + ((int64_t)s*H + hh)*hd; + float sc[8192]; + for (int t = 0; t <= qpos; t++) { + const float *kv = m->K[layer] + ((int64_t)kvh*m->max_t + t)*kvd; + float acc = 0; for (int dd = 0; dd < kvd; dd++) acc += qv[dd]*kv[dd]; + sc[t] = acc * scale; + } + softmax_row(sc, qpos+1); + float *cx = ctx + ((int64_t)s*H + hh)*hd; + for (int dd = 0; dd < kvd; dd++) cx[dd] = 0; + for (int t = 0; t <= qpos; t++) { + const float *vrow = m->V[layer] + ((int64_t)kvh*m->max_t + t)*kvd; + float a = sc[t]; for (int dd = 0; dd < kvd; dd++) cx[dd] += a * vrow[dd]; + } + } + } + /* apply attn_output_gate: attn_out *= sigmoid(gate) */ + float *ag = falloc((int64_t)S*H*hd); + for (int s = 0; s < S; s++) for (int hh = 0; hh < H; hh++) for (int dd = 0; dd < hd; dd++) { + int o = ((int64_t)s*H + hh)*hd + dd; + float g = gate_dim ? gate[o] : 0.f; + ag[o] = ctx[o] * (1.f / (1.f + expf(-g))); + } + matmul_d(out, ag, l->o, S, H*hd, D); + free(q); free(k); free(vv); free(query); free(gate); free(ctx); free(ag); +} + +/* MoE: grouped top-k routing (+ optional router bias) + shared expert. + * Mirrors HF Qwen3 MoE: softmax(gate), optional group-limited top-k, normalized + * weights, sum routed experts, then add the un-gated shared expert. */ +static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out) { + Cfg *c = &m->c; int D = c->hidden, E = c->n_experts, K = c->topk, I = c->inter; + float *logits = falloc((int64_t)S*E); + double _tr = tm_on() ? tm_now() : 0.0; + matmul_d(logits, x, l->gate, S, D, E); + if (tm_on()) tm_add(S, 4, tm_now()-_tr); + if (c->has_bias && l->gate_bias) { + for (int s = 0; s < S; s++) { float *pr = logits + (int64_t)s*E; for (int e = 0; e < E; e++) pr[e] += l->gate_bias[e]; } + } + memset(out, 0, (int64_t)S*D*sizeof(float)); + float *g = falloc(I), *u = falloc(I), *hh = falloc(D); + float *sh = falloc(I), *shu = falloc(I), *shd = falloc(D); /* shared expert scratch */ + for (int s = 0; s < S; s++) { + float *pr = logits + (int64_t)s*E; + if (m->momentum_logits && m->pilot_smooth > 0.f) { + float *ema = m->momentum_logits + (int64_t)layer * E; + int is_zero = 1; for (int e = 0; e < E; e++) if (ema[e] != 0.f) { is_zero = 0; break; } + if (is_zero) { for (int e = 0; e < E; e++) ema[e] = pr[e]; } + else { for (int e = 0; e < E; e++) ema[e] = (1.f - m->pilot_smooth)*pr[e] + m->pilot_smooth*ema[e]; } + } + softmax_row(pr, E); + /* group-limited top-k selection */ + uint8_t keep[1024]; int Ec = E < 1024 ? E : 1024; + if (c->n_group > 1 && c->n_group <= Ec) { + int per = E / c->n_group; + float gs[1024]; + for (int gi = 0; gi < c->n_group; gi++) { + float b1 = -1e30f, b2 = -1e30f; + for (int e = gi*per; e < gi*per+per; e++) { float v = pr[e]; if (v > b1) { b2=b1; b1=v; } else if (v > b2) b2=v; } + gs[gi] = b1 + b2; + } + uint8_t gkeep[1024] = {0}; + for (int kk = 0; kk < c->topk_group; kk++) { + int bg = -1; float bv = -1e30f; + for (int gi = 0; gi < c->n_group; gi++) { if (!gkeep[gi] && gs[gi] > bv) { bv = gs[gi]; bg = gi; } } + if (bg < 0) break; gkeep[bg] = 1; + } + for (int e = 0; e < Ec; e++) keep[e] = 0; + for (int gi = 0; gi < c->n_group; gi++) if (gkeep[gi]) for (int e = gi*per; e < gi*per+per; e++) keep[e] = 1; + } else { + for (int e = 0; e < Ec; e++) keep[e] = 1; + } + int idx[256]; float val[256]; + for (int kk = 0; kk < K; kk++) { + int best = -1; float bv = -1e30f; + for (int e = 0; e < E; e++) { + if (!keep[e]) continue; + int taken = 0; for (int j = 0; j < kk; j++) if (idx[j]==e){taken=1;break;} + if (!taken && pr[e] > bv) { bv = pr[e]; best = e; } + } + idx[kk] = best; val[kk] = bv; + } + if (m->resident_collecting) { + for (int kk = 0; kk < K; kk++) if (idx[kk] >= 0) m->seen[(int64_t)layer * E + idx[kk]] = 1; + } + /* HF renormalizes the top-k router weights unconditionally */ + { float sm=0; for (int kk=0;kk0) for (int kk=0;kkhot_pinned && m->freq) { + uint32_t *freq_l = m->freq + (int64_t)layer * E; + for (int kk = 0; kk < K; kk++) if (idx[kk] >= 0) freq_l[idx[kk]]++; + } + const float *xs = x + (int64_t)s*D; + { + for (int kk = 0; kk < K; kk++) { + Slot *e; expert_get(m, layer, idx[kk], &e); + matmul_q(g, xs, e->g, e->gs, D, I); + matmul_q(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_q(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]; + } + } + /* shared expert (SwiGLU), sigmoid-gated by shared_expert_gate */ + double _ts = 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()-_ts); + } + free(logits); free(g); free(u); free(hh); free(sh); free(shu); free(shd); +} + +/* Gated DeltaNet (linear_attention) forward — recurrent gated-delta-rule. + * Mirrors HF Qwen3_5MoeGatedDeltaNet with a carried causal-conv ring + recurrent + * state S[h]=[kdim,vdim]. The conv ring and S persist in m->DN_conv/rec[layer] + * across step() calls (prefill chunk -> decode tokens). Math validated + * torch-free against the prefill (zero-padded conv) path in tools/_ref_dn_stream.py. + * + * Per token: qkv=x@qkv^T; z=x@z^T; b=x@b^T; a=x@a^T; beta=sigmoid(b); + * g=-exp(A_log)*softplus(a+dt_bias); + * conv_out[c]=silu(sum_{kk} w[kk]*ring[kk] + w[convk-1]*qkv[c]); advance ring; + * split conv_out -> q_in/k_in/v_in; repeat_interleave q,k by rep; l2norm + * (q scaled by 1/sqrt(kdim)); recurrence S[h]*=exp(g); kv=k@S; delta=(v-kv)*beta; + * S+=k (x) delta; out=q@S; per-head Gated RMSNorm (plain weight) -> out_proj. */ +static void deltanet(Model *m, Layer *l, int layer, float *x, int S, int pos_base, float *out) { + (void)pos_base; + Cfg *c = &m->c; + int vh = c->dn_vheads, vk = c->dn_kheads, kdim = c->dn_kdim, vdim = c->dn_vdim; + int convk = c->dn_convk, conv_dim = c->dn_conv_dim; + int rep = vh / vk; + int key_dim_tot = vk * kdim; + int value_dim = vh * vdim; + float scale = 1.f / sqrtf((float)kdim); + int H = c->hidden; + + float *qkv = falloc(conv_dim); + float *z = falloc(value_dim); + float *b = falloc(vh); + float *a = falloc(vh); + float *beta= falloc(vh); + float *gg = falloc(vh); + float *conv_out = falloc(conv_dim); + float *q = falloc(vh * kdim); + float *k = falloc(vh * kdim); + float *outv = falloc(value_dim); + float *outr = falloc(value_dim); + float *kv = falloc(vdim); + float *delta = falloc(vdim); + + float *rec = m->DN_rec[layer]; /* [vh*kdim*vdim] */ + float *ring = m->DN_conv[layer]; /* [conv_dim*(convk-1)] */ + + for (int s = 0; s < S; s++) { + const float *xs = x + (int64_t)s * H; + extern double g_dn_sub[4]; + double _d0 = tm_on()? tm_now():0; + /* projections (single-token matmuls) */ + matmul_d(qkv, xs, l->dn_qkv, 1, H, conv_dim); + matmul_d(z, xs, l->dn_z, 1, H, value_dim); + matmul(b, xs, l->dn_b, 1, H, vh); + matmul(a, xs, l->dn_a, 1, H, vh); + if (tm_on() && S==1){ double t=tm_now(); g_dn_sub[0]+=t-_d0; _d0=t; } + for (int h = 0; h < vh; h++) { + beta[h] = 1.f / (1.f + expf(-b[h])); + gg[h] = -expf(l->dn_alog[h]) * softplus_f(a[h] + l->dn_dtbias[h]); + } + /* causal depthwise conv1d (groups=conv_dim, kernel=convk) with carried ring + * (serial: ~33k FLOP, an OpenMP fork/join would cost more) */ + for (int cc = 0; cc < conv_dim; cc++) { + const float *w = l->dn_conv + (int64_t)cc * convk; + const float *rg = ring + (int64_t)cc * (convk - 1); + float acc = 0.f; + for (int kk = 0; kk < convk - 1; kk++) acc += w[kk] * rg[kk]; + acc += w[convk - 1] * qkv[cc]; + conv_out[cc] = acc / (1.f + expf(-acc)); /* silu */ + } + /* advance ring: drop oldest, append current token's qkv */ + for (int cc = 0; cc < conv_dim; cc++) { + float *rg = ring + (int64_t)cc * (convk - 1); + for (int kk = 0; kk < convk - 2; kk++) rg[kk] = rg[kk + 1]; + rg[convk - 2] = qkv[cc]; + } + if (tm_on() && S==1){ double t=tm_now(); g_dn_sub[1]+=t-_d0; _d0=t; } + /* split into query/key (key_dim_tot each) + value (value_dim) */ + const float *q_in = conv_out; + const float *k_in = conv_out + key_dim_tot; + const float *v_in = conv_out + 2 * key_dim_tot; + /* repeat_interleave q/k by rep along head dim (vk heads -> vh heads). + * HF semantics (torch repeat_interleave): each key head is repeated + * `rep` consecutive times, so VALUE head h takes KEY head (h / rep). + * This is NOT h % vk. Verified against _ref_dn.py L245-247. */ + for (int h = 0; h < vh; h++) { + int vk_idx = h / rep; + memcpy(q + (int64_t)h * kdim, q_in + (int64_t)vk_idx * kdim, kdim * sizeof(float)); + memcpy(k + (int64_t)h * kdim, k_in + (int64_t)vk_idx * kdim, kdim * sizeof(float)); + } + /* per-head l2norm (+ scale q by 1/sqrt(kdim)); eps 1e-6 inside sqrt (HF default) */ + for (int oh = 0; oh < vh; oh++) { + float *qh = q + (int64_t)oh * kdim; + double sq = 1e-6; for (int d = 0; d < kdim; d++) sq += (double)qh[d] * qh[d]; + double nq = sqrt(sq); + for (int d = 0; d < kdim; d++) qh[d] = (float)((double)qh[d] / nq * scale); + float *kh = k + (int64_t)oh * kdim; + double sk = 1e-6; for (int d = 0; d < kdim; d++) sk += (double)kh[d] * kh[d]; + double nk = sqrt(sk); + for (int d = 0; d < kdim; d++) kh[d] = (float)((double)kh[d] / nk); + } + /* recurrent gated delta rule over the value heads (heads are + * independent -> parallel; kv/delta thread-local) */ + #pragma omp parallel for schedule(static) + for (int h = 0; h < vh; h++) { + float kvl[512], dl[512]; /* vdim <= 512 */ + float *Sh = rec + (int64_t)h * kdim * vdim; + float egh = expf(gg[h]); + for (int t = 0; t < kdim * vdim; t++) Sh[t] *= egh; + const float *kd = k + (int64_t)h * kdim; + const float *vd = v_in + (int64_t)h * vdim; + /* kv = kd @ Sh (length vdim) */ + for (int vv = 0; vv < vdim; vv++) kvl[vv] = 0.f; + for (int kk = 0; kk < kdim; kk++) { + float kkd = kd[kk]; const float *Sr = Sh + (int64_t)kk * vdim; + for (int vv = 0; vv < vdim; vv++) kvl[vv] += kkd * Sr[vv]; + } + /* delta = (v - kv) * beta */ + for (int vv = 0; vv < vdim; vv++) dl[vv] = (vd[vv] - kvl[vv]) * beta[h]; + /* Sh += outer(kd, delta) */ + for (int kk = 0; kk < kdim; kk++) { + float kkd = kd[kk]; float *Sr = Sh + (int64_t)kk * vdim; + for (int vv = 0; vv < vdim; vv++) Sr[vv] += kkd * dl[vv]; + } + /* out = qd @ Sh */ + const float *qd = q + (int64_t)h * kdim; + float *ov = outv + (int64_t)h * vdim; + for (int vv = 0; vv < vdim; vv++) ov[vv] = 0.f; + for (int kk = 0; kk < kdim; kk++) { + float qkd = qd[kk]; const float *Sr = Sh + (int64_t)kk * vdim; + for (int vv = 0; vv < vdim; vv++) ov[vv] += qkd * Sr[vv]; + } + } + if (tm_on() && S==1){ double t=tm_now(); g_dn_sub[2]+=t-_d0; _d0=t; } + /* per-head Gated RMSNorm (plain weight, r=1/sqrt(mean+eps)) then silu(z) gate, then out_proj. + * HF Qwen3_5MoeRMSNormGated: out = (o*r)*weight * silu(z) = (o*r)*weight * z/(1+e^-z). + * NB: it is silu (z in numerator), NOT sigmoid. */ + #pragma omp parallel for schedule(static) + for (int h = 0; h < vh; h++) { + const float *o = outv + (int64_t)h * vdim; + const float *zr = z + (int64_t)h * vdim; + const float *w = l->dn_norm; + double ms = 0; for (int d = 0; d < vdim; d++) ms += (double)o[d] * o[d]; + float r = 1.f / sqrtf((float)(ms / vdim) + c->eps); + for (int d = 0; d < vdim; d++) { + float val = o[d] * r * w[d]; + outr[(int64_t)h * vdim + d] = val * zr[d] / (1.f + expf(-zr[d])); + } + } + matmul_d(out + (int64_t)s * H, outr, l->dn_out, 1, value_dim, H); + if (tm_on() && S==1){ g_dn_sub[3]+=tm_now()-_d0; } + if (layer == 0 && s == 0 && getenv("DN_DBG")) { + FILE *dbg = fopen(getenv("DN_DBG"), "wb"); + if (dbg) { + fwrite(conv_out, sizeof(float), conv_dim, dbg); + fwrite(q, sizeof(float), (int64_t)vh * kdim, dbg); + fwrite(outv, sizeof(float), value_dim, dbg); + fwrite(z, sizeof(float), value_dim, dbg); + fwrite(outr, sizeof(float), value_dim, dbg); + fwrite(out + (int64_t)s * H, sizeof(float), H, dbg); + fwrite(b, sizeof(float), vh, dbg); + fwrite(a, sizeof(float), vh, dbg); + fwrite(beta, sizeof(float), vh, dbg); + fwrite(gg, sizeof(float), vh, dbg); + fclose(dbg); + } + } + } + free(qkv); free(z); free(b); free(a); free(beta); free(gg); + free(conv_out); free(q); free(k); free(outv); free(outr); free(kv); free(delta); +} + +static float *step(Model *m, const int *ids, int S, int pos_base) { + Cfg *c = &m->c; int D = c->hidden; + if (m->resident_mode && m->first_step) m->resident_collecting = 1; + /* Per-layer residual dump (last token) for torch-free cosine debugging. + * Set DUMP_LAYERS= to write n_layers * D raw float32 rows. */ + FILE *lf = NULL; const char *lfn = getenv("DUMP_LAYERS"); + if (lfn) { lf = fopen(lfn, "wb"); if (!lf) fprintf(stderr, "DUMP_LAYERS: cannot open %s\n", lfn); } + if (g_pilot && m->token_count > 0) { + pthread_mutex_lock(&g_pilot_mx); + memset(m->is_queued, 0, (size_t)c->n_layers * c->n_experts); + pthread_mutex_unlock(&g_pilot_mx); + } + float *x = falloc((int64_t)S*D); + for (int s = 0; s < S; s++) memcpy(x + (int64_t)s*D, m->embed + (int64_t)ids[s]*D, D*sizeof(float)); + float *nrm = falloc((int64_t)S*D), *tmp = falloc((int64_t)S*D); + for (int i = 0; i < c->n_layers; i++) { + Layer *l = &m->L[i]; + for (int s = 0; s < S; s++) rmsnorm_row(nrm + (int64_t)s*D, x + (int64_t)s*D, l->in_ln, D, c->eps); + double _t0 = tm_on() ? tm_now() : 0.0; + if (c->is_attn[i]) { + attention(m, l, i, nrm, S, pos_base, tmp); + if (tm_on()) tm_add(S, 1, tm_now()-_t0); + } else { + deltanet(m, l, i, nrm, S, pos_base, tmp); + if (tm_on()) tm_add(S, 0, tm_now()-_t0); + } + if (lf) fwrite(tmp + (int64_t)(S-1)*D, sizeof(float), D, lf); /* sublayer output */ + for (int64_t j = 0; j < (int64_t)S*D; j++) x[j] += tmp[j]; + if (lf) fwrite(x + (int64_t)(S-1)*D, sizeof(float), D, lf); /* post-deltanet residual */ + if (g_pilot >= 1 && S <= 8 && i + 1 < c->n_layers) + pilot_prefetch(m, i + 1, x, S); + for (int s = 0; s < S; s++) rmsnorm_row(nrm + (int64_t)s*D, x + (int64_t)s*D, l->post_ln, D, c->eps); + _t0 = tm_on() ? tm_now() : 0.0; + moe(m, l, i, nrm, S, tmp); + if (tm_on()) tm_add(S, 2, tm_now()-_t0); + for (int64_t j = 0; j < (int64_t)S*D; j++) x[j] += tmp[j]; + if (lf) fwrite(x + (int64_t)(S-1)*D, sizeof(float), D, lf); + if (g_pilot >= 2 && S <= 8 && i + 2 < c->n_layers) + pilot_prefetch(m, i + 2, x, S); + if (g_pilot >= 3 && S <= 8 && i + 3 < c->n_layers) + pilot_prefetch(m, i + 3, x, S); + } + m->token_count += S; m->freq_token_count += S; + if (!m->hot_pinned && m->hot_n > 0 && m->freq_token_count >= m->warmup_tokens) pin_hot_experts(m); + m->kv_len = pos_base + S; + float *last = falloc(D); + rmsnorm_row(last, x + (int64_t)(S-1)*D, m->final_norm, D, c->eps); + float *logit = falloc(c->vocab); + double _th = tm_on() ? tm_now() : 0.0; + matmul_d(logit, last, m->lm_head, 1, D, c->vocab); + if (tm_on()) { tm_add(S, 5, tm_now()-_th); if (S==1) g_tm_dec_tokens++; else g_tm_pre_tokens += S; } + free(x); free(nrm); free(tmp); free(last); + if (lf) fclose(lf); + if (m->resident_collecting) { + int prefill_end = m->first_step; + apply_resident(m, prefill_end ? 0 : 1); /* always report after prefill; quiet mid-decode */ + if (prefill_end) { + m->first_step = 0; + if (m->resident_mode < 2) m->resident_collecting = 0; /* mode 1: stop after prefill */ + /* mode 2: keep collecting through decode for incremental pin */ + } + } + return logit; +} + +static void pilot_realload(Model *m, int layer, int eid) { + LCache *lc = &m->cache[layer]; Cfg *c = &m->c; + pthread_mutex_lock(&g_pilot_mx); + if (!m->is_queued[layer * c->n_experts + eid]) { pthread_mutex_unlock(&g_pilot_mx); return; } + for (int i = 0; i < lc->n; i++) if (lc->slots[i].eid == eid) { m->is_queued[layer*c->n_experts+eid]=0; pthread_mutex_unlock(&g_pilot_mx); return; } + Slot *s; + if (lc->n < lc->cap) { s = &lc->slots[lc->n++]; slot_ensure_allocated(m, s); } + else { + int lru = -1; + for (int i = 0; i < lc->n; i++) { if (lc->slots[i].pinned || lc->slots[i].eid < 0) continue; if (lru < 0 || lc->slots[i].used < lc->slots[lru].used) lru = i; } + if (lru < 0) { m->is_queued[layer*c->n_experts+eid]=0; pthread_mutex_unlock(&g_pilot_mx); return; } + s = &lc->slots[lru]; s->pinned = 0; + } + s->eid = -1; s->used = ++m->clock; + pthread_mutex_unlock(&g_pilot_mx); + load_expert_merged(m, layer, eid, s); + pthread_mutex_lock(&g_pilot_mx); + s->eid = eid; s->pinned = m->is_pinned[layer*c->n_experts+eid]; s->used = ++m->clock; + m->is_queued[layer*c->n_experts+eid] = 0; pthread_mutex_unlock(&g_pilot_mx); +} + +static void *pilot_worker(void *arg) { + (void)arg; + while (1) { + unsigned r = __atomic_load_n(&pilot_r, __ATOMIC_ACQUIRE); + unsigned w = __atomic_load_n(&pilot_w, __ATOMIC_ACQUIRE); + if (r == w) { sleep_ms(1); continue; } + int layer = pilot_q[r & 4095].l, eid = pilot_q[r & 4095].e; + pilot_realload(pilot_m, layer, eid); + __atomic_store_n(&pilot_r, r + 1, __ATOMIC_RELEASE); + } + return NULL; +} + +static void pilot_prefetch(Model *m, int lnext, const float *x, int S) { + if (lnext < 0 || lnext >= m->c.n_layers) return; + Cfg *c = &m->c; int D = c->hidden, E = c->n_experts; + ensure_pilot_worker_started(m); + float *logits = falloc((int64_t)S * E); + Layer *l = &m->L[lnext]; + float *nrm_x = falloc((int64_t)S * D); + for (int s = 0; s < S; s++) rmsnorm_row(nrm_x + (int64_t)s*D, x + (int64_t)s*D, l->post_ln, D, c->eps); + matmul_d(logits, nrm_x, l->gate, S, D, E); /* int8 copy (f32 may be freed) */ + free(nrm_x); + for (int s = 0; s < S; s++) { + float *pr = logits + (int64_t)s*E; + float *blended = pr; + float *ema = m->momentum_logits + (int64_t)lnext*E; + if (m->pilot_smooth > 0.f) { + blended = falloc(E); int is_zero = 1; + for (int e = 0; e < E; e++) if (ema[e] != 0.f) { is_zero = 0; break; } + if (is_zero) { for (int e = 0; e < E; e++) { ema[e] = pr[e]; blended[e] = pr[e]; } } + else { for (int e = 0; e < E; e++) { blended[e] = (1.f-m->pilot_smooth)*pr[e] + m->pilot_smooth*ema[e]; ema[e] = blended[e]; } } + } + int cand = 0; int idx[128]; + float max_logit = -1e30f; for (int e = 0; e < E; e++) if (blended[e] > max_logit) max_logit = blended[e]; + float *exps = falloc(E); float sum_exps = 0.f; + for (int e = 0; e < E; e++) { exps[e] = expf(blended[e] - max_logit); sum_exps += exps[e]; } + float cum_sum = 0.f; int min_cand = c->topk; int max_cand = c->topk * g_wide; + if (max_cand < min_cand) max_cand = min_cand; if (max_cand > 128) max_cand = 128; if (max_cand > E) max_cand = E; + for (int kk = 0; kk < max_cand; kk++) { + int best = -1; float bv = -1.f; + for (int e = 0; e < E; e++) { int taken = 0; for (int j = 0; j < kk; j++) if (idx[j]==e){taken=1;break;} if (!taken && exps[e] > bv) { bv = exps[e]; best = e; } } + if (best < 0) break; + idx[kk] = best; cum_sum += bv; cand++; + if (cum_sum >= m->pilot_conf_limit * sum_exps && cand >= min_cand) break; + } + free(exps); + if (blended != pr) free(blended); + for (int a = 0; a < cand-1; a++) for (int b = a+1; b < cand; b++) + 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; } + pthread_mutex_unlock(&g_pilot_mx); + if (!found) { + int gidx = lnext*E + eid; + pthread_mutex_lock(&g_pilot_mx); int already_queued = m->is_queued[gidx]; + if (!already_queued) m->is_queued[gidx] = 1; + pthread_mutex_unlock(&g_pilot_mx); + if (!already_queued) { + unsigned w2 = __atomic_load_n(&pilot_w, __ATOMIC_RELAXED); + unsigned r2 = __atomic_load_n(&pilot_r, __ATOMIC_ACQUIRE); + if (w2 - r2 < 4096) { pilot_q[w2 & 4095].l = lnext; pilot_q[w2 & 4095].e = eid; __atomic_store_n(&pilot_w, w2+1, __ATOMIC_RELEASE); } + else { pthread_mutex_lock(&g_pilot_mx); m->is_queued[gidx] = 0; pthread_mutex_unlock(&g_pilot_mx); } + } + } + } + } + free(logits); +} + +/* When DUMP= is set, generate() copies the last-token logits here so main() + * can write them to (raw float32, length = vocab). Lets a torch-free + * cosine comparison against tools/_ref_dn.py's numpy logits validate the port. */ +static float *g_last_logit = NULL; + +/* Zero the DeltaNet recurrent state so a new request doesn't inherit the + * previous conversation's hidden state. Must be called at the start of every + * generation (the CLI runs once, so this is also correct there). */ +static void reset_recurrent(Model *m){ + Cfg *c = &m->c; + for (int i = 0; i < c->n_layers; i++){ + if (c->is_attn[i]) continue; + if (m->DN_rec[i]) memset(m->DN_rec[i], 0, (size_t)c->dn_vheads * c->dn_kdim * c->dn_vdim * sizeof(float)); + if (m->DN_conv[i]) memset(m->DN_conv[i], 0, (size_t)c->dn_conv_dim * (c->dn_convk - 1) * sizeof(float)); + } +} + +/* Allocate (once) or reuse the KV cache across requests. Grows only when a + * longer context is needed; never shrinks. Frees the previous buffers on + * growth so the server doesn't leak KV memory across requests. */ +static void ensure_kv(Model *m){ + Cfg *c = &m->c; + if (m->kv_cap >= m->max_t && m->K) return; + if (m->K){ + for (int i = 0; i < c->n_layers; i++){ if (m->K[i]) free(m->K[i]); if (m->V[i]) free(m->V[i]); } + free(m->K); free(m->V); m->K = NULL; m->V = NULL; + } + m->K = calloc(c->n_layers, sizeof(float*)); m->V = calloc(c->n_layers, sizeof(float*)); + for (int i = 0; i < c->n_layers; i++){ + if (c->is_attn[i]){ + m->K[i] = falloc((int64_t)c->kv_heads * m->max_t * c->k_head_dim); + m->V[i] = falloc((int64_t)c->kv_heads * m->max_t * c->k_head_dim); + } else { m->K[i] = NULL; m->V[i] = NULL; } + } + m->kv_cap = m->max_t; +} + +static void generate(Model *m, const int *prompt, int np, int n_new, int *out) { + Cfg *c = &m->c; + m->max_t = np + n_new; + reset_recurrent(m); + ensure_kv(m); + m->kv_len = 0; + for (int i = 0; i < np; i++) out[i] = prompt[i]; + float *logit = step(m, prompt, np, 0); + int len = np; + for (int s = 0; s < n_new; s++) { + int best = 0; float bv = logit[0]; + for (int i = 1; i < c->vocab; i++) if (logit[i] > bv) { bv = logit[i]; best = i; } + if (s == 0 && g_ttft < 0) g_ttft = now_s() - g_gen_t0; /* record TTFT */ + if (g_stream) { stream_token(best); fflush(stdout); } + if (s == n_new - 1) { + if (getenv("DUMP")) { + g_last_logit = malloc((size_t)c->vocab * sizeof(float)); + memcpy(g_last_logit, logit, (size_t)c->vocab * sizeof(float)); + } + free(logit); out[len++] = best; break; + } + free(logit); out[len++] = best; + int one = best; + { extern double g_tm_step; double _s0 = tm_on()? tm_now():0; + logit = step(m, &one, 1, len - 1); + if (tm_on()) g_tm_step += tm_now()-_s0; } + } +} + +static int tf_nll(Model *m, const int *full, int nfull, int np, double *nll_out) { + Cfg *c = &m->c; + m->max_t = nfull; + reset_recurrent(m); + ensure_kv(m); + m->kv_len = 0; + double nll = 0; int scored = 0; + float *logit = step(m, full, np, 0); + for (int i = np; i < nfull; i++) { + float mx = logit[0]; for (int v = 1; v < c->vocab; v++) if (logit[v] > mx) mx = logit[v]; + double Z = 0; for (int v = 0; v < c->vocab; v++) Z += exp((double)logit[v] - mx); + nll += -((double)logit[full[i]] - mx - log(Z)); + scored++; + free(logit); logit = NULL; + if (i == nfull - 1) break; + logit = step(m, &full[i], 1, i); + } + if (logit) free(logit); + *nll_out = nll / scored; + return scored; +} + +static int *read_int_array(jval *o, const char *key, int *n_out) { + jval *a = json_get(o, key); + if (!a || a->t != J_ARR) { fprintf(stderr, "ref.json: missing array \"%s\"\n", key); exit(1); } + int *r = malloc(a->len * sizeof(int)); + for (int i = 0; i < a->len; i++) r[i] = (int)a->kids[i]->num; + *n_out = a->len; return r; +} + +#ifndef QWEN36_NO_MAIN +int main(int argc, char **argv) { + const char *snap = getenv("SNAP"); + if (!snap) { fprintf(stderr, "set SNAP=\n"); return 1; } + g_pilot = getenv("PILOT") ? atoi(getenv("PILOT")) : 0; + g_wide = getenv("WIDE") ? atoi(getenv("WIDE")) : 1; + if (g_wide < 1) g_wide = 1; if (g_wide > 4) g_wide = 4; + if (getenv("OPENAI")) g_openai = 1; /* OpenAI-compatible output */ + const char *mv = getenv("MODEL"); if (mv && *mv) g_model = mv; + int hot_n = getenv("HOT") ? atoi(getenv("HOT")) : 0; + int cap = argc > 1 ? atoi(argv[1]) : 16; + int bits = argc > 2 ? atoi(argv[2]) : 4; + if (bits < 2 || bits > 8) { fprintf(stderr, "quant_bits must be 2..8 (got %d)\n", bits); return 1; } + const char *refpath = argc > 3 ? argv[3] : "ref.json"; + + float smooth = getenv("SMOOTH") ? (float)atof(getenv("SMOOTH")) : 0.3f; + float conf = getenv("CONF_LIMIT") ? (float)atof(getenv("CONF_LIMIT")) : 0.92f; + + fprintf(stderr, "== qwen36 Phase-2 engine | cache=%d/layer bits=%d pilot=%d wide=%d hot=%d smooth=%.2f conf=%.2f ==\n", + cap, bits, g_pilot, g_wide, hot_n, smooth, conf); + + + int is_ref = 0; + int rplen = (int)strlen(refpath); + if (rplen>=5 && strcmp(refpath+rplen-5, ".json")==0) is_ref = 1; + + int *prompt=NULL, *full=NULL, *out=NULL; + int np=0, nfull=0, n_new=0; + char *buf=NULL, *arena=NULL; + + /* load tokenizer early so text-prompt mode can encode before model_init */ + { + const char *tokpath = getenv("TOK"); + if (tokpath && *tokpath) load_tokenizer(tokpath); + else if (argc > 4 && argv[4] && *argv[4]) load_tokenizer(argv[4]); + else { char tpb[2048]; snprintf(tpb,sizeof tpb,"%s/tokenizer.json",snap); load_tokenizer(tpb); } + } + + if (is_ref) { + FILE *f = fopen(refpath, "rb"); if (!f) { perror(refpath); return 1; } + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + buf=malloc(n+1); if (fread(buf,1,n,f)!=(size_t)n) {} buf[n]=0; fclose(f); + jval *ref = json_parse(buf, &arena); + prompt = read_int_array(ref,"prompt_ids",&np); + full = read_int_array(ref,"full_ids",&nfull); + n_new = nfull - np; + } else { + /* text-prompt mode: read file as raw text, encode in C */ + FILE *f = fopen(refpath, "rb"); if (!f) { perror(refpath); return 1; } + fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + char *txt=malloc(n+1); if (fread(txt,1,n,f)!=(size_t)n) {} txt[n]=0; fclose(f); + if (!g_tok) { fprintf(stderr, "[enc] no tokenizer loaded; cannot encode text. Put tokenizer.json in SNAP or set TOK.\n"); free(txt); return 1; } + encode_text(txt, &prompt, &np); + free(txt); + n_new = getenv("N_NEW") ? atoi(getenv("N_NEW")) : 64; + if (n_new < 1) n_new = 1; + fprintf(stderr, "[enc] prompt tokens: %d | generating %d new tokens\n", np, n_new); + if (getenv("ENC_DEBUG") && np <= 300) { fprintf(stderr, "[enc] prompt ids: "); for (int i=0;ihidden; + int q_out = qc->q_heads * qc->q_head_dim, kv_out = qc->kv_heads * qc->k_head_dim; + for (int i = 0; i < qc->n_layers; i++) { + Layer *l = &m.L[i]; + qdw_register(l->q, D2, q_out); qdw_register(l->k, D2, kv_out); + qdw_register(l->v, D2, kv_out); qdw_register(l->o, qc->o_in, D2); + qdw_register(l->gate, D2, qc->n_experts); + qdw_register(l->sh_g, D2, qc->shared_inter); qdw_register(l->sh_u, D2, qc->shared_inter); + qdw_register(l->sh_d, qc->shared_inter, D2); + qdw_register(l->dn_qkv, D2, qc->dn_conv_dim); + qdw_register(l->dn_z, D2, qc->dn_vheads * qc->dn_vdim); + qdw_register(l->dn_out, qc->dn_vheads * qc->dn_vdim, D2); + } + qdw_register(m.lm_head, D2, qc->vocab); + /* Free the f32 originals -- the pointers only serve as lookup keys in + * matmul_d from here on (never dereferenced again). + * COLI_KEEP_F32=1 keeps them (debug). */ + double freed = 0; + if (!getenv("COLI_KEEP_F32")) { + for (int i = 0; i < g_qdw_n; i++) { + freed += (double)g_qdw[i].I * g_qdw[i].O * sizeof(float); + free((void*)g_qdw[i].w); + } + } + fprintf(stderr, "[dense-i8] %d matrices quantized in %.1f s, %.1f GB f32 freed\n", + g_qdw_n, now_s()-tq, freed/1073741824.0); + } + + + if (is_ref && getenv("PPL") && atoi(getenv("PPL")) == 1) { + double nll; double t = now_s(); + int scored = tf_nll(&m, full, nfull, np, &nll); + double dt = now_s() - t; + double tot = m.hits + m.miss; + printf("TF-NLL: %.4f nats/token over %d tokens | ppl = %.2f\n", nll, scored, exp(nll)); + printf("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); + printf("Speed: %.2f tok/s (%.1fs for %d tokens) | PEAK RSS: %.2f GB\n", scored/dt, dt, scored, rss_gb()); + free(buf); free(arena); return 0; + } + + out = malloc((np + n_new) * sizeof(int)); + /* timing + OpenAI id setup (before generation) */ + g_ttft = -1; g_gen_t0 = now_s(); + if (g_openai){ + g_oa_created = (long)time(NULL); + snprintf(g_oa_id, sizeof g_oa_id, "chatcmpl-%ld%04d", g_oa_created, (int)(now_s()*1000) % 10000); + } + /* streaming text: emit tokens as they are produced (text mode + tokenizer only) */ + if (!is_ref && g_tok && !getenv("NOSTREAM")) { + g_stream = 1; g_sbn = 0; + if (g_openai){ + char jb[320]; + snprintf(jb, sizeof jb, + "{\"id\":\"%s\",\"object\":\"chat.completion.chunk\",\"created\":%ld,\"model\":\"%s\"," + "\"choices\":[{\"index\":0,\"delta\":{\"role\":\"assistant\"},\"finish_reason\":null}]}", + g_oa_id, g_oa_created, g_model); + sse_chunk(jb); + } else { + fprintf(stderr, "Generated (%d new tokens):\nText : ", n_new); fflush(stderr); + } + } + double t = now_s(); + generate(&m, prompt, np, n_new, out); + double dt = now_s() - t; + + /* DUMP=: write last-token logits (raw float32, vocab) for a torch-free + * cosine comparison against tools/_ref_dn.py --dump. */ + if (g_last_logit) { + const char *dp = getenv("DUMP"); + FILE *df = fopen(dp && *dp ? dp : "qwen36_logits.f32", "wb"); + if (df) { fwrite(g_last_logit, sizeof(float), (size_t)m.c.vocab, df); fclose(df); + fprintf(stderr, "[dump] wrote %d logits -> %s\n", m.c.vocab, dp && *dp ? dp : "qwen36_logits.f32"); } + else fprintf(stderr, "[dump] cannot open %s\n", dp ? dp : "qwen36_logits.f32"); + } + + if (is_ref) { + int match = 0; + printf("\nReference: "); for (int i=np;i= 0) fprintf(stderr, "TTFT: %.2f s (time to first token)\n", g_ttft); + tm_report(); + 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); + fprintf(stderr, "Speed: %.2f tok/s (%.1fs for %d tokens)\n", n_new/dt, dt, n_new); + free(buf); free(arena); + return 0; +} +#endif /* QWEN36_NO_MAIN */ diff --git a/c/tools/convert_qwen36.py b/c/tools/convert_qwen36.py new file mode 100644 index 000000000..c5126c003 --- /dev/null +++ b/c/tools/convert_qwen36.py @@ -0,0 +1,572 @@ +#!/usr/bin/env python3 +"""Convert Qwen3.6-35B-A3B (or any Qwen3.5/3.6 MoE) HF checkpoint -> colibri container. + +Phase-2 converter: it converts ALL layers (GATED-ATTENTION full_attention AND the +Gated DeltaNet linear_attention layers). Every layer also carries its MoE/MLP block. +DeltaNet layers are stored under `model.layers.{i}.linear_attn.*` (in_proj_qkv/z/b/a, +conv1d.weight, dt_bias, A_log, norm.weight, out_proj). The engine implements the +recurrent gated-delta-rule for these layers; in Phase 1 they were skipped as identity. + +Design notes (must stay in sync with c/qwen36.c): + * The colibri container is just a directory of safetensors shards (+ config.json + + qwen36_meta.json). The engine reads them lazily via st_init, paging expert weights + from disk -- that is the whole "run a huge model in little RAM" trick. + * Expert weights are stored per-expert as `model.layers.{a}.mlp.experts.{e}.merged_weight` + + `.qs` (f32 scales). For --ebits >= 5 the merged_weight is int8 (layout [gate|up|down]); + for --ebits <= 4 it is TRUE 4-bit packed uint8 (2 elements/byte, half size) -- see + pack_int4(). The qs layout (per-row f32 scales) is identical either way. This matches + what c/qwen36.c's load_expert_merged expects: it detects int4 by ON-DISK SIZE (N/2 bytes) + and unpacks in-place to int8, so the rest of the MoE path is unchanged. + * Attention + router + shared-expert + norms stay f16 (they are tiny vs experts). + * Real Qwen3.6 is a vision-language checkpoint: config dims live under `text_config`, + and weight keys are prefixed `model.language_model.`. Both are handled transparently. + * The fused expert tensors `mlp.experts.gate_up_proj` / `down_proj` are split per expert + into the merged_weight layout (gate_up = [gate; up] along dim 0). + * Head dims are derived from the actual weight shapes (authoritative), not from config + heuristics, because Qwen3.6's qk split (q_head_dim=512, k/v_head_dim=256, rope over + 256 dims) does not match naive head_dim=hidden/n_heads. + +Usage (cloud, e.g. Colab free CPU): + python tools/convert_qwen36.py --repo Qwen/Qwen3.6-35B-A3B --out ./qwen36_i4 --ebits 4 \\ + --low-disk --stream-upload --upload-repo minne100/qwen36-35b-a3b-colibri-i4 + +Usage (local tiny model for end-to-end testing): + python tools/convert_qwen36.py --model ../qwen36_tiny --out ../qwen36_tiny_i4 --ebits 8 +""" + +import argparse, json, math, os, struct, sys +from pathlib import Path + +# Windows: force UTF-8 output +if sys.platform == "win32": + for s in (sys.stdout, sys.stderr): + try: + s.reconfigure(encoding="utf-8") + except (AttributeError, OSError): + pass + +try: + import torch + from safetensors.numpy import safe_open as safe_open_np + from safetensors.torch import safe_open as safe_open_pt + from safetensors.torch import save_file +except ImportError as exc: + sys.exit(f"Missing dependencies: {exc}. Install: pip install torch safetensors") + + +def quantize_row(w: torch.Tensor, bits: int) -> tuple[torch.Tensor, torch.Tensor]: + """Row-wise symmetric quantization to `bits` (2..8). Storage is int8 regardless of + bits (engine dequantizes q*scale); for bits<8 values are confined to a smaller range.""" + qmax = (1 << (bits - 1)) - 1 + w_f32 = w.float() + row_max = w_f32.abs().amax(dim=1, keepdim=True).clamp(min=1e-12) + scales = row_max / qmax + q = (w_f32 / scales).round().clamp(-qmax - 1, qmax).to(torch.int8) + return q, scales.squeeze(1) + + +def pack_int4(q: "torch.Tensor") -> "torch.Tensor": + """Pack a 1D int-tensor (values already confined to the signed-4-bit range [-8,7]) + into a uint8 blob, 2 values per byte (LOW nibble = element 2k, HIGH nibble = element 2k+1). + + Convention (kept in sync with the c/qwen36.c int4 dequant path, TODO): + byte k = (signed4(element[2k+1]) << 4) | (signed4(element[2k]) & 0xF) + Each nibble is a SIGNED 4-bit integer in [-8,7]; dequant = nibble * scale[row]. + The whole expert merged blob (gate||up||down, flattened) is packed as one contiguous + uint8 array, so its length is exactly N/2 (N = element count, always even here). + """ + q = q.to(torch.int16).clamp(-8, 7) + n = q.numel() + if n & 1: + q = torch.cat([q, torch.zeros(1, dtype=torch.int16)]) + lo = (q[0::2] & 0xF).to(torch.uint8) + hi = ((q[1::2] & 0xF) << 4).to(torch.uint8) + return (lo | hi).contiguous() + + +def make_merged(gate, up, down, ebits): + """gate/up: [inter, H]; down: [H, inter] (torch, any fp). + -> (merged_weight, qs f32 1D). + + Storage format depends on ebits: + * ebits >= 5 : merged_weight is int8 (1 byte / element), as before. + * ebits <= 4 : merged_weight is TRUE 4-bit packed uint8 (2 elements / byte, + half the size) -- see pack_int4 for the nibble convention. + The engine (c/qwen36.c) currently reads ebits>=5 (int8); the int4 path is WIP. + qs (per-row f32 scales) is identical in both cases. + """ + def q(t): + qt, s = quantize_row(t.reshape(t.shape[0], -1), ebits) # rows along dim0 + return qt, s + gq, gs = q(gate) + uq, us = q(up) + dq, ds = q(down) + mw_i8 = torch.cat([gq.flatten(), uq.flatten(), dq.flatten()]).contiguous() + if ebits <= 4: + mw = pack_int4(mw_i8) # uint8, 2x4-bit per byte -> HALF size (true int4) + else: + mw = mw_i8.to(torch.int8) # int8 storage (ebits 5..8) + qs = torch.cat([gs, us, ds]).contiguous().float() + return mw, qs + + +def _unpack_int4(packed: "torch.Tensor") -> "torch.Tensor": + """Inverse of pack_int4: uint8 blob -> signed int8 values (HIGH nibble = element 2k+1).""" + b = packed.to(torch.uint8) + lo = (b & 0xF).to(torch.int8) + hi = ((b >> 4) & 0xF).to(torch.int8) + # sign-extend 4-bit -> 8-bit + lo = torch.where(lo >= 8, lo - 16, lo) + hi = torch.where(hi >= 8, hi - 16, hi) + return torch.stack([lo, hi], dim=1).flatten() + + +def _selftest(): + """Round-trip check for the true-int4 packing used by --ebits<=4.""" + import random + print("=== int4 pack/unpack selftest ===") + ok = True + for ebits in (2, 3, 4): + # small synthetic expert: gate[inter,H] up[inter,H] down[H,inter] + inter, H = 8, 16 + g = torch.randn(inter, H) * 3 + u = torch.randn(inter, H) * 3 + d = torch.randn(H, inter) * 3 + mw, qs = make_merged(g, u, d, ebits) + # serialize to safetensors + reload (exercises the real dtype path) + import tempfile, os + td = tempfile.mkdtemp() + from safetensors.torch import save_file, load_file + save_file({"merged_weight": mw, "qs": qs}, os.path.join(td, "e0.safetensors")) + back = load_file(os.path.join(td, "e0.safetensors")) + mw_r = back["merged_weight"] + # rebuild the int8 reference exactly as make_merged did (pre-pack) + def qref(t): + qt, _ = quantize_row(t.reshape(t.shape[0], -1), ebits) + return qt.flatten() + ref = torch.cat([qref(g), qref(u), qref(d)]).contiguous().to(torch.int8) + if mw_r.dtype != torch.uint8: + print(f" ebits={ebits}: FAIL dtype={mw_r.dtype} (expected uint8)") + ok = False + continue + rec = _unpack_int4(mw_r).to(torch.int8)[:ref.numel()] + maxerr = (rec.int() - ref.int()).abs().max().item() + int8_bytes = ref.numel() + int4_bytes = mw_r.numel() + exact = bool(maxerr == 0) + halved = bool(int4_bytes * 2 == int8_bytes) + print(f" ebits={ebits}: maxerr={maxerr} exact={exact} " + f"int4_bytes={int4_bytes} int8_bytes={int8_bytes} halved={halved}") + ok = ok and exact and halved + # also confirm ebits>=5 still stores int8 + g = torch.randn(8, 16) * 3; u = torch.randn(8, 16) * 3; d = torch.randn(16, 8) * 3 + mw8, _ = make_merged(g, u, d, 8) + print(f" ebits=8: dtype={mw8.dtype} (expected int8), bytes={mw8.numel()}") + ok = ok and (mw8.dtype == torch.int8) + print("SELFTEST", "PASS" if ok else "FAIL") + sys.exit(0 if ok else 1) + + +def resolve_prefix(keys): + """Return the layer-weight prefix so that `prefix + "layers.{i}."` matches the + actual weight keys. Handles the common layouts: + * 'model.' for `model.layers.0...` (standard HF Qwen3 MoE) + * 'model.language_model.' for `model.language_model.layers.0...` (some VL checkpoints) + """ + for k in keys: + if k.startswith("model.language_model.layers."): + return "model.language_model." + for k in keys: + if k.startswith("model.layers."): + return "model." + return "" + + +def main(): + ap = argparse.ArgumentParser(description="Convert Qwen3.6 MoE HF checkpoint -> colibri container") + src = ap.add_mutually_exclusive_group(required=False) + src.add_argument("--repo", help="HuggingFace repo ID (will be downloaded)") + src.add_argument("--model", help="Local HF checkpoint directory") + ap.add_argument("--out", required=False, help="Output container directory") + ap.add_argument("--ebits", type=int, default=4, help="Expert quant bits (2..8, default 4)") + ap.add_argument("--upload-repo", help="Push the finished container to this HF repo ID") + ap.add_argument("--hf-token", help="HF token (defaults to HF_TOKEN env / cached login)") + ap.add_argument("--low-disk", action="store_true", + help="Download source shards one at a time, delete each after its layers are " + "converted (only the generated container is kept).") + ap.add_argument("--stream-upload", action="store_true", + help="Upload each generated shard to --upload-repo immediately, then delete it " + "locally. Keeps local disk at ~one shard.") + ap.add_argument("--no-readme", action="store_true", help="Skip README.md generation") + ap.add_argument("--selftest", action="store_true", + help="Validate the int4 pack/unpack round-trip on synthetic weights, then exit " + "(no model conversion). Checks that packing recovers the exact quantized " + "integers and that int4 halves the storage vs int8.") + args = ap.parse_args() + + if args.selftest: + return _selftest() + if not (args.repo or args.model): + sys.exit("error: need --repo or --model (or --selftest)") + if not args.out: + sys.exit("error: --out is required") + + if not 2 <= args.ebits <= 8: + sys.exit(f"--ebits must be 2..8 (got {args.ebits})") + + token = args.hf_token or os.environ.get("HF_TOKEN") + if args.repo: + from huggingface_hub import hf_hub_download, HfApi + api = HfApi(token=token) + else: + api = None + + # ---- locate source ---- + if args.repo: + # download config + index only first (cheap) to learn structure + cfg_path = hf_hub_download(args.repo, "config.json", token=token) + try: + idx_path = hf_hub_download(args.repo, "model.safetensors.index.json", token=token) + except Exception: + idx_path = None + src_dir = None + else: + src_dir = Path(args.model) + if not (src_dir / "config.json").is_file(): + sys.exit(f"config.json missing in {src_dir}") + cfg_path = str(src_dir / "config.json") + idx_path = str(src_dir / "model.safetensors.index.json") if (src_dir / "model.safetensors.index.json").is_file() else None + + cfg_full = json.load(open(cfg_path, encoding="utf-8")) + # VL checkpoints nest dims under text_config + mcfg = cfg_full.get("text_config", cfg_full) + out = Path(args.out) + out.mkdir(parents=True, exist_ok=True) + import shutil + # The engine reads FLAT keys from config.json, but VL checkpoints nest them + # under text_config -> write the flattened text config (the original goes to + # config.hf.json for reference). Fixes "missing hidden_size" on VL models. + with open(out / "config.json", "w", encoding="utf-8") as cf: + json.dump(mcfg, cf, indent=1) + shutil.copy2(cfg_path, out / "config.hf.json") + print(f"config -> {out / 'config.json'} (flat; model_type={cfg_full.get('model_type')})") + # The engine needs tokenizer.json next to the shards (TOK= override exists, + # but a container should run out of the box). + try: + if args.repo: + tok_path = hf_hub_download(args.repo, "tokenizer.json", token=token) + else: + tok_path = str(src_dir / "tokenizer.json") + if Path(tok_path).is_file(): + shutil.copy2(tok_path, out / "tokenizer.json") + print(f"tokenizer -> {out / 'tokenizer.json'}") + else: + print("WARNING: tokenizer.json not found; the engine will need TOK=") + except Exception as e: + print(f"WARNING: could not fetch tokenizer.json ({e}); the engine will need TOK=") + + # ---- build weight map (key -> shard file) ---- + if idx_path: + wm = json.load(open(idx_path, encoding="utf-8"))["weight_map"] + else: + # no index: enumerate local safetensors + if args.repo: + sys.exit("remote repo without index.json is unsupported; download the repo first") + wm = {} + for sh in sorted(src_dir.glob("*.safetensors")): + with safe_open_np(str(sh), framework="np") as f: + for k in f.keys(): + wm[k] = sh.name + + # ---- detect prefix + layer count from weights ---- + sample_layer_keys = [k for k in wm if ".layers.0." in k or ".layers.0." in k] + prefix = resolve_prefix(wm.keys()) + # find max layer index present + import re + layer_ids = set(int(m) for m in re.findall(r"\." + re.escape(prefix) + r"layers\.(\d+)\.", " ".join(wm.keys()))) + n_layers_weight = (max(layer_ids) + 1) if layer_ids else mcfg.get("num_hidden_layers", 0) + + layer_types = mcfg.get("layer_types") + if layer_types is None: + # infer: every 4th layer is full_attention (Qwen3.5/3.6 convention) + n = mcfg.get("num_hidden_layers", n_layers_weight) + layer_types = ["full_attention" if i % 4 == 3 else "linear_attention" for i in range(n)] + full_idx = [i for i, lt in enumerate(layer_types) if lt == "full_attention"] + all_idx = list(range(len(layer_types))) + print(f"layers: {len(layer_types)} total, {len(full_idx)} full_attention (Phase 2: all {len(all_idx)} stored)") + + # ---- shard cache / eviction ---- + tmp = out / ".src_tmp" + tmp.mkdir(parents=True, exist_ok=True) + local_shards = {} # shard name -> local path + shard_handles = {} # shard name -> safe_open handle + shard_layers = {} # shard name -> set of layer ids needing it (-1 = globals) + + def layer_keys(i): + p = f"{prefix}layers.{i}." + return [k for k in wm if k.startswith(p)] + + for i in all_idx: + for k in layer_keys(i): + shard_layers.setdefault(wm[k], set()).add(i) + globals_keys = [k for k in (f"{prefix}embed_tokens.weight", + "lm_head.weight", f"{prefix}lm_head.weight", + f"{prefix}norm.weight") + if k in wm] + for k in globals_keys: + shard_layers.setdefault(wm[k], set()).add(-1) + + def ensure_shard(name): + if name in local_shards: + return local_shards[name] + if args.repo: + p = hf_hub_download(args.repo, name, token=token, local_dir=str(tmp), + local_dir_use_symlinks=False) + else: + p = str(src_dir / name) + local_shards[name] = p + return p + + def get_tensor(name): + shard = wm[name] + ensure_shard(shard) + if shard not in shard_handles: + # torch backend returns a native torch tensor (bf16 works natively); + # numpy backend returns a bf16 numpy array that torch.from_numpy rejects. + shard_handles[shard] = safe_open_pt(local_shards[shard], framework="pt") + return shard_handles[shard].get_tensor(name).contiguous() + + def evict_layer(i): + for shard, lays in list(shard_layers.items()): + lays.discard(i) + if not lays: + if args.repo: # only delete files we downloaded; for local keep them on disk + h = shard_handles.pop(shard, None) + try: + if h is not None: + h.__exit__(None, None, None) + except Exception: + pass + p = local_shards.pop(shard, None) + if p is not None: + try: + os.remove(p) + except OSError: + pass + # For local source we keep the handle open (the file stays on disk and + # the post-loop meta head-dim derivation re-reads it). Only the remote + # path pops/closes, since it deletes the downloaded shard. + shard_layers.pop(shard, None) + + stream_api = None + if args.stream_upload or args.upload_repo: + from huggingface_hub import HfApi + stream_api = HfApi(token=token) + if args.upload_repo: + stream_api.create_repo(args.upload_repo, repo_type="model", exist_ok=True, private=False) + + def upload_local(path: Path): + if stream_api and args.upload_repo: + stream_api.upload_file(repo_id=args.upload_repo, repo_type="model", + path_in_repo=path.name, path_or_fileobj=str(path)) + print(f" streamed -> {args.upload_repo}/{path.name}") + + # ---- resume: discover already-uploaded shards so a restart skips them ---- + done_files = set() + if args.upload_repo and stream_api is not None: + try: + done_files = set(stream_api.list_repo_files(args.upload_repo, repo_type="model")) + print(f"[resume] {len(done_files)} file(s) already on {args.upload_repo}; " + f"will skip finished layers.") + except Exception as e: + done_files = set() + print(f"[resume] could not list repo ({e}); running full conversion.") + + # ---- globals shard (embed / lm_head / final norm) as f16 ---- + if "model-globals.safetensors" in done_files: + print("[globals] already on HF, skip") + else: + g_out = {} + for k in globals_keys: + arr = get_tensor(k).half() + newk = k.replace("model.language_model.", "model.") + g_out[newk] = arr + gpath = out / "model-globals.safetensors" + save_file(g_out, str(gpath)) + print(f"[globals] {gpath.name} ({len(g_out)} tensors)") + if args.stream_upload: + upload_local(gpath) + gpath.unlink() + + # ---- per layer: Gated-Attention (full_attention) AND Gated DeltaNet (linear_attention) ---- + # Every layer also carries its MoE/MLP block. DeltaNet layers export linear_attn.* which + # the generic f16 copy below handles automatically; the engine implements the recurrence. + active = 0 + for i in all_idx: + a = active + active += 1 + out_name = f"model-{a:05d}.safetensors" + if out_name in done_files: + print(f"[layer {i} -> active {a}] {out_name} already on HF, skip") + evict_layer(i) + continue + tens = {} + lk = layer_keys(i) + # ---- MoE experts: handle BOTH layouts the source may use ---- + # (a) FUSED: model.layers.i.mlp.experts.gate_up_proj [E,2*inter,H] + # + model.layers.i.mlp.experts.down_proj [E,H,inter] + # (b) SEPARATE (standard HF Qwen3 MoE, incl. the real 35B): + # model.layers.i.mlp.experts.{e}.gate_proj [inter,H] + # model.layers.i.mlp.experts.{e}.up_proj [inter,H] + # model.layers.i.mlp.experts.{e}.down_proj [H,inter] + sep = {} # e -> {'gate':Tensor,'up':Tensor,'down':Tensor} + for k in lk: + if re.search(r"\.experts\.gate_up_proj(?:\.weight)?$", k): + gu = get_tensor(k).float() # [E, 2*inter, H] + E, twoI, H = gu.shape + inter = twoI // 2 + gate = gu[:, :inter, :] + up = gu[:, inter:, :] + dk = k.replace("gate_up_proj", "down_proj") + down = get_tensor(dk).float() # [E, H, inter] + for e in range(E): + mw, qs = make_merged(gate[e], up[e], down[e], args.ebits) + tens[f"model.layers.{a}.mlp.experts.{e}.merged_weight"] = mw + tens[f"model.layers.{a}.mlp.experts.{e}.qs"] = qs + continue + ms = re.search(r"\.experts\.(\d+)\.(gate_proj|up_proj|down_proj)(?:\.weight)?$", k) + if ms: + e = int(ms.group(1)); proj = ms.group(2) + sep.setdefault(e, {})[proj] = get_tensor(k).float() + continue + for e in sorted(sep): + d = sep[e] + if "gate_proj" in d and "up_proj" in d and "down_proj" in d: + mw, qs = make_merged(d["gate_proj"], d["up_proj"], d["down_proj"], args.ebits) + tens[f"model.layers.{a}.mlp.experts.{e}.merged_weight"] = mw + tens[f"model.layers.{a}.mlp.experts.{e}.qs"] = qs + else: + print(f"[layer {i}] WARN expert {e} missing some proj " + f"(have {sorted(d.keys())}) — skipped") + for k in lk: + newk = k.replace(f"{prefix}layers.{i}.", f"model.layers.{a}.") + if "mlp.experts." in newk: + continue # already handled above + # everything else stays f16 + tens[newk] = get_tensor(k).half() + out_path = out / f"model-{a:05d}.safetensors" + save_file(tens, str(out_path)) + print(f"[layer {i} -> active {a}] {out_path.name} ({len(tens)} tensors)") + if args.stream_upload: + upload_local(out_path) + out_path.unlink() + evict_layer(i) + + # ---- derive authoritative head dims from a real full-attention layer's weights ---- + meta = { + "model_type": cfg_full.get("model_type"), + "hidden": int(mcfg["hidden_size"]), + "n_layers": int(mcfg["num_hidden_layers"]), + "n_active": len(all_idx), + "layer_types": layer_types, + "num_experts": int(mcfg["num_experts"]), + "topk": int(mcfg["num_experts_per_tok"]), + "moe_inter": int(mcfg.get("moe_intermediate_size", mcfg.get("intermediate_size", 0) // 2)), + "shared_inter": int(mcfg.get("shared_expert_intermediate_size", mcfg.get("moe_intermediate_size", 0))), + "rms_eps": float(mcfg.get("rms_norm_eps", 1e-6)), + "ebits": args.ebits, + "scoring_func": mcfg.get("scoring_func", "softmax"), + "n_group": int(mcfg.get("n_group", 1)), + "topk_group": int(mcfg.get("topk_group", 1)), + "norm_topk_prob": bool(mcfg.get("norm_topk_prob", False)), + "attn_output_gate": bool(mcfg.get("attn_output_gate", False)), + "rope_theta": float((mcfg.get("rope_parameters") or {}).get("rope_theta", 10000000.0)), + "mrope_section": (mcfg.get("rope_parameters") or {}).get("mrope_section", [16, 16, 16]), + "partial_rotary_factor": float(mcfg.get("partial_rotary_factor", 0.25)), + } + # derive head dims from first full-attention layer weights (authoritative). + # Read from SOURCE (not local output) so this works even when layer 0 was + # skipped during a resume and its local shard was already stream-deleted. + if full_idx: + fi = full_idx[0] + def shp_src(nm): + k = f"{prefix}layers.{fi}.self_attn.{nm}.weight" + if k in wm: + return tuple(get_tensor(k).shape) + return None + qp = shp_src("q_proj") + kp = shp_src("k_proj") + vp = shp_src("v_proj") + op = shp_src("o_proj") + qn = shp_src("q_norm") + n_q = int(mcfg["num_attention_heads"]) + n_kv = int(mcfg["num_key_value_heads"]) + meta["q_heads"] = n_q + meta["kv_heads"] = n_kv + if qp is not None: + meta["q_head_dim"] = qp[0] // n_q + if kp is not None: + meta["k_head_dim"] = kp[0] // n_kv + if vp is not None: + meta["v_head_dim"] = vp[0] // n_kv + if op is not None: + meta["o_in"] = op[1] + if qn is not None: + meta["qk_rope_head_dim"] = qn[0] + meta["head_dim"] = meta.get("k_head_dim", meta.get("q_head_dim", 256)) + meta["rope_dim"] = meta.get("qk_rope_head_dim", meta["head_dim"] // 4) + # ---- DeltaNet (linear_attention) dims. From config (authoritative for dn; unlike the + # attention qk split, dn dims are explicit in config and never evicted shards). ---- + meta["dn_vheads"] = int(mcfg.get("linear_num_value_heads", mcfg.get("num_value_heads", 0))) + meta["dn_kheads"] = int(mcfg.get("linear_num_key_heads", mcfg.get("num_key_heads", 0))) + meta["dn_kdim"] = int(mcfg.get("linear_key_head_dim", mcfg.get("key_head_dim", 0))) + meta["dn_vdim"] = int(mcfg.get("linear_value_head_dim", mcfg.get("value_head_dim", 0))) + meta["dn_convk"] = int(mcfg.get("linear_conv_kernel_dim", mcfg.get("conv_kernel_size", 0))) + meta["dn_conv_dim"] = meta["dn_kheads"] * meta["dn_kdim"] * 2 + meta["dn_vheads"] * meta["dn_vdim"] + # close remaining source handles now that we've finished reading for meta + for h in shard_handles.values(): + try: + h.__exit__(None, None, None) + except Exception: + pass + (out / "qwen36_meta.json").write_text(json.dumps(meta, indent=2), encoding="utf-8") + print(f"[meta] {out / 'qwen36_meta.json'}") + + # ---- README ---- + if not args.no_readme: + lines = ["---", "tags:", " - colibri", " - qwen3.6", " - qwen3.6-35b-a3b", " - moe", + "library_name: colibri", "---", "", + "colibri container for Qwen3.6-35B-A3B (Phase 2: all layers, incl. Gated DeltaNet).", + "Every layer (Gated-Attention + Gated DeltaNet linear_attention) carries its", + "MoE/MLP block. DeltaNet weights live under model.layers.{i}.linear_attn.* and are", + "run by the recurrent gated-delta-rule in the colibri `qwen36` engine.", + "", "Engine: https://github.com/JustVugg/colibri (c/qwen36.c)"] + (out / "README.md").write_text("\n".join(lines) + "\n", encoding="utf-8") + + # ---- final upload (non-stream path) + cleanup ---- + if args.upload_repo: + print(f"\nUploading to HF Hub: {args.upload_repo}") + for f in sorted(out.glob("*.safetensors")) + [out / "config.json", out / "config.hf.json", + out / "tokenizer.json", + out / "qwen36_meta.json", out / "README.md"]: + if f.is_file(): + stream_api.upload_file(repo_id=args.upload_repo, repo_type="model", + path_in_repo=f.name, path_or_fileobj=str(f)) + print(f"Uploaded. Pull on your laptop with:\n" + f" python -c \"from huggingface_hub import snapshot_download; " + f"print(snapshot_download('{args.upload_repo}'))\"") + # never delete the generated container + try: + if tmp.is_dir(): + import shutil as _sh + _sh.rmtree(tmp, ignore_errors=True) + except Exception: + pass + + print(f"\nDone. Container at: {out}") + print(f"All {len(all_idx)} layers stored (incl. Gated DeltaNet). n_active={len(all_idx)}.") + print(f"Run (engine): SNAP={out} ./qwen36 16 {args.ebits} ref_qwen36.json") + + +if __name__ == "__main__": + main() diff --git a/c/tools/make_qwen36_oracle.py b/c/tools/make_qwen36_oracle.py new file mode 100644 index 000000000..6a82b3a35 --- /dev/null +++ b/c/tools/make_qwen36_oracle.py @@ -0,0 +1,139 @@ +#!/usr/bin/env python3 +"""Generate reference token IDs for Qwen3.6 (or a Qwen3-MoE-shaped model). + +Two modes: + --mode attention_only (default, for Phase 1) + Replace each DeltaNet layer's self_attn + mlp with a zero-returning + module, so the layer is a clean identity (residual + 0). The C engine + (qwen36.c) runs Gated Attention + MoE ONLY on the attention layers + (i%4==3) and treats the others as identity, so its output must match. + --mode full (for the eventual full hybrid engine) + Run the whole model unmodified. + +For local validation on a 24 GB laptop, pair this with tools/make_qwen36_tiny.py +(a tiny Qwen3.6-shaped model that fits in RAM) -- the 35B bf16 checkpoint needs +~70 GB and cannot be loaded here. + +Usage: + python tools/make_qwen36_oracle.py --model ./qwen36_tiny --out ref_qwen36.json + python tools/make_qwen36_oracle.py --repo Qwen/Qwen3.6-35B-A3B --out ref_qwen36.json + python tools/make_qwen36_oracle.py --model ./qwen36_tiny --out ref_qwen36.json --prompt "The capital of France is" +""" +import argparse, json, sys +from pathlib import Path + +if sys.platform == "win32": + for s in (sys.stdout, sys.stderr): + try: + s.reconfigure(encoding="utf-8") + except (AttributeError, OSError): + pass + +try: + import torch + import torch.nn as nn + from transformers import AutoTokenizer, AutoModelForCausalLM +except ImportError as exc: + sys.exit(f"Missing deps: {exc}. Run: pip install torch transformers") + + +class Zero(nn.Module): + """Returns zeros_like(input) regardless of extra args. + + Replacing a decoder layer's self_attn/mlp with this makes the layer a clean + identity: h = residual + self_attn(norm(h)) = residual + 0, and likewise for + the MLP branch. The residual/layernorm wiring is preserved, so the layer's + return tuple stays well-formed for the surrounding model. + """ + + def forward(self, hidden_states, *args, **kwargs): + if torch.is_tensor(hidden_states): + return torch.zeros_like(hidden_states) + # some layers forward a tuple; zero the first tensor element + if isinstance(hidden_states, (tuple, list)): + return tuple( + torch.zeros_like(t) if torch.is_tensor(t) else t for t in hidden_states + ) + return hidden_states + + +def detect_layer_type(model, i: int) -> str: + """A layer is 'attention' iff its index i%4==3 (Qwen3.6 layout).""" + return "attn" if i % 4 == 3 else "delta" + + +def main(): + ap = argparse.ArgumentParser(description="Qwen3.6 reference token generator") + src = ap.add_mutually_exclusive_group(required=True) + src.add_argument("--repo", help="HuggingFace repo ID") + src.add_argument("--model", help="Local HF checkpoint directory") + ap.add_argument("--out", required=True, help="Output ref JSON path") + ap.add_argument("--mode", choices=["attention_only", "full"], default="attention_only") + ap.add_argument("--prompt", default="The capital of France is") + ap.add_argument("--max-new-tokens", type=int, default=16) + args = ap.parse_args() + + if args.repo: + from huggingface_hub import snapshot_download + print(f"Resolving {args.repo} ...") + try: + mdir = snapshot_download(args.repo, local_files_only=True) + except Exception: + mdir = snapshot_download(args.repo) + else: + mdir = args.model + + print(f"Loading tokenizer from {mdir} ...") + tokenizer = AutoTokenizer.from_pretrained(mdir) + + print("Encoding prompt ...") + enc = tokenizer(args.prompt, return_tensors="pt") + prompt_ids = enc["input_ids"][0].tolist() + print(f" Prompt IDs ({len(prompt_ids)}): {prompt_ids}") + + print(f"Loading model from {mdir} ...") + print(" (large models need a lot of RAM -- be patient)") + model = AutoModelForCausalLM.from_pretrained( + mdir, torch_dtype=torch.bfloat16, device_map="cpu", low_cpu_mem_usage=True + ) + model.eval() + print(" Model loaded!") + + if args.mode == "attention_only": + n = len(model.model.layers) + replaced = 0 + for i in range(n): + if detect_layer_type(model, i) == "delta": + layer = model.model.layers[i] + layer.self_attn = Zero() + layer.mlp = Zero() + replaced += 1 + print(f" attention_only: replaced {replaced} DeltaNet layers with identity") + + print(f"Generating {args.max_new_tokens} tokens ...") + with torch.no_grad(): + out_ids = model.generate( + enc["input_ids"], + max_new_tokens=args.max_new_tokens, + do_sample=False, + use_cache=True, + ) + full_ids = out_ids[0].tolist() + gen_ids = full_ids[len(prompt_ids):] + + print(f"Prompt IDs : {prompt_ids}") + print(f"Full IDs : {full_ids}") + print(f"Generated : {gen_ids}") + try: + print(f"Text : {tokenizer.decode(gen_ids, skip_special_tokens=True)!r}") + except Exception: + pass + + payload = {"prompt_ids": prompt_ids, "full_ids": full_ids, + "mode": args.mode, "model": mdir} + Path(args.out).write_text(json.dumps(payload, indent=2)) + print(f"\nSaved reference to {args.out}") + + +if __name__ == "__main__": + main() diff --git a/c/tools/make_qwen36_tiny.py b/c/tools/make_qwen36_tiny.py new file mode 100644 index 000000000..5336df028 --- /dev/null +++ b/c/tools/make_qwen36_tiny.py @@ -0,0 +1,158 @@ +#!/usr/bin/env python3 +"""Build a tiny Qwen3.6-SHAPED model for local validation of qwen36.c (Phase 0/1). + +The real Qwen3.6-35B-A3B is ~70 GB in bf16 and cannot be loaded on a 24 GB +laptop. This script synthesizes a *tiny* model with the SAME hybrid layout +(10 x (3 x Gated DeltaNet -> MoE, 1 x Gated Attention -> MoE)) and the SAME +tensor names, but with toy dimensions: + hidden=64, n_layers=8 (-> 2 attention layers at idx 3,7), + q_heads=4, kv_heads=2, head_dim=16, rope_dim=8, + n_experts=8, topk=2, n_group=1, topk_group=1, inter=32, vocab=320. + +Because the layout is identical, convert_qwen36.py + qwen36.c treat it exactly +like the big model, so you can validate token-exactness end-to-end on a laptop. + +It also emits ref.json in attention_only mode (DeltaNet layers -> identity), +matching what qwen36.c Phase 1 computes. No tokenizer needed. + +Usage: + python tools/make_qwen36_tiny.py --out ./qwen36_tiny + python tools/make_qwen36_tiny.py --out ./qwen36_tiny --emit-ref ref_qwen36.json +""" +import argparse, json, sys +from pathlib import Path + +if sys.platform == "win32": + for s in (sys.stdout, sys.stderr): + try: + s.reconfigure(encoding="utf-8") + except (AttributeError, OSError): + pass + +try: + import torch + import torch.nn as nn +except ImportError as exc: + sys.exit(f"Missing deps: {exc}. Run: pip install torch transformers") + + +def get_classes(): + """Resolve the Qwen3-MoE model/config classes across transformers versions. + + Uses each model class's declared `config_class` (NOT a name guess): in + transformers >=5 the text model's config is `Qwen3_5MoeTextConfig`, while the + same-named `Qwen3_5MoeConfig` is the vision-language wrapper and lacks the + text fields (vocab_size, head_dim, ...). Guessing by name picks the wrong one. + """ + candidates = [ + "Qwen3_5MoeForCausalLM", + "Qwen3MoeForCausalLM", + "Qwen3NextMoeForCausalLM", + ] + import transformers + for mcls in candidates: + mc = getattr(transformers, mcls, None) + if mc is not None: + cc = getattr(mc, "config_class", None) + if cc is not None: + return mc, cc + sys.exit("No Qwen3-MoE model class found in this transformers build. Upgrade transformers.") + + +class Zero(nn.Module): + def forward(self, hidden_states, *args, **kwargs): + if torch.is_tensor(hidden_states): + return torch.zeros_like(hidden_states) + if isinstance(hidden_states, (tuple, list)): + return tuple(torch.zeros_like(t) if torch.is_tensor(t) else t for t in hidden_states) + return hidden_states + + +def build(out: Path, hidden=64, n_layers=8, q_heads=4, kv_heads=2, + head_dim=16, rope_dim=8, n_experts=8, topk=2, inter=32, + vocab=320, max_new=16, prompt_ids=None, emit_ref=None): + ModelCls, ConfigCls = get_classes() + layer_types = ["full_attention" if i % 4 == 3 else "linear_attention" + for i in range(n_layers)] + base = dict( + vocab_size=vocab, hidden_size=hidden, intermediate_size=hidden * 2, + num_hidden_layers=n_layers, num_attention_heads=q_heads, + num_key_value_heads=kv_heads, num_experts=n_experts, + num_experts_per_tok=topk, moe_intermediate_size=inter, + shared_expert_intermediate_size=inter, max_position_embeddings=512, + rms_norm_eps=1e-6, rope_theta=10000.0, tie_word_embeddings=False, + head_dim=head_dim, linear_conv_kernel_dim=4, + linear_key_head_dim=8, linear_value_head_dim=8, + linear_num_key_heads=q_heads, linear_num_value_heads=q_heads * 2, + layer_types=layer_types, hidden_act="silu", + attention_bias=False, attention_dropout=0.0, use_cache=True, + rope_parameters={"rope_type": "default", "rope_theta": 10000.0}, + ) + # Qwen3_5MoeConfig uses **kwargs, so pass everything; fall back to filtered + # only if a build rejects an unknown key. + try: + cfg = ConfigCls(**base) + except TypeError: + import inspect + allowed = set(inspect.signature(ConfigCls.__init__).parameters) - {"self"} + cfg = ConfigCls(**{k: v for k, v in base.items() if k in allowed}) + + # Some transformers builds require pad/bos/eos token ids on Qwen3-MoE configs; + # if missing, save_pretrained()/generate() crash on attribute access. + for _name, _val in (("pad_token_id", 0), ("bos_token_id", 1), ("eos_token_id", vocab - 1)): + if not hasattr(cfg, _name): + try: + setattr(cfg, _name, _val) + except Exception: + pass + + model = ModelCls(cfg) + model.eval() + out.mkdir(parents=True, exist_ok=True) + model.save_pretrained(str(out)) + print(f"Tiny model saved at {out} (params ~ {sum(p.numel() for p in model.parameters())/1e6:.1f}M)") + + if emit_ref is not None: + # attention_only: replace DeltaNet layers (i%4!=3) with identity + replaced = 0 + for i in range(n_layers): + if i % 4 != 3: + model.model.layers[i].self_attn = Zero() + model.model.layers[i].mlp = Zero() + replaced += 1 + print(f"attention_only: replaced {replaced} DeltaNet layers with identity") + if prompt_ids is None: + prompt_ids = [1, 2, 3, 4, 5] + input_ids = torch.tensor([prompt_ids]) + with torch.no_grad(): + out_ids = model.generate(input_ids, max_new_tokens=max_new, do_sample=False, + use_cache=True) + full = out_ids[0].tolist() + payload = {"prompt_ids": prompt_ids, "full_ids": full, + "mode": "attention_only", "model": "qwen36_tiny"} + Path(emit_ref).write_text(json.dumps(payload, indent=2)) + print(f"ref.json -> {emit_ref}") + print(f" prompt_ids={prompt_ids}") + print(f" full_ids ={full}") + + +def main(): + ap = argparse.ArgumentParser(description="Build a tiny Qwen3.6-shaped model") + ap.add_argument("--out", required=True, help="Output model dir") + ap.add_argument("--emit-ref", default="ref_qwen36.json", + help="Also emit this ref.json (attention_only). Set '' to skip.") + ap.add_argument("--max-new", type=int, default=16) + ap.add_argument("--prompt-ids", default=None, + help="Comma-separated token ids for the prompt (default 1,2,3,4,5)") + args = ap.parse_args() + + prompt_ids = None + if args.prompt_ids: + prompt_ids = [int(x) for x in args.prompt_ids.split(",") if x.strip() != ""] + emit = args.emit_ref if args.emit_ref else None + + build(Path(args.out), max_new=args.max_new, prompt_ids=prompt_ids, emit_ref=emit) + + +if __name__ == "__main__": + main() diff --git a/docs/qwen36-phase01.md b/docs/qwen36-phase01.md new file mode 100644 index 000000000..4926b121b --- /dev/null +++ b/docs/qwen36-phase01.md @@ -0,0 +1,97 @@ +# Qwen3.6-35B-A3B in colibri — Phase 0 + Phase 1 + +把 Qwen3.6 加进 colibri 的分阶段实现。**Phase 0** = 权重转换 + 校验预言机;**Phase 1** = Gated Attention + 流式 MoE 的 C 引擎,DeltaNet 层先当 identity 跳过。 + +## 为什么这样切 + +Qwen3.6 不是标准 MoE,是混合架构: + +``` +10 x ( 3 x (Gated DeltaNet -> MoE) , 1 x (Gated Attention -> MoE) ) +# = 40 层。Gated Attention 层 = 索引 i%4==3 (3,7,11,...,39)，共 10 层 = 25%。 +# Gated DeltaNet 层 = 其余 30 层 = 75% (Phase 2 再实现)。 +``` + +关键尺寸(以官方 `Qwen/Qwen3.6-35B-A3B` 为准): +- hidden=2048, vocab=248320, layers=40 +- **Gated Attention**: q_heads=16, kv_heads=2, head_dim=256, **rope_dim=64**(只对每头前 64 维做旋转) +- MoE: n_experts=256, topk=8 (路由) + 1 共享, inter=512 +- 激活参数 ~3B;整模型 ~35B + +Phase 1 只跑通 25% 的注意力层 + 流式 MoE 全链路(容器/分词/流式/KV/分组路由/共享专家),**用把 DeltaNet 层替换为恒等的 HF 预言机做 token 级对齐**。这把 80% 风险(自定义线性注意力)隔离到 Phase 2。 + +## 本机端到端验证(无需 35B 权重) + +35B bf16 预言机要 ~70GB,24GB 本子跑不了。用同布局的 **tiny 模型** 验证引擎逻辑: + +```bash +cd c +# 1) 造一个 tiny Qwen3.6 形模型 + 直接吐 ref.json (attention_only) +python tools/make_qwen36_tiny.py --out ../qwen36_tiny +# -> ../qwen36_tiny/ (模型) + ref_qwen36.json (attention_only 参考) + +# 2) 转成 colibri 容器 (专家 int8;Phase1 用 ebits=8 保真) +python tools/convert_qwen36.py --model ../qwen36_tiny --out ../qwen36_tiny_i8 --ebits 8 + +# 3) 编译引擎 (MinGW / gcc) +make qwen36 + +# 4) 跑,对 token +SNAP=../qwen36_tiny_i8 ./qwen36.exe 8 8 ref_qwen36.json +# 期望: Matching tokens: N/N +``` + +tiny 模型参数约几 MB,bf16 预言机仅几百 MB,24GB 本子轻松跑。对齐即证明注意力 + 流式 MoE + 分组路由 + 共享专家 + 跳过 DeltaNet 这整条链路正确。 + +## 真模型流程(需要能跑 HF 预言机的机器,如 GPU 服务器) + +```bash +# 转换 (建议 ebits=4 给 16GB 本留余量;对齐验证用 ebits=8) +python tools/convert_qwen36.py --repo Qwen/Qwen3.6-35B-A3B --out /data/qwen36_i4 --ebits 4 + +# 预言机 (在能加载 35B 的机器上跑,attention_only 模式) +python tools/make_qwen36_oracle.py --repo Qwen/Qwen3.6-35B-A3B \ + --out ref_qwen36.json --mode attention_only --max-new-tokens 32 + +# 引擎 (在 16/24GB 目标机) +SNAP=/data/qwen36_i4 ./qwen36.exe 16 4 ref_qwen36.json +``` + +`--mode full` 留给 Phase 2 之后(全混合)用。 + +## 云端转换(推荐:笔记本只下载成品容器) + +见 [`docs/qwen36-cloud-convert.md`](./qwen36-cloud-convert.md)。转换器已支持 `--repo` 拉权重 + `--upload-repo` 推回 Hub + `--low-disk` 流式省磁盘。笔记本端只 `snapshot_download` 那 18GB `.coli` 容器即可。 + +## 容器格式 / 张量约定 + +`convert_qwen36.py` 输出一个 **safetensors 分片目录**(colibri 的"容器"就是目录): +- 稠密权重(embed/attn qkv o/qk_norm/RMSNorm/router gate/shared expert/lm_head/final norm)保持原 dtype(F32 载入)。 +- 每个专家三矩阵合并为 `model.layers.{l}.mlp.experts.{e}.merged_weight`(int8,拼接顺序 **g|u|d**) + `...qs`(f32 scale,顺序 gs|us|ds)——**完全复用 olmoe.c 的 `load_expert_merged`/`Slot`,零改动**。 +- 附带 `qwen36_meta.json`:引擎所需的全部 Qwen3.6 专属维度(attn head_dim、rope_dim、MoE 分组、attn 层列表等)。缺失时引擎回退到 `i%4==3` + 默认值。 + +## 引擎要点 (`c/qwen36.c`) + +- `attention()`: Qwen3 GQA + 每头 q/k_norm(RMSNorm) + **partial RoPE**(仅每头前 `rope_dim` 维旋转,其余不变)。KV cache 按 `kv_heads` 分配。 +- `moe()`:HF Qwen3 MoE 路由——softmax(gate) → 可选 group-limited top-k(`n_group`/`topk_group`)→ 权重归一化 → 路由专家加权和 → **加上无门控共享专家**(SwiGLU)。可选 router bias(`e_score_correction_bias`)。 +- `step()`:遍历 40 层,**`is_attn[i]==0`(DeltaNet)直接跳过 → 恒等**。pilot 预取仅对"下一层也是 attn"触发。 +- 专家磁盘流式 + LRU + PILOT 预取线程 + HOT 热固定:与 olmoe 同套(`PILOT/HOT/WARMUP/WIDE/SMOOTH/CONF_LIMIT` 环境变量)。 + +## 必须人工核对的项 + +`convert_qwen36.py` 从 `config.json` + safetensors 头读出所有维度并打印。跑转换后,**核对 `qwen36_meta.json` 这些值与上方官方值一致**,尤其: +- `attn.head_dim`(应为 256)、`attn.rope_dim`(应为 64) +- `moe.n_group` / `moe.topk_group`(官方未明示,转换脚本默认 `n_group=1`;若 HF config 里有 `n_group`/`topk_group` 会被自动采用,否则无分组路由) +- `moe.has_bias`(是否含 `e_score_correction_bias`) + +若 `rope_dim` 解析为 `head_dim/4` 的回退值,需确认 config 里真正的 rope 维名字(已尝试 `rope_dim`/`rotary_emb_dim`/`qk_rope_head_dim`)。 + +## 已知限制 / 下一步 + +- DeltaNet 层(75%)在 Phase 1 是恒等 —— **只验证注意力层子集**,不是完整模型。 +- 未实现:27 层 ViT 视觉编码器、MTP 投机头(都属后续阶段)。 +- **Phase 2**:纯 C 实现 Gated DeltaNet(递归状态 + delta rule + 门控),先 tiny 对齐再接入,届时把 `step()` 里的 `continue` 换成真正的 DeltaNet 前向,并把预言机切到 `--mode full`。 + +## 校验门 + +任意修改后:重跑 tiny 流程,**`Matching tokens: N/N`(至少前若干 token 全中)**。贪心生成对量化误差敏感,`ebits=8` 应全中;`ebits=4` 可能在若干 token 后漂移(属预期,后续用 KL/perplexity 松绑)。 diff --git a/docs/qwen36-phase02.md b/docs/qwen36-phase02.md new file mode 100644 index 000000000..d5b5a4bee --- /dev/null +++ b/docs/qwen36-phase02.md @@ -0,0 +1,103 @@ +# Qwen3.6-35B-A3B in colibri — Phase 2: Gated DeltaNet in C + +Phase 1 跑通了 25% 的 Gated Attention 层 + 流式 MoE,把 DeltaNet 层当恒等跳过。 +**Phase 2 把 75% 的 Gated DeltaNet 层(线性回归注意力)也在 `c/qwen36.c` 里实现了**, +所以现在的引擎能跑**完整混合模型**:每层(无论 Attention 还是 DeltaNet)都带自己的 MoE/MLP。 + +本文件只讲 Phase 2 相对 Phase 1 的增量。Phase 1 的架构、容器格式、tiny 验证套路见 +[`qwen36-phase01.md`](./qwen36-phase01.md)。 + +## DeltaNet 数学(已在 numpy 里 torch-free 证明) + +Gated DeltaNet 子层(`HF Qwen3_5MoeGatedDeltaNet.forward`,`tools/hf_qwen3_5_moe.py` +为权威来源)逐 token: + +1. 投影:`qkv = x@qkv^T` `[S, conv_dim]`;`z = x@z^T` `[S, value_dim]`;`b=x@b^T`、`a=x@a^T` `[S, vh]`。 +2. **因果深度卷积**(groups=conv_dim, kernel=convk, silu):`conv_out[c] = silu(Σ_kk w[kk]·x_ring[kk] + w[convk-1]·qkv[c])`。`x_ring` 是**跨 chunk 携带的环形缓冲**(最近 `convk-1` 个原始输入),实现左填充因果卷积。 +3. 切分:`q_in/k_in = conv_out[:2·key_dim_tot]`(各 `key_dim_tot`),`v_in = conv_out[2·key_dim_tot:]`(`value_dim`)。 +4. `beta = sigmoid(b)`; `g = -exp(A_log)·softplus(a + dt_bias)`(都按 value head)。 +5. `q/k` 沿 head 维 `repeat_interleave` 把 `vk` 头扩成 `vh` 头;**每头 l2norm**(eps 1e-6 放进 sqrt),**q 再乘 `1/sqrt(kdim)`**,k 不乘。 +6. **递归 gated delta rule**(每个 value head `h`,状态 `S_h ∈ [kdim, vdim]`): + `S_h *= exp(g)`; `kv = k·S_h`(长度 `vdim`,**向量点乘 S 的第一轴**);`delta = (v - kv)·beta`; + `S_h += k (⊗) delta`(外积);`out = q·S_h`。 +7. **Gated RMSNorm**(普通 weight,**无 +1**,rms 里 `/vdim`):`o = (o·rsqrt(mean(o²)+eps))·dn_norm·silu(z)`。 +8. `out_proj`:`outr @ dn_out^T` → `[S, hidden]`。 + +> 关键坑(已在 Phase-2 前修掉,这里记一笔):递归里向量必须点乘 S 的**第一轴** +> (`k·S`、`q·S`),不是 `S·k`。只有当 `kdim==vdim` 时两者才巧合相等。 +> Gated RMSNorm 是**普通 weight**(HF `RMSNormGated` 无 `+1`),不是通用 RMSNorm 的 `(1+weight)`。 + +## C 实现(`c/qwen36.c`) + +- **新增 `deltanet()`**:完整实现上面 1–8,对应 `tools/_ref_dn_stream.py` 里已验证的 + `deltanet_stream`(streaming==prefill,cos≈1.0)。`matmul` 单 token 投射 + 手写递归。 +- **逐层携带状态**(放在 `Model`):`DN_rec[layer]` = 递归状态 `[vh, kdim, vdim]`; + `DN_conv[layer]` = 卷积环形缓冲 `[conv_dim, convk-1]`。两者在 `step()` 的 + prefill chunk → decode token 之间**跨调用保留**(与 HF 的 recurrent 语义一致)。 +- **`Cfg`** 新增 `dn_vheads/dn_kheads/dn_kdim/dn_vdim/dn_convk/dn_conv_dim`, + 由 `qwen36_meta.json` 的 `dn_*` 字段读入(`load_meta`)。 +- **`Layer`** 新增 `dn_qkv/dn_z/dn_b/dn_a/dn_conv/dn_dtbias/dn_alog/dn_norm/dn_out`。 +- **`step()`**:不再 `continue` 跳过 DeltaNet 层。每层 `rmsnorm(in_ln)` 后 + `is_attn[i] ? attention() : deltanet()`,残差,`rmsnorm(post_ln)`,`moe()`。 + pilot 预取改为「下一层存在即可」(现在每层都有 MoE)。 +- **KV cache 只为 attention 层分配**(DeltaNet 用递归状态,不占 KV cache,给 16GB 目标省 ~2GB)。 +- **容器**:所有层都按**原始索引** `model.layers.{i}` 存放(`active_of` 现在是恒等映射), + `linear_attn.*` 由 `convert_qwen36.py` 的通用 f16 拷贝自动导出, + `st_read_f32` 自动转成 f32。`meta` 多出 `dn_*` 维度字段。 + +## 验证策略 + +### 本地(torch-free,已完成) +- `tools/_ref_dn_stream.py`:证明 **streaming(有状态)== prefill(零填充卷积)** 整条 + DeltaNet 数学 cos≈1.0(一次性整段、逐 token、3-then-rest 三种切分全过)。这是 C `deltanet()` + 镜像的精确算法。 +- `tools/_ref_dn.py` 里的 numpy DeltaNet 与上面的 matmul 方向 / Gated RMSNorm / g 公式 + 完全一致(交叉核对过),并且它本身跑 **numpy 全前向 vs HF transformers** 的 cosine 信号。 + +### 真机(GPU 服务器 / Spark,有 gcc + torch) +```bash +cd c +# 1) HF 预言机:numpy 全前向 vs HF,给出三层 cosine 信号(权威证明 DeltaNet 正确) +python tools/_ref_dn.py --hf ../qwen36_tiny --prompt 1,2,3,4,5 +# [1] LOGIT cos [1] PER-LAYER cos [2] DELTANET-ISOLATION cos —— 都应 >0.999 + +# 2) 编译 +make qwen36 + +# 3) C 引擎 logits dump(与 numpy 同权重,做 torch-free 余弦比对) +SNAP=../qwen36_tiny_i8 DUMP=qwen36_logits.f32 ./qwen36.exe 8 8 ref_qwen36.json +# -> 写 qwen36_logits.f32 (vocab 个 float32) + +# 4) numpy 侧同样 dump(不需要 torch 也能跑前向;torch 只在 [1]/[2] 信号里用) +python tools/_ref_dn.py --hf ../qwen36_tiny --prompt 1,2,3,4,5 --dump ref_logits.f32 + +# 5) 比对余弦(应 ~0.999+;f16 容器 vs f32 源有微小误差) +python -c "import numpy as np; a=np.fromfile('qwen36_logits.f32',np.float32); b=np.fromfile('ref_logits.f32',np.float32); print('cos=', float(np.dot(a,b)/(np.linalg.norm(a)*np.linalg.norm(b))))" +``` + +### 关于 token 级对齐(重要) +tiny 模型是**随机初始化权重**(非训练权重),贪心生成的 argmax 对 f16/运算序误差极敏感, +**token 级全中在随机权重下不可能**(Phase 1 的 `ref_qwen36.json` 是 attention_only 参考, +Phase 2 引擎算的是完整模型,两者不再一致,不要再拿它比对 token)。正确判据是 **logits +余弦 ≥0.999**(上面第 5 步)。真 35B 模型用 `--mode full` 预言机时同理看余弦。 + +## 真模型流程(16/24GB 目标机) +```bash +# 转换(已支持全层;ebits=4 给 16GB 留余量,ebits=8 保真对齐) +python tools/convert_qwen36.py --repo Qwen/Qwen3.6-35B-A3B --out /data/qwen36_i4 --ebits 4 +# -> 含 qwen36_meta.json(dn_* 维度) +SNAP=/data/qwen36_i4 ./qwen36.exe 16 4 ref_qwen36.json # 全混合跑通 +``` + +## 必须人工核对的项(真模型) +转换后核对 `qwen36_meta.json`: +- `dn_vheads/dn_kheads/dn_kdim/dn_vdim/dn_convk/dn_conv_dim` 与官方 config 的 + `linear_num_value_heads / linear_num_key_heads / linear_key_head_dim / + linear_value_head_dim / linear_conv_kernel_dim` 一致(35B 默认 + vheads=??、kdim=??、convk=4 —— 以 config 实际值为准)。 +- 若 `dn_*` 缺失,引擎对 DeltaNet 层会报 "dn dims missing from meta" 并退出(正确行为)。 + +## 已知限制 / 下一步 +- 视觉编码器(27 层 ViT)、MTP 投机头仍未实现(属后续阶段,文本可跳过)。 +- 真 35B 的 `--mode full` HF 预言机应在能加载 35B 的 GPU 服务器上跑,做最终余弦校验。 +- Phase 3+ :完整混合调优 + 16GB 内存工作集压测;可选视觉/MTP。 From 9c901b54031f1aa7ae6db0631ab699dc636e69f0 Mon Sep 17 00:00:00 2001 From: minne100 Date: Sat, 1 Aug 2026 20:21:08 +0800 Subject: [PATCH 02/17] feat(qwen36): make engine drivable by coli + SERVE=1 gateway protocol Addresses the main blocker on #712: the engine built clean and touched no existing files, but coli could not reach it. c/coli: - model_arch(): return 'qwen' for any model_type containing 'qwen' - engine_for(): resolve 'qwen36' binary - need_model(): qwen36 build target - chat banner model_id: qwen36-colibri c/qwen36.c (SERVE=1 mode): - implements the colibri gateway wire protocol (READY + STAT + SUBMIT/ CANCEL + DATA + DONE + PROF), mirroring kimi_k3.c / inkling.c so 'coli chat' / 'coli web' / 'coli serve' can drive the engine - READY/STAT handshake, SUBMIT payload -> encode_text -> step() prefill + per-token step() decode, stream decoded bytes as DATA frames, stop on EOS (Qwen3 151645, Q36_EOS override), DONE/STAT on completion - temperature + top-p sampler (ported from kimi_k3.c, qsort O(V log V)) - Windows binary-mode fix inlined (#748 hang: CRT rewrites n->rn and the gateway never matches READY). Uses _setmode(_fileno(stdout/stdin), _O_BINARY) before the sentinel; switch to compat.h's coli_serve_binary_mode() once #749 lands (helper not in base yet). --- c/Makefile | 1 + c/family_registry.py | 54 ++++++++++++++ c/qwen36.c | 123 ++++++++++++++++++++++++++++++++ c/tests/test_family_registry.py | 10 ++- 4 files changed, 187 insertions(+), 1 deletion(-) diff --git a/c/Makefile b/c/Makefile index c6bfa7d65..f131359fc 100644 --- a/c/Makefile +++ b/c/Makefile @@ -1305,6 +1305,7 @@ install: colibri$(EXE) inkling$(EXE) kimi_k3$(EXE) olmoe$(EXE) \ $(INSTALL) -m 755 inkling$(EXE) $(DESTDIR)$(LIBEXECDIR)/inkling$(EXE) $(INSTALL) -m 755 kimi_k3$(EXE) $(DESTDIR)$(LIBEXECDIR)/kimi_k3$(EXE) $(INSTALL) -m 755 olmoe$(EXE) $(DESTDIR)$(LIBEXECDIR)/olmoe$(EXE) + $(INSTALL) -m 755 qwen36$(EXE) $(DESTDIR)$(LIBEXECDIR)/qwen36$(EXE) @if [ -f deepseek_v4$(EXE) ]; then \ $(INSTALL) -m 755 deepseek_v4$(EXE) $(DESTDIR)$(LIBEXECDIR)/deepseek_v4$(EXE); \ fi diff --git a/c/family_registry.py b/c/family_registry.py index ee26c4c00..8c0580839 100644 --- a/c/family_registry.py +++ b/c/family_registry.py @@ -126,6 +126,29 @@ def _glm_geometry(config, context, _model_dir): return PlannerGeometry(state, 0, workspace, experts) +def _qwen36_geometry(config, context, _model_dir): + """Hybrid: only the full_attention layers hold a KV cache; the linear + (DeltaNet) layers carry a recurrent state whose size does not depend on the + context at all. One scaled context term would over-promise on a model where + 30 of 40 layers never grow.""" + layers = _required_int(config, "num_hidden_layers", "qwen36") + kinds = config.get("layer_types") + if not isinstance(kinds, list) or len(kinds) != layers: + raise ValueError("qwen36: missing or invalid planning key 'layer_types'") + full = sum(kind == "full_attention" for kind in kinds) + kv = (full * context * _required_int(config, "num_key_value_heads", "qwen36") * + _required_int(config, "head_dim", "qwen36") * 2 * 4) + key_heads = _required_int(config, "linear_num_key_heads", "qwen36") + key_dim = _required_int(config, "linear_key_head_dim", "qwen36") + value_heads = _required_int(config, "linear_num_value_heads", "qwen36") + value_dim = _required_int(config, "linear_value_head_dim", "qwen36") + conv_k = _required_int(config, "linear_conv_kernel_dim", "qwen36", 2) + conv_dim = key_heads * key_dim * 2 + value_heads * value_dim + fixed = (layers - full) * (value_heads * key_dim * value_dim + + conv_dim * (conv_k - 1)) * 4 + return PlannerGeometry(kv, fixed, 0, _required_int(config, "num_experts", "qwen36")) + + _GLM_EXPERT = re.compile( r"(?:^|\.)model\.layers\.(\d+)\.mlp\.experts\.(\d+)\." ) @@ -263,6 +286,37 @@ def _inkling_expert_inventory(name, size, config): has_cli_adapter=True, tune_prompt_template="<|user|>\n{prompt}\n<|assistant|>\n", ), + FamilyDescriptor( + id="qwen36", + model_types=("qwen3_5_moe", "qwen3_5_moe_text"), + display_name="Qwen3.6-35B-A3B", + display_scale="35B", + engine_artifact="qwen36", + engine_aliases=(), + engine_group="qwen36", + internal_arch="qwen36", + build_target="qwen36", + process_names=("qwen36",), + default_model_id="qwen3.6-colibri", + cli_adapter="qwen36", + gateway_adapter="qwen36", + planner_id="qwen36_hybrid", + planner_geometry=_qwen36_geometry, + planner_unsupported_reason="", + expert_inventory=_individual_expert_inventory(_GLM_EXPERT), + config_section="text_config", + limits=FamilyLimits(8192, 262144, 1024, 8192, 1, 8, "Q36_MAXT"), + capabilities=FamilyCapabilities(False, False, False, True), + has_gateway_adapter=True, + # coli run stays unwired on purpose: cmd_run dispatches per arch after + # this gate, and without a qwen36 branch the engine would inherit GLM's + # prompt template -- a wrong template does not fail loudly, it degrades + # the answer. False gives the user "use coli chat or coli serve", which + # is true and actionable; chat/serve/web all work through the gateway. + has_cli_adapter=False, + tune_prompt_template=( + "<|im_start|>user\n{prompt}<|im_end|>\n<|im_start|>assistant\n\n"), + ), FamilyDescriptor( id="deepseek_v4", model_types=("deepseek_v4",), diff --git a/c/qwen36.c b/c/qwen36.c index fa9b0702f..99da321c6 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -1839,6 +1839,123 @@ static int *read_int_array(jval *o, const char *key, int *n_out) { } #ifndef QWEN36_NO_MAIN + +/* ===================== coli serve mode (SERVE=1) ===================== * + * Implements the colibri gateway wire protocol so `coli chat` / `coli web` / + * `coli serve` can drive this engine. Without it the engine is unreachable: + * users run `coli chat`, not the binary directly. + * Protocol (matches kimi_k3.c / inkling.c, the other non-GLM engines): + * engine: \x01\x01READY\x01\x01\n + * STAT 0 0.00 0.0 \n + * gateway: SUBMIT \n \n + * engine: ACCEPT \n + * DATA \n \n (repeated per decoded chunk) + * DONE STAT \n + * gateway: CANCEL (abort current turn) + * Windows: stdout/stdin must go binary BEFORE the READY sentinel or the CRT + * rewrites the trailing \n as \r\n and the gateway never matches it -> the + * session hangs forever (#748). compat.h's coli_serve_binary_mode (#749) + * carries that fix for every engine; see its comment for the full story. */ + +typedef struct { char id[64]; int max_tok; float temp, top_p; char *payload; int plen; } ServeReq; + +static int serve_read_req(ServeReq *q){ + char line[512], cmd[16], id[64]; + if(!fgets(line,sizeof(line),stdin)) return -1; + if(sscanf(line,"%15s %63s",cmd,id)<2) return 0; + if(!strcmp(cmd,"CANCEL")||!strcmp(cmd,"STOP")) return 0; + if(strcmp(cmd,"SUBMIT")) return 0; + int slot, plen, max_tok; float temp, top_p; + if(sscanf(line,"%*s %*s %d %d %d %f %f",&slot,&plen,&max_tok,&temp,&top_p)!=5 || + plen<0||plen>(1<<24)||max_tok<1){ + printf("ERROR %s bad submit header\n",id); fflush(stdout); return 0; + } + (void)slot; + char *payload=malloc((size_t)plen+1); + if(!payload){ printf("ERROR %s out of memory\n",id); fflush(stdout); return 0; } + if(fread(payload,1,(size_t)plen,stdin)!=(size_t)plen){ free(payload); return -1; } + (void)fgetc(stdin); payload[plen]=0; + snprintf(q->id,sizeof(q->id),"%s",id); + q->max_tok=max_tok; q->temp=temp; q->top_p=top_p; + q->payload=payload; q->plen=plen; + return 2; +} + +static void serve_data(const char *id, const char *p, int n){ + if(n<=0) return; + printf("DATA %s %d\n",id,n); + fwrite(p,1,(size_t)n,stdout); fputc('\n',stdout); fflush(stdout); +} + +/* temperature + top-p sampler (ported from kimi_k3.c; vocab ~250k -> qsort O(V log V) per token) */ +typedef struct { float p; int id; } SampleProb; +static int sample_prob_desc(const void *a, const void *b){ + float pa=((const SampleProb*)a)->p, pb=((const SampleProb*)b)->p; + return (pb>pa)-(pa>pb); +} +static int serve_sample(const float *lo, int V, float temp, float top_p){ + if(temp<=0.f){ int b=0; for(int i=1;ilo[b]) b=i; return b; } + SampleProb *rank=malloc((size_t)V*sizeof(SampleProb)); float mx=lo[0]; + if(!rank){ fprintf(stderr,"OOM sampling\n"); exit(1); } + for(int i=1;imx) mx=lo[i]; + double sum=0; + for(int i=0;i0.f&&top_p<1.f)?top_p*sum:sum, kept=0; int n=0; + while(n=r){ pick=rank[i].id; break; } } + free(rank); return pick; +} + +static void serve_one(Model *m, ServeReq *q){ + int *ids=NULL, np=0; + encode_text(q->payload, &ids, &np); /* payload is raw prompt text; qwen36 adds no BOS */ + int max_ctx = getenv("Q36_MAXT")?atoi(getenv("Q36_MAXT")):8192; + if(np<1 || np+q->max_tok>max_ctx){ + printf("ERROR %s CONTEXT_EXCEEDED prompt_tokens=%d requested=%d capacity=%d\n",q->id,np,q->max_tok,max_ctx); + fflush(stdout); free(ids); return; + } + printf("ACCEPT %s %d\n",q->id,np); fflush(stdout); + m->max_t = np + q->max_tok; + reset_recurrent(m); ensure_kv(m); m->kv_len = 0; + float *lo = step(m, ids, np, 0); + int gen=0, limited=1; + int eos_id = getenv("Q36_EOS")?atoi(getenv("Q36_EOS")):151645; /* Qwen3 <|im_end|> */ + double t0=now_s(); + unsigned char sbuf[16]; int sbn=0; + for(int s=0;smax_tok;s++){ + int tk = serve_sample(lo, m->c.vocab, q->temp, q->top_p); + free(lo); lo=NULL; + if(tk==eos_id){ limited=0; break; } + unsigned char tmp[256]; int tn=0; decode_id_to_bytes(tk, tmp, &tn); + unsigned char chunk[256]; int cn=0; utf8_drain(sbuf,&sbn,tmp,tn,chunk,&cn); + if(cn>0) serve_data(q->id,(char*)chunk,cn); + gen++; + lo = step(m, &tk, 1, np+s); + } + if(sbn>0) serve_data(q->id,(char*)sbuf,sbn); /* flush trailing partial UTF-8 */ + free(lo); free(ids); + double dt=now_s()-t0; + printf("DONE %s STAT %d %.3f %.1f %.2f %d %d\n",q->id,gen, + dt>0?gen/dt:0.0,0.0,rss_gb(),np,limited); + fflush(stdout); +} + +static void serve_loop(Model *m){ + coli_serve_binary_mode(); + setvbuf(stdin,NULL,_IONBF,0); + fputs("\x01\x01READY\x01\x01\n",stdout); + printf("STAT 0 0.00 0.0 %.2f\n",rss_gb()); + fflush(stdout); + for(;;){ + ServeReq q={0}; int r; + do r=serve_read_req(&q); while(r==0); + if(r<0) return; + if(r==2){ serve_one(m,&q); free(q.payload); } + } +} + int main(int argc, char **argv) { const char *snap = getenv("SNAP"); if (!snap) { fprintf(stderr, "set SNAP=\n"); return 1; } @@ -1931,6 +2048,12 @@ 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 */ + 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); + return 0; + } if (is_ref && getenv("PPL") && atoi(getenv("PPL")) == 1) { double nll; double t = now_s(); diff --git a/c/tests/test_family_registry.py b/c/tests/test_family_registry.py index 190ae257d..f9aa461d4 100644 --- a/c/tests/test_family_registry.py +++ b/c/tests/test_family_registry.py @@ -136,7 +136,10 @@ def test_qwen_fixture_models_gqa_and_fixed_deltanet_state(self): "linear_key_head_dim": 8, "linear_value_head_dim": 8, "linear_conv_kernel_dim": 4, } - by_id, by_type = _build_registry(FAMILIES + (QWEN36_FIXTURE,)) + # qwen36 is a registered family now, so the fixture would collide on its + # model_type alias. Assert against the production descriptor instead -- + # which is the stronger test: it pins the shipped geometry, not a copy. + by_id, by_type = _build_registry(FAMILIES) family = by_type[config["model_type"]] self.assertEqual(family, by_id["qwen36"]) resolved = type("R", (), {"descriptor": family, "family_config": config, @@ -224,6 +227,11 @@ def test_tuning_replay_prompts_are_registry_owned(self): "inkling": "<|user|>hello {world}<|assistant|>", "kimi": "K3CHAT1\nM user 13\nhello {world}G 0\n\n", "olmoe": "<|user|>\nhello {world}\n<|assistant|>\n", + # Qwen3.6's generation prompt MUST open : the model was + # never trained on the bare "assistant\\n" state and greedy argmax + # there lands on an EOS special (measured gen=0). + "qwen36": "<|im_start|>user\nhello {world}<|im_end|>\n" + "<|im_start|>assistant\n\n", "deepseek_v4": "hello {world}", } self.assertEqual( From 5a0715e744da1e36f713ddf684d7bdd71e8210eb Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Sat, 1 Aug 2026 17:12:42 +0200 Subject: [PATCH 03/17] =?UTF-8?q?feat(qwen36):=20group-scaled=20int4=20(gs?= =?UTF-8?q?64)=20containers=20=E2=80=94=20converter=20flag=20+=20CPU=20rea?= =?UTF-8?q?d=20path?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit convert_qwen36.py --gs N writes one f32 scale per N input elements per row (scales [O, ceil(I/gs)], expert_gs in qwen36_meta.json) instead of one per row. The engine reads expert_gs and dispatches expert GEMVs through matmul_qe: per-row scales unchanged, grouped scales via matmul_q_gs. Same trade as GLM's gs64 containers: ~9% more container bytes for scales that track short rows (moe_intermediate_size=512) much more closely. Validated on the tiny fixture: grouped beats per-row on cosine to the int8 anchor. Full-model A/B on the published containers is in the PR thread. --- c/qwen36.c | 74 +++++++++++++++++++++++++++++++++++---- c/tools/convert_qwen36.py | 31 ++++++++++++++-- 2 files changed, 95 insertions(+), 10 deletions(-) diff --git a/c/qwen36.c b/c/qwen36.c index 99da321c6..567a20429 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -543,6 +543,7 @@ typedef struct { uint8_t *is_attn; /* [n_layers] 1 if Gated Attention layer, 0 if DeltaNet */ /* Gated DeltaNet (linear_attention) dims, read from qwen36_meta.json. */ int dn_vheads, dn_kheads, dn_kdim, dn_vdim, dn_convk, dn_conv_dim; + int expert_gs; /* expert scale group size along input dim; 0 = per-row */ } Cfg; /* ---------- per-layer dense weights ---------- */ @@ -748,6 +749,58 @@ static void matmul_q(float *y, const float *x, const int8_t *q, const float *sca #endif } +/* Group-scaled int8 GEMV: one f32 scale per `gs` input elements per row + * (gs64 expert containers). Row layout of `scale`: [O][I/gs] row-major. */ +static int g_expert_gs = 0; /* set from qwen36_meta.json (expert_gs) at load */ +static void matmul_q_gs(float *y, const float *x, const int8_t *q, const float *scale, + int I, int O, int gs) { + int ng = (I + gs - 1) / gs; +#if defined(__AVX2__) && defined(__FMA__) + if ((gs & 31) == 0) { + #pragma omp parallel for schedule(static) if(O >= 256) + for (int o = 0; o < O; o++) { + const int8_t *w = q + (int64_t)o * I; + const float *sc = scale + (int64_t)o * ng; + float acc = 0.f; + for (int gi = 0; gi < ng; gi++) { + __m256 a0 = _mm256_setzero_ps(), a1 = _mm256_setzero_ps(); + int base = gi * gs, end = base + gs; if (end > I) end = I; + for (int i = base; i + 16 <= end; i += 16) { + __m128i b0 = _mm_loadu_si128((const __m128i*)(w + i)); + a0 = _mm256_fmadd_ps(_mm256_loadu_ps(x+i), _mm256_cvtepi32_ps(_mm256_cvtepi8_epi32(b0)), a0); + a1 = _mm256_fmadd_ps(_mm256_loadu_ps(x+i+8), _mm256_cvtepi32_ps(_mm256_cvtepi8_epi32(_mm_srli_si128(b0,8))), a1); + } + a0 = _mm256_add_ps(a0, a1); + __m128 s = _mm_add_ps(_mm256_castps256_ps128(a0), _mm256_extractf128_ps(a0,1)); + s = _mm_add_ps(s, _mm_movehl_ps(s,s)); + s = _mm_add_ss(s, _mm_shuffle_ps(s,s,1)); + acc += _mm_cvtss_f32(s) * sc[gi]; + } + y[o] = acc; + } + return; + } +#endif + #pragma omp parallel for schedule(static) if(O >= 256) + for (int o = 0; o < O; o++) { + const int8_t *w = q + (int64_t)o * I; + const float *sc = scale + (int64_t)o * ng; + float acc = 0.f; + for (int gi = 0; gi < ng; gi++) { + int base = gi * gs, end = base + gs; if (end > I) end = I; + float part = 0.f; + for (int i = base; i < end; i++) part += x[i] * (float)w[i]; + acc += part * sc[gi]; + } + y[o] = acc; + } +} +/* Expert-GEMV dispatch: per-row scales (classic) or grouped (gs64 container). */ +static void matmul_qe(float *y, const float *x, const int8_t *q, const float *scale, int I, int O) { + if (g_expert_gs) matmul_q_gs(y, x, q, scale, I, O, g_expert_gs); + else matmul_q(y, x, q, scale, I, O); +} + /* ---- Dense int8: per-row quantized copies of the large f32 matrices. * matmul_d dispatches via pointer lookup to matmul_q; COLI_DENSE_I8=0 falls * back to f32 (reference path for parity tests). ~4x less memory traffic. */ @@ -852,6 +905,7 @@ static void load_meta(Cfg *c, const char *snap) { G("q_heads", q_heads); G("kv_heads", kv_heads); G("head_dim", head_dim); G("q_head_dim", q_head_dim); G("k_head_dim", k_head_dim); G("v_head_dim", v_head_dim); G("o_in", o_in); G("rope_dim", rope_dim); G("qk_rope_head_dim", rope_dim); + G("expert_gs", expert_gs); G("num_experts", n_experts); G("topk", topk); G("moe_inter", inter); G("shared_inter", shared_inter); G("n_group", n_group); G("topk_group", topk_group); @@ -1007,6 +1061,10 @@ static void model_init(Model *m, const char *snap, int cap, int bits) { m->dense_load_s = now_s() - t0; } +/* scale counts per expert matrix: per-row (gs=0) or grouped along input dim */ +static int64_t scale_count_gu(const Cfg *c){ return c->expert_gs ? (int64_t)c->inter * ((c->hidden + c->expert_gs - 1) / c->expert_gs) : c->inter; } +static int64_t scale_count_d (const Cfg *c){ return c->expert_gs ? (int64_t)c->hidden * ((c->inter + c->expert_gs - 1) / c->expert_gs) : c->hidden; } + static void slot_ensure_allocated(Model *m, Slot *s) { if (s->g) return; Cfg *c = &m->c; @@ -1017,10 +1075,10 @@ static void slot_ensure_allocated(Model *m, Slot *s) { s->g = w_block; s->u = w_block + ng; s->d = w_block + ng + ng; - float *s_block = falloc(c->inter + c->inter + c->hidden); + float *s_block = falloc(2*scale_count_gu(c) + scale_count_d(c)); s->gs = s_block; - s->us = s_block + c->inter; - s->ds = s_block + c->inter + c->inter; + s->us = s_block + scale_count_gu(c); + 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) */ @@ -1034,7 +1092,7 @@ static void load_expert_merged(Model *m, int layer, int eid, Slot *s) { Cfg *cc = &m->c; int64_t ng = (int64_t)cc->inter * cc->hidden, nd = (int64_t)cc->hidden * cc->inter; int64_t want_w = ng + ng + nd; - int64_t want_s = (int64_t)cc->inter + cc->inter + cc->hidden; + int64_t want_s = 2*scale_count_gu(cc) + scale_count_d(cc); st_tensor *tw = st_find(&m->S, nm), *ts = st_find(&m->S, qsnm); if (!tw || (tw->nbytes != want_w && tw->nbytes != want_w / 2)) { fprintf(stderr, "%s: expert weight is %lld bytes — expected %lld (int8) or %lld (int4)\n", @@ -1389,10 +1447,10 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out) { { for (int kk = 0; kk < K; kk++) { Slot *e; expert_get(m, layer, idx[kk], &e); - matmul_q(g, xs, e->g, e->gs, D, I); - matmul_q(u, xs, e->u, e->us, D, I); + 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]; } - matmul_q(hh, g, e->d, e->ds, I, D); + 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]; @@ -2016,6 +2074,8 @@ int main(int argc, char **argv) { } Model m; model_init(&m, snap, cap, bits); + g_expert_gs = m.c.expert_gs; + if (g_expert_gs) fprintf(stderr, "[qwen36] group-scaled experts: gs=%d\n", g_expert_gs); fprintf(stderr, "resident weights loaded in %.1fs | RSS after load: %.2f GB\n", m.dense_load_s, rss_gb()); /* quantize the large dense matrices to int8 (COLI_DENSE_I8=0 disables) */ if (dense_i8_on()) { diff --git a/c/tools/convert_qwen36.py b/c/tools/convert_qwen36.py index c5126c003..a694080d3 100644 --- a/c/tools/convert_qwen36.py +++ b/c/tools/convert_qwen36.py @@ -84,7 +84,25 @@ def pack_int4(q: "torch.Tensor") -> "torch.Tensor": return (lo | hi).contiguous() -def make_merged(gate, up, down, ebits): +def quantize_row_grouped(w: "torch.Tensor", bits: int, gs: int): + """Group-wise symmetric quantization: one f32 scale per `gs` input elements + per row (like GLM's gs64 containers). Returns (q int8 [O,I], scales f32 + [O, ceil(I/gs)] flattened row-major).""" + qmax = (1 << (bits - 1)) - 1 + O, I = w.shape[0], w.reshape(w.shape[0], -1).shape[1] + w_f32 = w.reshape(O, -1).float() + pad = (-I) % gs + if pad: + w_f32 = torch.nn.functional.pad(w_f32, (0, pad)) + ng = w_f32.shape[1] // gs + g = w_f32.view(O, ng, gs) + scales = g.abs().amax(dim=2, keepdim=True).clamp(min=1e-12) / qmax + q = (g / scales).round().clamp(-qmax - 1, qmax).to(torch.int8).view(O, -1)[:, :I] + return q, scales.view(O, ng) + + +def make_merged(gate, up, down, ebits, gs=0): + gsz = gs # local `gs` is rebound to the gate scales below -- keep the group size safe """gate/up: [inter, H]; down: [H, inter] (torch, any fp). -> (merged_weight, qs f32 1D). @@ -96,6 +114,9 @@ def make_merged(gate, up, down, ebits): qs (per-row f32 scales) is identical in both cases. """ def q(t): + if gsz: + qt, s = quantize_row_grouped(t, ebits, gsz) # [O, ng] scales + return qt, s.reshape(-1) qt, s = quantize_row(t.reshape(t.shape[0], -1), ebits) # rows along dim0 return qt, s gq, gs = q(gate) @@ -189,6 +210,9 @@ def main(): src.add_argument("--model", help="Local HF checkpoint directory") ap.add_argument("--out", required=False, help="Output container directory") ap.add_argument("--ebits", type=int, default=4, help="Expert quant bits (2..8, default 4)") + ap.add_argument("--gs", type=int, default=0, + help="Group size for expert scales (e.g. 64). 0 = per-row (default). " + "Group-scaled containers need engine support (expert_gs in meta).") ap.add_argument("--upload-repo", help="Push the finished container to this HF repo ID") ap.add_argument("--hf-token", help="HF token (defaults to HF_TOKEN env / cached login)") ap.add_argument("--low-disk", action="store_true", @@ -430,7 +454,7 @@ def upload_local(path: Path): dk = k.replace("gate_up_proj", "down_proj") down = get_tensor(dk).float() # [E, H, inter] for e in range(E): - mw, qs = make_merged(gate[e], up[e], down[e], args.ebits) + mw, qs = make_merged(gate[e], up[e], down[e], args.ebits, gs=args.gs) tens[f"model.layers.{a}.mlp.experts.{e}.merged_weight"] = mw tens[f"model.layers.{a}.mlp.experts.{e}.qs"] = qs continue @@ -442,7 +466,7 @@ def upload_local(path: Path): for e in sorted(sep): d = sep[e] if "gate_proj" in d and "up_proj" in d and "down_proj" in d: - mw, qs = make_merged(d["gate_proj"], d["up_proj"], d["down_proj"], args.ebits) + mw, qs = make_merged(d["gate_proj"], d["up_proj"], d["down_proj"], args.ebits, gs=args.gs) tens[f"model.layers.{a}.mlp.experts.{e}.merged_weight"] = mw tens[f"model.layers.{a}.mlp.experts.{e}.qs"] = qs else: @@ -513,6 +537,7 @@ def shp_src(nm): meta["o_in"] = op[1] if qn is not None: meta["qk_rope_head_dim"] = qn[0] + meta["expert_gs"] = args.gs # 0 = per-row scales; >0 = group size along input dim meta["head_dim"] = meta.get("k_head_dim", meta.get("q_head_dim", 256)) meta["rope_dim"] = meta.get("qk_rope_head_dim", meta["head_dim"] // 4) # ---- DeltaNet (linear_attention) dims. From config (authoritative for dn; unlike the From 24d3c89e55204a7f5650bd85eff7401a86cc6d47 Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Mon, 3 Aug 2026 21:23:33 +0200 Subject: [PATCH 04/17] =?UTF-8?q?docs(qwen36):=20user-facing=20page=20?= =?UTF-8?q?=E2=80=94=20prebuilt=20containers=20as=20the=20easy=20path?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Mirrors docs/inkling.md's structure: pre-converted containers first (gs64 recommended, per-row available), converter as the DIY route, coli chat as the primary interface now that the engine speaks the serve protocol. The container recommendation follows the GLM #455 experience; the per-row-vs-gs64 measurement backing it is in the PR thread. --- docs/qwen36.md | 62 ++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 62 insertions(+) create mode 100644 docs/qwen36.md diff --git a/docs/qwen36.md b/docs/qwen36.md new file mode 100644 index 000000000..134f59d26 --- /dev/null +++ b/docs/qwen36.md @@ -0,0 +1,62 @@ +# Qwen3.6-35B-A3B on colibri + +`c/qwen36.c` runs [Qwen/Qwen3.6-35B-A3B](https://huggingface.co/Qwen/Qwen3.6-35B-A3B) +(35B total / ~3B active, Apache 2.0) — a hybrid architecture: 25% Gated +Attention layers, 75% Gated DeltaNet (linear attention) layers, each followed +by a streamed MoE block (256 experts, top-8 + 1 shared). Dense weights stay +resident; routed experts stream from the container through an LRU + pinned +cache. Development notes live in `docs/qwen36-phase01.md` / +`qwen36-phase02.md`. + +Architecture-identical checkpoints (same config geometry, e.g. +KAT-Coder-V2.5-Dev) run on this engine unchanged. + +## Quickstart + +Pre-converted containers (int4 experts, self-contained, ~20 GB): + +```sh +# group-scaled int4 (gs64) — recommended, see "Which container" below +hf download Kreuzzelg/qwen36-35b-a3b-colibri-i4-gs64 --local-dir ~/Models/qwen36_i4_gs64 + +# per-row int4 +hf download Kreuzzelg/qwen36-35b-a3b-colibri-i4 --local-dir ~/Models/qwen36_i4 +``` + +or convert the original bf16 checkpoint yourself (~70 GB download): + +```sh +python3 c/tools/convert_qwen36.py --repo Qwen/Qwen3.6-35B-A3B --out ~/Models/qwen36_i4_gs64 --gs 64 +``` + +Build and chat: + +```sh +make -C c qwen36 +COLI_MODEL=~/Models/qwen36_i4_gs64 ./c/coli chat +``` + +`coli` reads the model's `config.json` (`model_type` contains `qwen`), picks +this engine, and drives it over the serve protocol — `coli web` and +`coli serve` (OpenAI-compatible API) work the same way. + +Direct invocation without the gateway: + +```sh +SNAP=~/Models/qwen36_i4_gs64 TOK=~/Models/qwen36_i4_gs64/tokenizer.json \ +N_NEW=200 ./c/qwen36 256 4 prompt.txt +``` + +Requirements: ~30 GB RAM for comfortable expert caching, NVMe storage for the +container. CPU-only in this build; the CUDA VRAM expert tier is a separate PR +(`docs/qwen36-cuda-tier.md`). + +## Which container? + +The gs64 container carries one scale per 64-weight group instead of one per +row. On GLM, per-row int4 was the root cause of think-mode loops and +never-terminating generations (#455), and group scales fixed them in +controlled A/Bs — with `moe_intermediate_size=512`, Qwen's rows are short, so +per-row quantization error concentrates the same way. The gs64 container costs +~1.7 GB more on disk and a few percent on cold-start; warm decode speed is the +same or slightly better. From 2ff73e0494194c19f08032d1465dbc0d0f8fcb88 Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Mon, 3 Aug 2026 21:47:58 +0200 Subject: [PATCH 05/17] fix(qwen36): serve mode must not require an argv prompt file MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit main loaded the prompt (argv[3], default ref.json) before the SERVE branch, so an engine launched the way coli launches it — SERVE=1 and no prompt argument — exited 1 on the missing ref.json before ever emitting READY. Serve gets its prompts over the wire; skip the argv prompt load in serve mode. Found by running the serve handshake end-to-end. --- c/qwen36.c | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/c/qwen36.c b/c/qwen36.c index 567a20429..ec4ed12af 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -2042,6 +2042,11 @@ int main(int argc, char **argv) { int *prompt=NULL, *full=NULL, *out=NULL; int np=0, nfull=0, n_new=0; char *buf=NULL, *arena=NULL; + /* serve mode gets its prompts over the wire: skip the argv prompt file + * entirely, or the default "ref.json" kills the engine before serve_loop + * is ever reached — which is exactly how `coli` launches it (SERVE=1, no + * prompt argument). */ + int serve_mode = getenv("SERVE") && getenv("SERVE")[0]=='1'; /* load tokenizer early so text-prompt mode can encode before model_init */ { @@ -2051,7 +2056,9 @@ int main(int argc, char **argv) { else { char tpb[2048]; snprintf(tpb,sizeof tpb,"%s/tokenizer.json",snap); load_tokenizer(tpb); } } - if (is_ref) { + if (serve_mode) { + /* no argv prompt to load */ + } else if (is_ref) { FILE *f = fopen(refpath, "rb"); if (!f) { perror(refpath); return 1; } fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); buf=malloc(n+1); if (fread(buf,1,n,f)!=(size_t)n) {} buf[n]=0; fclose(f); From 342f1042c1aef28cb65e4b0eefb792f8d7e7854e Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Mon, 3 Aug 2026 22:24:38 +0200 Subject: [PATCH 06/17] fix(qwen36): serve EOS ids from the tokenizer, not the 151k-vocab constant MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit serve_one compared against a hardcoded 151645 — <|im_end|> in the older 151k-token Qwen vocabs. Qwen3.6's vocab is 248320 tokens with the specials at 248044+ (<|endoftext|> 248044, <|im_end|> 248046), so EOS never matched and every serve turn ran to max_tok; in the worst case 151645 is an ordinary BPE piece here and could truncate a completion mid-sentence. Resolve <|im_end|> and <|endoftext|> from the tokenizer's added_tokens (same spirit as kimi_k3's config-driven eos list); Q36_EOS still overrides. Found by the serve-path termination check the A/B comment called for. --- c/qwen36.c | 20 ++++++++++++++++++-- 1 file changed, 18 insertions(+), 2 deletions(-) diff --git a/c/qwen36.c b/c/qwen36.c index ec4ed12af..9fd721e57 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -1966,6 +1966,21 @@ static int serve_sample(const float *lo, int V, float temp, float top_p){ free(rank); return pick; } +/* Chat turns end on <|im_end|>, base completions on <|endoftext|>. Resolve + * both ids from the tokenizer's added_tokens: Qwen3.6's 248320-token vocab + * puts them at 248044+, so the old hardcoded 151645 (the 151k-vocab Qwen id) + * silently never matched and every serve turn ran into max_tok. Q36_EOS + * still overrides for experiments. */ +static int serve_eos_ids(int *ids, int cap){ + int n=0; + if(getenv("Q36_EOS")){ ids[n++]=atoi(getenv("Q36_EOS")); return n; } + for(int k=0;k")||!strcmp(g_sp_str[k],"<|endoftext|>")) + ids[n++]=g_sp_id[k]; + if(!n) ids[n++]=151645; /* tokenizer without added_tokens: old default */ + return n; +} + static void serve_one(Model *m, ServeReq *q){ int *ids=NULL, np=0; encode_text(q->payload, &ids, &np); /* payload is raw prompt text; qwen36 adds no BOS */ @@ -1979,13 +1994,14 @@ static void serve_one(Model *m, ServeReq *q){ reset_recurrent(m); ensure_kv(m); m->kv_len = 0; float *lo = step(m, ids, np, 0); int gen=0, limited=1; - int eos_id = getenv("Q36_EOS")?atoi(getenv("Q36_EOS")):151645; /* Qwen3 <|im_end|> */ + int eos_ids[4]; int n_eos=serve_eos_ids(eos_ids,4); double t0=now_s(); unsigned char sbuf[16]; int sbn=0; for(int s=0;smax_tok;s++){ int tk = serve_sample(lo, m->c.vocab, q->temp, q->top_p); free(lo); lo=NULL; - if(tk==eos_id){ limited=0; break; } + int is_eos=0; for(int e=0;e0) serve_data(q->id,(char*)chunk,cn); From 6a9f2930c45863957a9ebe36acdef6c9a0928c33 Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Mon, 3 Aug 2026 22:31:44 +0200 Subject: [PATCH 07/17] feat(qwen36): qwen chat-template family in the gateway MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit openai_server.py knew glm/inkling/kimi only, so a qwen model fell to the GLM renderer and the engine never saw its native frames. render_chat_qwen mirrors Qwen3.6's chat_template byte for byte in the text-only subset (inkling's approach): <|im_start|>role\n...<|im_end|>\n frames and a generation prompt that OPENS the mandatory block — the model was never trained on a bare 'assistant\n' state, and greedy argmax there is an EOS special (measured: gen=0 on the very first token). Thinking disabled pre-closes the block, exactly like the official jinja. Arch detection, --arch choice, renderer dispatch and the single-KV-slot guard follow the kimi pattern. Spawn needs nothing new: argv is (cap), bits defaults to 4, tokenizer comes from SNAP. --- c/openai_server.py | 35 ++++++++++++++++++++++++++++++++++- 1 file changed, 34 insertions(+), 1 deletion(-) diff --git a/c/openai_server.py b/c/openai_server.py index fa258df89..19108c4ba 100644 --- a/c/openai_server.py +++ b/c/openai_server.py @@ -484,7 +484,8 @@ def _tool_hold(): return max(len(m) for m in _tool_stream_markers()) - 1 -ARCH = "glm" # set in main(): glm | inkling | kimi | deepseek_v4 +ARCH = "glm" # set in main(): a family id from family_registry (glm | inkling | + # kimi | olmoe | qwen36 | deepseek_v4) INK_THINK, INK_TEXT = "<|content_thinking|>", "<|content_text|>" @@ -917,6 +918,37 @@ def render_chat_olmoe(messages, enable_thinking=False, reasoning_effort=None, to return "".join(parts) +def render_chat_qwen(messages, enable_thinking=False, reasoning_effort=None, tools=None, + tool_choice=None): + """Text-only subset of Qwen3.6's chat_template: <|im_start|>role\\n ... + <|im_end|>\\n frames, then the generation prompt. The official template + opens a mandatory block after `<|im_start|>assistant\\n` — the + model was never trained on the bare `assistant\\n` state, and greedy + argmax there lands on an EOS special (measured: gen=0). With thinking + disabled the template pre-closes the block instead; both branches are + mirrored here byte for byte.""" + if not isinstance(messages, list) or not messages: + raise APIError(400, "`messages` must be a non-empty array.", "messages") + if tools or tool_choice not in (None, "none"): + raise APIError(400, "Tool use is not wired up for the qwen36 engine yet.", + "tools", "unsupported_parameter") + parts = [] + for index, message in enumerate(messages): + if not isinstance(message, dict): + raise APIError(400, "Each message must be an object.", f"messages.{index}") + role = message.get("role") + if role == "developer": + role = "system" + if role not in ("system", "user", "assistant"): + raise APIError(400, f"Unsupported role {role!r}.", f"messages.{index}.role") + raw = message.get("content") + text = content_text(raw, f"messages.{index}.content") if raw is not None else "" + parts.append(f"<|im_start|>{role}\n{text}<|im_end|>\n") + parts.append("<|im_start|>assistant\n") + parts.append("\n" if enable_thinking else "\n\n\n\n") + return "".join(parts) + + def render_chat_inkling(messages, enable_thinking=False, reasoning_effort=None, tools=None, tool_choice=None, audio_out=None): """Text-only subset of Inkling's chat_template.jinja: role tokens with @@ -1092,6 +1124,7 @@ def render_chat_for_arch(messages, enable_thinking=False, reasoning_effort=None, return render_chat_inkling(messages, enable_thinking, reasoning_effort, tools, tool_choice, audio_out=audio_out) renderer = (render_chat_kimi if ARCH == "kimi" else + render_chat_qwen if ARCH == "qwen36" else render_chat_v4 if ARCH == "deepseek_v4" else render_chat_olmoe if ARCH == "olmoe" else render_chat) return renderer(messages, enable_thinking, reasoning_effort, tools, tool_choice) From c30c5457bbe89c1a7349f8794c8f4ab0c04750a1 Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Thu, 13 Aug 2026 22:55:23 +0200 Subject: [PATCH 08/17] =?UTF-8?q?fix(qwen36):=20heap=20per-thread=20attent?= =?UTF-8?q?ion=20score=20rows=20=E2=80=94=208k=20context=20smashed=20the?= =?UTF-8?q?=20stack?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit attention() kept one score per attended key in a fixed `float sc[8192]` on the stack and filled it with `for (t = 0; t <= qpos; t++)`, so any position past 8191 wrote off the end. serve_one() carried the same literal as the Q36_MAXT default and refused longer requests, but generate() and tf_nll() set max_t straight from the caller and never consulted it — and raising Q36_MAXT moved the guard without moving the buffer, turning a refusal into a stack smash. Same bug and the same fix as #110 in the MLA path (found again independently in #122): the score rows move to a per-thread heap buffer sized from max_t, allocated in ensure_kv() where the other context-sized buffers already grow — not per call, which would put a malloc in the hot parallel loop. The capacity is now named once, QWEN36_ATTN_MAX_CTX (262144 = max_position_embeddings), and enforced on all three paths that set max_t. Q36_MAXT still defaults to a conservative 8192 because context costs 40 KB per token here (10 attention layers, f32 KV; the 30 DeltaNet layers keep a fixed-size recurrent state instead), but it can now be raised safely. tests/test_qwen36_ctx.c gates it without a model file: KV layout, growth across requests, per-thread score rows at twice the old buffer, and that Q36_MAXT can lower the ceiling but never raise it past the capacity. Greedy output is bit-identical before and after. --- c/Makefile | 5 ++ c/qwen36.c | 58 +++++++++++++- c/tests/test_qwen36_ctx.c | 154 ++++++++++++++++++++++++++++++++++++++ 3 files changed, 215 insertions(+), 2 deletions(-) create mode 100644 c/tests/test_qwen36_ctx.c diff --git a/c/Makefile b/c/Makefile index f131359fc..21f9e8e7d 100644 --- a/c/Makefile +++ b/c/Makefile @@ -882,6 +882,11 @@ olmoe$(EXE): olmoe.c st.h json.h compat.h sample.h tok.h tok_unicode.h tok_unico qwen36$(EXE): qwen36.c st.h json.h compat.h $(CC) $(CFLAGS) qwen36.c -o qwen36$(EXE) $(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) + 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 9fd721e57..138568947 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -26,9 +26,33 @@ #include #include #include +#ifdef _OPENMP +#include +#endif #if defined(__AVX2__) #include #endif + +/* Hard context ceiling: the model's max_position_embeddings. Every buffer that + * scales with position (KV cache, attention score row) is allocated from max_t, + * so this is a policy limit, not a buffer limit -- but it is ONE limit, named + * once. It used to be the literal 8192 in two unrelated places: the size of a + * stack array in attention() and the default of Q36_MAXT in serve_one(). They + * agreed by luck, and raising Q36_MAXT moved the guard without moving the + * buffer, so a longer prompt overran the stack instead of being refused. + * Context costs 40 KB/token in KV (10 attention layers, f32) -- 128k is 5.0 GiB + * -- which is why Q36_MAXT still defaults far below this. */ +#define QWEN36_ATTN_MAX_CTX 262144 +#define QWEN36_DEFAULT_MAX_CTX 8192 + +/* Effective ceiling: Q36_MAXT if set and sane, the conservative default + * otherwise; never above the hard limit. */ +static int qwen36_max_ctx(void) { + const char *e = getenv("Q36_MAXT"); + int v = (e && *e) ? atoi(e) : QWEN36_DEFAULT_MAX_CTX; + if (v < 1) v = QWEN36_DEFAULT_MAX_CTX; + return v > QWEN36_ATTN_MAX_CTX ? QWEN36_ATTN_MAX_CTX : v; +} #if defined(__APPLE__) || defined(__linux__) || defined(__FreeBSD__) #include #include @@ -574,6 +598,8 @@ typedef struct { float **DN_conv; /* [n_layers] conv ring [conv_dim, convk-1] for DeltaNet layers (NULL for attn) */ uint64_t clock, hits, miss; float **K, **V; int kv_len, max_t, kv_cap; + float *attn_sc; /* [attn_sc_thr * kv_cap] score rows, one per thread */ + int attn_sc_thr; double dense_load_s; uint32_t *freq; int freq_token_count, hot_pinned, hot_n, warmup_tokens, token_count; @@ -1353,7 +1379,11 @@ static void attention(Model *m, Layer *l, int layer, float *x, int S, int pos_ba int kvh = hh / q_per_kv; int qpos = pos_base + s; const float *qv = query + ((int64_t)s*H + hh)*hd; - float sc[8192]; + int tid = 0; +#ifdef _OPENMP + tid = omp_get_thread_num(); +#endif + float *sc = m->attn_sc + (int64_t)tid * m->kv_cap; for (int t = 0; t <= qpos; t++) { const float *kv = m->K[layer] + ((int64_t)kvh*m->max_t + t)*kvd; float acc = 0; for (int dd = 0; dd < kvd; dd++) acc += qv[dd]*kv[dd]; @@ -1834,11 +1864,30 @@ static void ensure_kv(Model *m){ m->V[i] = falloc((int64_t)c->kv_heads * m->max_t * c->k_head_dim); } else { m->K[i] = NULL; m->V[i] = NULL; } } + /* Attention scores: one row per thread, indexed by absolute position, so + * each row must hold max_t entries. Sized here rather than in attention() + * because it grows with the context exactly like the KV cache does, and + * because a per-call allocation would run 10x per token. */ + free(m->attn_sc); + m->attn_sc_thr = 1; +#ifdef _OPENMP + m->attn_sc_thr = omp_get_max_threads(); + if (m->attn_sc_thr < 1) m->attn_sc_thr = 1; +#endif + m->attn_sc = falloc((int64_t)m->attn_sc_thr * m->max_t); m->kv_cap = m->max_t; } static void generate(Model *m, const int *prompt, int np, int n_new, int *out) { Cfg *c = &m->c; + /* Same ceiling serve_one() enforces. Past max_position_embeddings the RoPE + * positions leave the range the model was trained on, so this is a + * correctness limit, not just a memory one. */ + if (np + n_new > QWEN36_ATTN_MAX_CTX) { + fprintf(stderr, "[ctx] prompt %d + %d new exceeds the %d-token ceiling\n", + np, n_new, QWEN36_ATTN_MAX_CTX); + exit(1); + } m->max_t = np + n_new; reset_recurrent(m); ensure_kv(m); @@ -1868,6 +1917,11 @@ static void generate(Model *m, const int *prompt, int np, int n_new, int *out) { static int tf_nll(Model *m, const int *full, int nfull, int np, double *nll_out) { Cfg *c = &m->c; + if (nfull > QWEN36_ATTN_MAX_CTX) { + fprintf(stderr, "[ctx] %d tokens exceed the %d-token ceiling\n", + nfull, QWEN36_ATTN_MAX_CTX); + exit(1); + } m->max_t = nfull; reset_recurrent(m); ensure_kv(m); @@ -1984,7 +2038,7 @@ static int serve_eos_ids(int *ids, int cap){ static void serve_one(Model *m, ServeReq *q){ int *ids=NULL, np=0; encode_text(q->payload, &ids, &np); /* payload is raw prompt text; qwen36 adds no BOS */ - int max_ctx = getenv("Q36_MAXT")?atoi(getenv("Q36_MAXT")):8192; + int max_ctx = qwen36_max_ctx(); if(np<1 || np+q->max_tok>max_ctx){ printf("ERROR %s CONTEXT_EXCEEDED prompt_tokens=%d requested=%d capacity=%d\n",q->id,np,q->max_tok,max_ctx); fflush(stdout); free(ids); return; diff --git a/c/tests/test_qwen36_ctx.c b/c/tests/test_qwen36_ctx.c new file mode 100644 index 000000000..f6e036708 --- /dev/null +++ b/c/tests/test_qwen36_ctx.c @@ -0,0 +1,154 @@ +/* Context-size gates for the qwen36 hybrid engine. + * + * Regression test for the fixed-size score buffer: attention() used to keep one + * score per attended key in a stack `float sc[8192]`, filled with + * `for (t = 0; t <= qpos; t++)`. serve_one() carried the same literal as the + * Q36_MAXT default, and generate()/tf_nll() checked nothing at all. The two + * literals agreed only by luck: raising Q36_MAXT moved the guard without moving + * the buffer, so a longer prompt smashed the stack instead of being refused. + * Same class of bug as #122/#110 in the MLA path, and fixed the same way -- a + * per-thread heap buffer sized from max_t, plus one named capacity. + * + * The 40-layer hybrid also keeps two kinds of state, and only one of them + * grows: the 10 Gated Attention layers hold a KV cache, while the 30 Gated + * DeltaNet layers hold a recurrent state of fixed size and get no KV cache at + * all (ensure_kv leaves K[i]/V[i] NULL). Mixing those up silently allocates 4x + * the memory or indexes a NULL layer, so the layout is worth pinning too. + * + * No model file: ensure_kv only reads n_layers/is_attn/kv_heads/k_head_dim. + */ +#define main qwen36_main_unused +#include "../qwen36.c" +#undef main + +static int fails = 0; + +/* Same pattern as test_omp_tune.c / test_stops.c, and for the same reason: + * compat.h maps setenv() to SetEnvironmentVariableA, which updates the Win32 + * environment block -- but getenv() reads the CRT's own copy and never sees + * it, so the value under test would silently not arrive. _putenv_s updates + * the copy getenv() reads. */ +static void env_set(const char *name, const char *value) { +#ifdef _WIN32 + _putenv_s(name, value); +#else + setenv(name, value, 1); +#endif +} + +static void env_unset(const char *name) { +#ifdef _WIN32 + _putenv_s(name, ""); +#else + unsetenv(name); +#endif +} + +static void ck(int cond, const char *what) { + if (cond) { printf(" ok %s\n", what); return; } + printf(" FAIL %s\n", what); + fails++; +} + +/* the shipping Qwen3.6-35B-A3B geometry, from qwen36_meta.json */ +#define REAL_LAYERS 40 +#define REAL_KVHEADS 2 +#define REAL_KHEADDIM 256 +#define REAL_ATTN_LAYERS 10 /* i % 4 == 3 */ + +static void shape_model(Model *m, int n_layers, int kv_heads, int k_head_dim) { + memset(m, 0, sizeof(*m)); + m->c.n_layers = n_layers; + m->c.kv_heads = kv_heads; + m->c.k_head_dim = k_head_dim; + m->c.is_attn = calloc(n_layers, 1); + for (int i = 0; i < n_layers; i++) m->c.is_attn[i] = (i % 4 == 3); +} + +/* --- the layout every run depends on ------------------------------------- */ +static void case_layout(void) { + Model m; shape_model(&m, REAL_LAYERS, REAL_KVHEADS, REAL_KHEADDIM); + printf("KV layout at a short context\n"); + + m.max_t = 4096; + ensure_kv(&m); + ck(m.kv_cap == 4096, "kv_cap follows max_t"); + + int with_cache = 0, without = 0; + for (int i = 0; i < m.c.n_layers; i++) { + if (m.c.is_attn[i]) { if (m.K[i] && m.V[i]) with_cache++; } + else if (!m.K[i] && !m.V[i]) without++; + } + ck(with_cache == REAL_ATTN_LAYERS, "every attention layer has K and V"); + ck(without == REAL_LAYERS - REAL_ATTN_LAYERS, + "no DeltaNet layer has a KV cache (recurrent state instead of history)"); + + /* 2 kv heads * 256 dims * 4 B * 2 (K and V) * 10 layers */ + ck(2 * (int64_t)REAL_KVHEADS * REAL_KHEADDIM * (int64_t)sizeof(float) + * REAL_ATTN_LAYERS == 40960, "context costs 40 KB per token"); + + /* under ASan this is the actual out-of-bounds gate, not a formality */ + for (int i = 0; i < m.c.n_layers; i++) { + if (!m.c.is_attn[i]) continue; + int64_t last = (int64_t)m.c.kv_heads * m.max_t * m.c.k_head_dim - 1; + m.K[i][last] = 1.f; m.V[i][last] = 2.f; + } + ck(1, "last element of every cache is writable"); +} + +/* --- serve reuses one cache across requests of different lengths ---------- */ +static void case_growth(void) { + Model m; shape_model(&m, REAL_LAYERS, REAL_KVHEADS, REAL_KHEADDIM); + printf("growth across requests\n"); + + m.max_t = 512; ensure_kv(&m); + float *first = m.K[3]; + m.max_t = 256; ensure_kv(&m); + ck(m.K[3] == first && m.kv_cap == 512, + "a shorter request leaves the cache alone (never shrinks)"); + + m.max_t = 2048; ensure_kv(&m); + ck(m.kv_cap == 2048, "a longer request grows the cache"); + int64_t last = (int64_t)m.c.kv_heads * m.max_t * m.c.k_head_dim - 1; + m.K[3][last] = 1.f; /* the re-allocation must be sized to the NEW max_t */ + ck(1, "the grown cache is writable to its last element"); +} + +/* --- the ceiling itself --------------------------------------------------- */ +static void case_ceiling(void) { + printf("attention capacity\n"); + /* One capacity, named once, enforced on every path that sets max_t -- + * instead of a stack array in attention() and a second literal in + * serve_one() that happened to agree. */ + ck(QWEN36_ATTN_MAX_CTX >= 131072, + "declared capacity covers the model's context length"); + + /* The score rows must be sized from max_t and carry one row per thread, + * so a context past the old 8192 has somewhere to go. */ + Model m; shape_model(&m, 8, REAL_KVHEADS, 16); + m.max_t = 16384; /* twice the old fixed buffer */ + ensure_kv(&m); + ck(m.attn_sc != NULL && m.attn_sc_thr >= 1, "per-thread score rows allocated"); + for (int t = 0; t < m.attn_sc_thr; t++) + m.attn_sc[(int64_t)t * m.kv_cap + m.kv_cap - 1] = 1.f; + ck(1, "every thread's score row holds max_t entries"); + + /* Q36_MAXT may lower the ceiling but never raise it past the capacity. */ + env_set("Q36_MAXT", "999999999"); + ck(qwen36_max_ctx() == QWEN36_ATTN_MAX_CTX, + "Q36_MAXT cannot be raised past the capacity"); + env_set("Q36_MAXT", "4096"); + ck(qwen36_max_ctx() == 4096, "Q36_MAXT can lower the ceiling"); + env_unset("Q36_MAXT"); + ck(qwen36_max_ctx() == QWEN36_DEFAULT_MAX_CTX, + "unset Q36_MAXT falls back to the conservative default"); +} + +int main(void) { + case_layout(); + case_growth(); + case_ceiling(); + if (fails) { printf("FAILED %d\n", fails); return 1; } + printf("OK test_qwen36_ctx\n"); + return 0; +} From 3252675c9461df6ba1a1e2384eddfaa9cce23324 Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Sun, 16 Aug 2026 22:10:30 +0200 Subject: [PATCH 09/17] fix(qwen36): expert_get must not evict an in-flight slot (A5) expert_get's last resort was 'if (lru < 0) lru = 0;' -- take slot 0 no matter what. olmoe.c:471-486, which this cache derives from, deleted that exact line and replaced it with wait-and-rescan, and its comment names the consequence: the buffer is owned by an unlocked pread that will publish into it, so two writers race the same slab and whichever publishes last decides which expert the resident bytes answer to. Silently wrong weights, not a crash. Reachable in ordinary use, not only on a malformed container: PILOT queues up to 128 candidates per layer against cap slots, so any box small enough to run a low cache/layer can hit it. The wait-and-rescan loop is taken verbatim from olmoe, comment included -- the point of the fix is that the correction stops going missing when this cache is copied again. Also guards cap >= 1 (one row of the A2 table). Without it the new loop would wait for a publish that cannot come: cap=0 leaves every layer cache empty, so no slot ever becomes evictable. The old fallback turned that same input into a heap OOB via slots[0]. Neither is a failure mode to ship, and the guard belongs with the change that introduces the wait. --- c/qwen36.c | 30 +++++++++++++++++++++++++++++- 1 file changed, 29 insertions(+), 1 deletion(-) diff --git a/c/qwen36.c b/c/qwen36.c index 138568947..5d2bd738b 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -1199,15 +1199,39 @@ static void expert_get(Model *m, int layer, int eid, Slot **out) { Cfg *c = &m->c; Slot *s; if (lc->n < lc->cap) { s = &lc->slots[lc->n++]; slot_ensure_allocated(m, s); } else { + /* LRU eviction — skip pinned and in-flight (eid==-1) slots */ int lru = -1; for (int i = 0; i < lc->n; i++) { if (lc->slots[i].pinned || lc->slots[i].eid < 0) continue; if (lru < 0 || lc->slots[i].used < lc->slots[lru].used) lru = i; } if (lru < 0) { + /* All slots are pinned or in-flight; find the oldest non-in-flight + * slot (may be pinned, but never one currently being loaded). */ for (int i = 0; i < lc->n; i++) { if (lc->slots[i].eid < 0) continue; if (lru < 0 || lc->slots[i].used < lc->slots[lru].used) lru = i; } } - if (lru < 0) lru = 0; + while (lru < 0) { + /* EVERY slot is in flight: each buffer is owned by an unlocked pread + * in the pilot worker (or a demand load) that will publish into it. + * The old last resort (lru=0) stole such a slot mid-load — two writers + * racing the same slab, then whichever published last decided the + * expert id the resident bytes answered to. Wait for a publish instead + * and rescan; in-flight always drains because a load either finishes + * or the process is already dead in the water. + * + * Taken verbatim from olmoe.c, which this cache derives from and + * where this exact fallback was deleted for exactly this reason. + * Reachable whenever cap is smaller than the number of candidates a + * layer has in flight — PILOT queues up to 128 per layer — i.e. on + * any small-RAM box, and it corrupts silently rather than crashing. */ + pthread_mutex_unlock(&g_pilot_mx); + sleep_ms(1); + pthread_mutex_lock(&g_pilot_mx); + for (int i = 0; i < lc->n; i++) { + if (lc->slots[i].eid < 0) continue; + if (lru < 0 || lc->slots[i].used < lc->slots[lru].used) lru = i; + } + } s = &lc->slots[lru]; s->pinned = 0; } s->eid = -1; s->used = ++m->clock; @@ -2095,6 +2119,10 @@ int main(int argc, char **argv) { int hot_n = getenv("HOT") ? atoi(getenv("HOT")) : 0; int cap = argc > 1 ? atoi(argv[1]) : 16; int bits = argc > 2 ? atoi(argv[2]) : 4; + /* cap < 1 leaves every layer cache empty, so expert_get finds no slot to + * evict and waits for a publish that can never come. The old lru=0 fallback + * turned that into a heap OOB instead; neither is a failure mode to ship. */ + if (cap < 1) { fprintf(stderr, "cache/layer must be >= 1 (got %d)\n", cap); return 1; } if (bits < 2 || bits > 8) { fprintf(stderr, "quant_bits must be 2..8 (got %d)\n", bits); return 1; } const char *refpath = argc > 3 ? argv[3] : "ref.json"; From 9ae0c8a8c50b7c817cd748286870373f9e86765f Mon Sep 17 00:00:00 2001 From: kreuzzelg Date: Mon, 17 Aug 2026 20:44:38 +0200 Subject: [PATCH 10/17] ci: build and ship qwen36, and gate it against a tiny oracle (B5 + B5b) B5 -- the engine shipped nowhere. ci.yml built four engines, release.yml built, copied and asserted four; qwen36 was in none of them, so it was never compiled on Windows/macOS/aarch64 and a tagged archive would have carried a coli that dispatches to a missing binary. Third recurrence of #720/#858. Added at all four sites: ENGINES, the release build loop, the dist/ copy, and SIBLINGS. B5b -- a gate rather than a claim. qwen36-tiny-check generates the tiny Qwen3.6-shaped fixture, converts it, and requires token-exactness at cache capacities 1, 2 and 8. cap=1 evicts on every routed expert, which is where slot bookkeeping breaks -- the inkling job uses the same staffel for the same reason, and the in-flight eviction bug fixed in 4e79f0c is exactly that class. The run repeats under ASan+UBSan with PILOT=1. Token-exactness alone would not have caught the config-driven heap overflows in review section A: they do not necessarily change the output. Two things the wiring needed that the suggestion could not know: - The reference comes from make_qwen36_tiny.py, not make_qwen36_oracle.py. The oracle encodes a text prompt via AutoTokenizer.from_pretrained(); this fixture is synthetic and has no tokenizer. The tiny script already holds the model, so it gained --ref-mode full -- its ref.json was attention_only only, which replaces the 30 DeltaNet layers with identity and would have left three quarters of the engine untested. - The engine returned 0 regardless of the result in oracle mode, printing 'Matching tokens: n/m' and exiting green. inkling.c returns (match == ngen) ? 0 : 1 and its CI job depends on it. Without this the whole gate could not fail. Verified locally end to end: 16/16 at cap 1, 2 and 8; ASan+UBSan clean. A malformed-container case belongs here once A1/A2 land -- noted as a TODO, since a well-formed fixture cannot reproduce what ASan found. --- .github/workflows/ci.yml | 67 ++++++++++++++++++++++++++++++++++- .github/workflows/release.yml | 10 ++++-- c/qwen36.c | 7 ++++ c/tools/make_qwen36_tiny.py | 40 +++++++++++++++------ 4 files changed, 109 insertions(+), 15 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 321ab856c..805bd5b4e 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -190,7 +190,7 @@ jobs: run: | cd c rc=0 - ENGINES="colibri inkling kimi_k3 olmoe" + ENGINES="colibri inkling kimi_k3 olmoe qwen36" if [ "${{ matrix.v4 }}" = "1" ]; then ENGINES="$ENGINES deepseek-v4"; fi for t in $ENGINES; do echo "::group::$t" @@ -317,6 +317,71 @@ jobs: -Xcompiler=-Wall,-Wextra echo "inkling CUDA syntax check passed" + qwen36-tiny-check: + name: Qwen3.6 tiny oracle (token-exact + ASan/UBSan) + runs-on: ubuntu-latest + steps: + - uses: actions/checkout@v4 + - uses: actions/setup-python@v5 + with: + python-version: '3.12' + cache: pip + cache-dependency-path: c/tools/oracle-requirements.txt + - name: Install torch (CPU) + transformers + run: pip install -r c/tools/oracle-requirements.txt + - name: Tiny Qwen3.6-shaped fixture + full-hybrid oracle + run: | + cd c + # Same layout as the 35B model (10 x (3 x DeltaNet -> MoE, 1 x Attention + # -> MoE)) at toy dimensions, so the converter and the engine treat it + # exactly like the real one. --mode full exercises BOTH layer kinds; + # attention_only would leave the DeltaNet path untested. + # The reference comes from make_qwen36_tiny.py, not make_qwen36_oracle.py: + # the oracle encodes a text prompt through AutoTokenizer.from_pretrained(), + # and this fixture is synthetic -- weights, no tokenizer. The tiny script + # already has the model in memory and generates greedily from fixed ids. + # --ref-mode full leaves BOTH layer kinds active; the default + # attention_only replaces the 30 DeltaNet layers with identity, which is + # what Phase 1 computed and would leave three quarters of the engine + # untested. + python3 tools/make_qwen36_tiny.py --out qwen36_tiny --ref-mode full \ + --emit-ref qwen36_tiny/ref_full.json + # ebits=8 keeps the expert quantization error far below anything that + # could flip a greedy argmax, so a mismatch means the engine changed, + # not the container. + python3 tools/convert_qwen36.py --model qwen36_tiny --out qwen36_tiny_c --ebits 8 + - name: Token-exact against the oracle, at several cache capacities + run: | + cd c + make qwen36 + # cap=1 evicts on every routed expert, which is where slot bookkeeping + # breaks; cap=8 never would. That is not hypothetical here: the + # in-flight eviction fallback this engine inherited was exactly such a + # bug, and it only shows when cap is smaller than the number of loads + # in flight. The engine exits non-zero on any token mismatch. + for cap in 1 2 8; do + echo "::group::cap=$cap" + SNAP=qwen36_tiny_c ./qwen36 "$cap" 8 qwen36_tiny/ref_full.json + echo "::endgroup::" + done + - name: Same run under ASan + UBSan + run: | + cd c + # Token-exactness alone would not have caught the config-driven heap + # overflows this engine shipped with: they do not necessarily change + # the output. The sanitizers are the half of this gate that watches + # memory, and PILOT=1 puts concurrent expert loads against the cache + # so the slot paths are exercised, not just walked past. + make clean >/dev/null 2>&1 || true + make qwen36 EXTRA_CFLAGS="-fsanitize=address,undefined -fno-omit-frame-pointer -g" + ASAN_OPTIONS=detect_leaks=0 UBSAN_OPTIONS=halt_on_error=1 \ + SNAP=qwen36_tiny_c PILOT=1 WIDE=2 \ + ./qwen36 1 8 qwen36_tiny/ref_full.json + # TODO once A1/A2 land: add a malformed-container case here (config.json + # and qwen36_meta.json disagreeing on the layer count) and assert the + # engine REFUSES it. That is the case ASan caught, and a well-formed + # fixture cannot reproduce it -- this step guards the happy path only. + inkling-oracle: name: Inkling oracle (token-exact vs transformers) runs-on: ubuntu-latest diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index b7036e932..0c1579edb 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -62,10 +62,10 @@ jobs: - name: Build engines run: | cd c - for t in colibri inkling kimi_k3 olmoe; do + for t in colibri inkling kimi_k3 olmoe qwen36; do make $t ${{ matrix.make_args }} done - ls -lh colibri${{ matrix.ext }} inkling${{ matrix.ext }} kimi_k3${{ matrix.ext }} olmoe${{ matrix.ext }} + ls -lh colibri${{ matrix.ext }} inkling${{ matrix.ext }} kimi_k3${{ matrix.ext }} olmoe${{ matrix.ext }} qwen36${{ matrix.ext }} # DeepSeek V4 has its own Makefile and its own target name (`deepseek-v4`, # producing `deepseek_v4`), so it never joined the loop above -- and #858 is @@ -120,6 +120,10 @@ jobs: cp c/inkling${{ matrix.ext }} dist/inkling${{ matrix.ext }} cp c/kimi_k3${{ matrix.ext }} dist/kimi_k3${{ matrix.ext }} cp c/olmoe${{ matrix.ext }} dist/olmoe${{ matrix.ext }} + # qwen36 -- plain name, like the others: engine_for() resolves "qwen" to + # a "qwen36" binary next to coli. Unconditional: the engine is portable C + # with no platform gate, so every archive can carry it. + cp c/qwen36${{ matrix.ext }} dist/qwen36${{ matrix.ext }} # deepseek_v4 -- underscore, not the hyphen of the make target: engine_for() # looks for "deepseek_v4" next to coli. Guarded, not unconditional, because # the macos-arm64 archive has no such engine to ship (#858). @@ -168,7 +172,7 @@ jobs: # a green build, green tests, and an archive with no engine for the model the # user actually had. Presence is not enough -- assert they are executable, and # that the launcher's own resolver finds them where it looks (next to coli). - SIBLINGS="inkling kimi_k3 olmoe" + SIBLINGS="inkling kimi_k3 olmoe qwen36" # deepseek_v4 only where the engine builds (COLI_V4_SUPPORTED); asserting it # unconditionally would fail the macos-arm64 archive for shipping an engine # that platform cannot have. diff --git a/c/qwen36.c b/c/qwen36.c index 5d2bd738b..3f228e1a6 100644 --- a/c/qwen36.c +++ b/c/qwen36.c @@ -2267,12 +2267,14 @@ int main(int argc, char **argv) { else fprintf(stderr, "[dump] cannot open %s\n", dp ? dp : "qwen36_logits.f32"); } + int ref_match = 0; if (is_ref) { int match = 0; printf("\nReference: "); for (int i=np;i