From 364b9741e26a3127c40f560ca5d9e527857445be Mon Sep 17 00:00:00 2001 From: Nicholas Beerbower Date: Wed, 15 Jul 2026 23:10:10 -0400 Subject: [PATCH 01/71] tok: o200k pre-tokenizer support, auto-detected from tokenizer.json Inkling ships an o200k-family tokenizer (case-aware Split regex, GPT-4o lineage) rather than cl100k. tok_load now detects the family from the pattern itself (\p{Lu} appears only in the o200k regex) so GLM behavior is untouched, and encode dispatches to a new pretok_chunk_o200k that replays the regex engine's backtracking order exactly: greedy optional prefix, maximally-greedy uppercase run given back until the lowercase run can match, contractions attached to letter runs, \p{N}{1,3}, and the [\r\n/]* punctuation tail. tok_unicode_o200k.h adds the two range tables the new classes need (Lu+Lt and Lm+Lo+M), generated from Python unicodedata. Validated against HF tokenizers on 357 adversarial strings (case transitions, contractions, CJK, combining marks, emoji + modifiers, zero-width chars, 300 mixed-charset fuzz cases): 357/357 identical. Co-Authored-By: Claude Fable 5 --- c/tok.h | 116 ++++++++++++++++++++- c/tok_unicode_o200k.h | 228 ++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 343 insertions(+), 1 deletion(-) create mode 100644 c/tok_unicode_o200k.h diff --git a/c/tok.h b/c/tok.h index 7d1140b..6b957df 100644 --- a/c/tok.h +++ b/c/tok.h @@ -19,6 +19,7 @@ #include #include "json.h" #include "tok_unicode.h" +#include "tok_unicode_o200k.h" /* ---------- hash map (chiavi binarie con lunghezza) ---------- */ typedef struct { const char *k; int klen; int v; int used; } ment; @@ -50,6 +51,7 @@ typedef struct { Special *sp; int nsp; /* added tokens, ordinati per lunghezza decrescente */ uint32_t byte2cp[256]; int byte2cp_len[256]; char byte2str[256][3]; int16_t cp2byte[1024]; + int o200k; /* pre_tokenizer regex family: 0 = cl100k (GLM), 1 = o200k (Inkling) */ } Tok; /* ---------- UTF-8 ---------- */ @@ -144,6 +146,17 @@ static void tok_load(Tok *T, const char *path){ } qsort(T->sp,T->nsp,sizeof(Special),cmp_sp_len); /* match piu' lungo per primo */ } + /* pre_tokenizer family: the o200k Split regex is recognizable by its + * case-category classes (\p{Lu}...) which cl100k does not use */ + jval *pt=json_get(root,"pre_tokenizer"); + if(pt){ + jval *ps=json_get(pt,"pretokenizers"); + if(ps&&ps->t==J_ARR) for(int i=0;ilen;i++){ + jval *pat=json_get(ps->kids[i],"pattern"); + jval *rx=pat?json_get(pat,"Regex"):NULL; + if(rx&&rx->t==J_STR&&strstr(rx->str,"\\p{Lu}")) T->o200k=1; + } + } /* arena/buf restano allocati: le stringhe (j_dup) sono malloc indipendenti e ci servono vive */ (void)arena; } @@ -241,6 +254,104 @@ static void pretok_chunk(Tok *T, const unsigned char *p, int a, int b, int *out, free(cp); free(off); } +/* ---------- pre-tokenizer o200k (Inkling / GPT-4o family) ---------- + * Split regex: + * A: [^\r\n\p{L}\p{N}]?[\p{Lu}\p{Lt}\p{Lm}\p{Lo}\p{M}]*[\p{Ll}\p{Lm}\p{Lo}\p{M}]+(?i:'s|'t|'re|'ve|'m|'ll|'d)? + * B: [^\r\n\p{L}\p{N}]?[\p{Lu}\p{Lt}\p{Lm}\p{Lo}\p{M}]+[\p{Ll}\p{Lm}\p{Lo}\p{M}]*(?i:'s|'t|'re|'ve|'m|'ll|'d)? + * C: \p{N}{1,3} D: ' ?[^\s\p{L}\p{N}]+[\r\n/]*' E: \s*[\r\n]+ F: \s+(?!\S) G: \s+ + * S1 = Lu|Lt|Lm|Lo|M, S2 = Ll|Lm|Lo|M. The letter matcher below replays the + * regex engine's backtracking order exactly: A with greedy optional prefix and + * maximally-greedy S1* given back until S2+ can take >=1 char, then B. */ +#define O2_S1(c) (is_U(c)||is_X(c)) +#define O2_S2(c) (is_X(c)||(is_L(c)&&!is_U(c))) +static uint32_t o2_low(uint32_t c){ return (c>='A'&&c<='Z')?c+32:c; } +static int o2_contraction(const uint32_t *cp, int n, int k){ + if(k=0; pfx--){ + int j0=i; + if(pfx){ + uint32_t c=cp[i]; + if(c=='\r'||c=='\n'||is_L(c)||is_N(c)||i+1>=n) continue; + j0=i+1; + } + int m1=j0; while(m1=j0; s--){ + if(s=0; pfx--){ + int j0=i; + if(pfx){ + uint32_t c=cp[i]; + if(c=='\r'||c=='\n'||is_L(c)||is_N(c)||i+1>=n) continue; + j0=i+1; + } + int m1=j0; while(m1j0){ + int k=m1; while(ki){ i=e; bpe_piece(T,p,off[start],off[i],out,no,max); continue; } + } + /* C: \p{N}{1,3} */ + if(is_N(c)){ int j=i,k=0; while(ji){ int last=-1; for(int j=i;j=0){ i=last+1; bpe_piece(T,p,off[start],off[i],out,no,max); continue; } + int end = (r id (split sugli added token, poi pretok+BPE) ---------- */ static int tok_encode(Tok *T, const char *text, int len, int *out, int max){ const unsigned char *p=(const unsigned char*)text; int no=0; int i=0; @@ -254,7 +365,10 @@ static int tok_encode(Tok *T, const char *text, int len, int *out, int max){ } } int chunk_end = (hitpos<0) ? len : hitpos; - if(chunk_end>i) pretok_chunk(T,p,i,chunk_end,out,&no,max); + if(chunk_end>i){ + if(T->o200k) pretok_chunk_o200k(T,p,i,chunk_end,out,&no,max); + else pretok_chunk(T,p,i,chunk_end,out,&no,max); + } if(hitpos<0) break; if(no + +static const uint32_t uni_U[][2] = { + {0x41,0x5A},{0xC0,0xD6},{0xD8,0xDE},{0x100,0x100},{0x102,0x102},{0x104,0x104}, + {0x106,0x106},{0x108,0x108},{0x10A,0x10A},{0x10C,0x10C},{0x10E,0x10E},{0x110,0x110}, + {0x112,0x112},{0x114,0x114},{0x116,0x116},{0x118,0x118},{0x11A,0x11A},{0x11C,0x11C}, + {0x11E,0x11E},{0x120,0x120},{0x122,0x122},{0x124,0x124},{0x126,0x126},{0x128,0x128}, + {0x12A,0x12A},{0x12C,0x12C},{0x12E,0x12E},{0x130,0x130},{0x132,0x132},{0x134,0x134}, + {0x136,0x136},{0x139,0x139},{0x13B,0x13B},{0x13D,0x13D},{0x13F,0x13F},{0x141,0x141}, + {0x143,0x143},{0x145,0x145},{0x147,0x147},{0x14A,0x14A},{0x14C,0x14C},{0x14E,0x14E}, + {0x150,0x150},{0x152,0x152},{0x154,0x154},{0x156,0x156},{0x158,0x158},{0x15A,0x15A}, + {0x15C,0x15C},{0x15E,0x15E},{0x160,0x160},{0x162,0x162},{0x164,0x164},{0x166,0x166}, + {0x168,0x168},{0x16A,0x16A},{0x16C,0x16C},{0x16E,0x16E},{0x170,0x170},{0x172,0x172}, + {0x174,0x174},{0x176,0x176},{0x178,0x179},{0x17B,0x17B},{0x17D,0x17D},{0x181,0x182}, + {0x184,0x184},{0x186,0x187},{0x189,0x18B},{0x18E,0x191},{0x193,0x194},{0x196,0x198}, + {0x19C,0x19D},{0x19F,0x1A0},{0x1A2,0x1A2},{0x1A4,0x1A4},{0x1A6,0x1A7},{0x1A9,0x1A9}, + {0x1AC,0x1AC},{0x1AE,0x1AF},{0x1B1,0x1B3},{0x1B5,0x1B5},{0x1B7,0x1B8},{0x1BC,0x1BC}, + {0x1C4,0x1C5},{0x1C7,0x1C8},{0x1CA,0x1CB},{0x1CD,0x1CD},{0x1CF,0x1CF},{0x1D1,0x1D1}, + {0x1D3,0x1D3},{0x1D5,0x1D5},{0x1D7,0x1D7},{0x1D9,0x1D9},{0x1DB,0x1DB},{0x1DE,0x1DE}, + {0x1E0,0x1E0},{0x1E2,0x1E2},{0x1E4,0x1E4},{0x1E6,0x1E6},{0x1E8,0x1E8},{0x1EA,0x1EA}, + {0x1EC,0x1EC},{0x1EE,0x1EE},{0x1F1,0x1F2},{0x1F4,0x1F4},{0x1F6,0x1F8},{0x1FA,0x1FA}, + {0x1FC,0x1FC},{0x1FE,0x1FE},{0x200,0x200},{0x202,0x202},{0x204,0x204},{0x206,0x206}, + {0x208,0x208},{0x20A,0x20A},{0x20C,0x20C},{0x20E,0x20E},{0x210,0x210},{0x212,0x212}, + {0x214,0x214},{0x216,0x216},{0x218,0x218},{0x21A,0x21A},{0x21C,0x21C},{0x21E,0x21E}, + {0x220,0x220},{0x222,0x222},{0x224,0x224},{0x226,0x226},{0x228,0x228},{0x22A,0x22A}, + {0x22C,0x22C},{0x22E,0x22E},{0x230,0x230},{0x232,0x232},{0x23A,0x23B},{0x23D,0x23E}, + {0x241,0x241},{0x243,0x246},{0x248,0x248},{0x24A,0x24A},{0x24C,0x24C},{0x24E,0x24E}, + {0x370,0x370},{0x372,0x372},{0x376,0x376},{0x37F,0x37F},{0x386,0x386},{0x388,0x38A}, + {0x38C,0x38C},{0x38E,0x38F},{0x391,0x3A1},{0x3A3,0x3AB},{0x3CF,0x3CF},{0x3D2,0x3D4}, + {0x3D8,0x3D8},{0x3DA,0x3DA},{0x3DC,0x3DC},{0x3DE,0x3DE},{0x3E0,0x3E0},{0x3E2,0x3E2}, + {0x3E4,0x3E4},{0x3E6,0x3E6},{0x3E8,0x3E8},{0x3EA,0x3EA},{0x3EC,0x3EC},{0x3EE,0x3EE}, + {0x3F4,0x3F4},{0x3F7,0x3F7},{0x3F9,0x3FA},{0x3FD,0x42F},{0x460,0x460},{0x462,0x462}, + {0x464,0x464},{0x466,0x466},{0x468,0x468},{0x46A,0x46A},{0x46C,0x46C},{0x46E,0x46E}, + {0x470,0x470},{0x472,0x472},{0x474,0x474},{0x476,0x476},{0x478,0x478},{0x47A,0x47A}, + {0x47C,0x47C},{0x47E,0x47E},{0x480,0x480},{0x48A,0x48A},{0x48C,0x48C},{0x48E,0x48E}, + {0x490,0x490},{0x492,0x492},{0x494,0x494},{0x496,0x496},{0x498,0x498},{0x49A,0x49A}, + {0x49C,0x49C},{0x49E,0x49E},{0x4A0,0x4A0},{0x4A2,0x4A2},{0x4A4,0x4A4},{0x4A6,0x4A6}, + {0x4A8,0x4A8},{0x4AA,0x4AA},{0x4AC,0x4AC},{0x4AE,0x4AE},{0x4B0,0x4B0},{0x4B2,0x4B2}, + {0x4B4,0x4B4},{0x4B6,0x4B6},{0x4B8,0x4B8},{0x4BA,0x4BA},{0x4BC,0x4BC},{0x4BE,0x4BE}, + {0x4C0,0x4C1},{0x4C3,0x4C3},{0x4C5,0x4C5},{0x4C7,0x4C7},{0x4C9,0x4C9},{0x4CB,0x4CB}, + {0x4CD,0x4CD},{0x4D0,0x4D0},{0x4D2,0x4D2},{0x4D4,0x4D4},{0x4D6,0x4D6},{0x4D8,0x4D8}, + {0x4DA,0x4DA},{0x4DC,0x4DC},{0x4DE,0x4DE},{0x4E0,0x4E0},{0x4E2,0x4E2},{0x4E4,0x4E4}, + {0x4E6,0x4E6},{0x4E8,0x4E8},{0x4EA,0x4EA},{0x4EC,0x4EC},{0x4EE,0x4EE},{0x4F0,0x4F0}, + {0x4F2,0x4F2},{0x4F4,0x4F4},{0x4F6,0x4F6},{0x4F8,0x4F8},{0x4FA,0x4FA},{0x4FC,0x4FC}, + {0x4FE,0x4FE},{0x500,0x500},{0x502,0x502},{0x504,0x504},{0x506,0x506},{0x508,0x508}, + {0x50A,0x50A},{0x50C,0x50C},{0x50E,0x50E},{0x510,0x510},{0x512,0x512},{0x514,0x514}, + {0x516,0x516},{0x518,0x518},{0x51A,0x51A},{0x51C,0x51C},{0x51E,0x51E},{0x520,0x520}, + {0x522,0x522},{0x524,0x524},{0x526,0x526},{0x528,0x528},{0x52A,0x52A},{0x52C,0x52C}, + {0x52E,0x52E},{0x531,0x556},{0x10A0,0x10C5},{0x10C7,0x10C7},{0x10CD,0x10CD},{0x13A0,0x13F5}, + {0x1C90,0x1CBA},{0x1CBD,0x1CBF},{0x1E00,0x1E00},{0x1E02,0x1E02},{0x1E04,0x1E04},{0x1E06,0x1E06}, + {0x1E08,0x1E08},{0x1E0A,0x1E0A},{0x1E0C,0x1E0C},{0x1E0E,0x1E0E},{0x1E10,0x1E10},{0x1E12,0x1E12}, + {0x1E14,0x1E14},{0x1E16,0x1E16},{0x1E18,0x1E18},{0x1E1A,0x1E1A},{0x1E1C,0x1E1C},{0x1E1E,0x1E1E}, + {0x1E20,0x1E20},{0x1E22,0x1E22},{0x1E24,0x1E24},{0x1E26,0x1E26},{0x1E28,0x1E28},{0x1E2A,0x1E2A}, + {0x1E2C,0x1E2C},{0x1E2E,0x1E2E},{0x1E30,0x1E30},{0x1E32,0x1E32},{0x1E34,0x1E34},{0x1E36,0x1E36}, + {0x1E38,0x1E38},{0x1E3A,0x1E3A},{0x1E3C,0x1E3C},{0x1E3E,0x1E3E},{0x1E40,0x1E40},{0x1E42,0x1E42}, + {0x1E44,0x1E44},{0x1E46,0x1E46},{0x1E48,0x1E48},{0x1E4A,0x1E4A},{0x1E4C,0x1E4C},{0x1E4E,0x1E4E}, + {0x1E50,0x1E50},{0x1E52,0x1E52},{0x1E54,0x1E54},{0x1E56,0x1E56},{0x1E58,0x1E58},{0x1E5A,0x1E5A}, + {0x1E5C,0x1E5C},{0x1E5E,0x1E5E},{0x1E60,0x1E60},{0x1E62,0x1E62},{0x1E64,0x1E64},{0x1E66,0x1E66}, + {0x1E68,0x1E68},{0x1E6A,0x1E6A},{0x1E6C,0x1E6C},{0x1E6E,0x1E6E},{0x1E70,0x1E70},{0x1E72,0x1E72}, + {0x1E74,0x1E74},{0x1E76,0x1E76},{0x1E78,0x1E78},{0x1E7A,0x1E7A},{0x1E7C,0x1E7C},{0x1E7E,0x1E7E}, + {0x1E80,0x1E80},{0x1E82,0x1E82},{0x1E84,0x1E84},{0x1E86,0x1E86},{0x1E88,0x1E88},{0x1E8A,0x1E8A}, + {0x1E8C,0x1E8C},{0x1E8E,0x1E8E},{0x1E90,0x1E90},{0x1E92,0x1E92},{0x1E94,0x1E94},{0x1E9E,0x1E9E}, + {0x1EA0,0x1EA0},{0x1EA2,0x1EA2},{0x1EA4,0x1EA4},{0x1EA6,0x1EA6},{0x1EA8,0x1EA8},{0x1EAA,0x1EAA}, + {0x1EAC,0x1EAC},{0x1EAE,0x1EAE},{0x1EB0,0x1EB0},{0x1EB2,0x1EB2},{0x1EB4,0x1EB4},{0x1EB6,0x1EB6}, + {0x1EB8,0x1EB8},{0x1EBA,0x1EBA},{0x1EBC,0x1EBC},{0x1EBE,0x1EBE},{0x1EC0,0x1EC0},{0x1EC2,0x1EC2}, + {0x1EC4,0x1EC4},{0x1EC6,0x1EC6},{0x1EC8,0x1EC8},{0x1ECA,0x1ECA},{0x1ECC,0x1ECC},{0x1ECE,0x1ECE}, + {0x1ED0,0x1ED0},{0x1ED2,0x1ED2},{0x1ED4,0x1ED4},{0x1ED6,0x1ED6},{0x1ED8,0x1ED8},{0x1EDA,0x1EDA}, + {0x1EDC,0x1EDC},{0x1EDE,0x1EDE},{0x1EE0,0x1EE0},{0x1EE2,0x1EE2},{0x1EE4,0x1EE4},{0x1EE6,0x1EE6}, + {0x1EE8,0x1EE8},{0x1EEA,0x1EEA},{0x1EEC,0x1EEC},{0x1EEE,0x1EEE},{0x1EF0,0x1EF0},{0x1EF2,0x1EF2}, + {0x1EF4,0x1EF4},{0x1EF6,0x1EF6},{0x1EF8,0x1EF8},{0x1EFA,0x1EFA},{0x1EFC,0x1EFC},{0x1EFE,0x1EFE}, + {0x1F08,0x1F0F},{0x1F18,0x1F1D},{0x1F28,0x1F2F},{0x1F38,0x1F3F},{0x1F48,0x1F4D},{0x1F59,0x1F59}, + {0x1F5B,0x1F5B},{0x1F5D,0x1F5D},{0x1F5F,0x1F5F},{0x1F68,0x1F6F},{0x1F88,0x1F8F},{0x1F98,0x1F9F}, + {0x1FA8,0x1FAF},{0x1FB8,0x1FBC},{0x1FC8,0x1FCC},{0x1FD8,0x1FDB},{0x1FE8,0x1FEC},{0x1FF8,0x1FFC}, + {0x2102,0x2102},{0x2107,0x2107},{0x210B,0x210D},{0x2110,0x2112},{0x2115,0x2115},{0x2119,0x211D}, + {0x2124,0x2124},{0x2126,0x2126},{0x2128,0x2128},{0x212A,0x212D},{0x2130,0x2133},{0x213E,0x213F}, + {0x2145,0x2145},{0x2183,0x2183},{0x2C00,0x2C2E},{0x2C60,0x2C60},{0x2C62,0x2C64},{0x2C67,0x2C67}, + {0x2C69,0x2C69},{0x2C6B,0x2C6B},{0x2C6D,0x2C70},{0x2C72,0x2C72},{0x2C75,0x2C75},{0x2C7E,0x2C80}, + {0x2C82,0x2C82},{0x2C84,0x2C84},{0x2C86,0x2C86},{0x2C88,0x2C88},{0x2C8A,0x2C8A},{0x2C8C,0x2C8C}, + {0x2C8E,0x2C8E},{0x2C90,0x2C90},{0x2C92,0x2C92},{0x2C94,0x2C94},{0x2C96,0x2C96},{0x2C98,0x2C98}, + {0x2C9A,0x2C9A},{0x2C9C,0x2C9C},{0x2C9E,0x2C9E},{0x2CA0,0x2CA0},{0x2CA2,0x2CA2},{0x2CA4,0x2CA4}, + {0x2CA6,0x2CA6},{0x2CA8,0x2CA8},{0x2CAA,0x2CAA},{0x2CAC,0x2CAC},{0x2CAE,0x2CAE},{0x2CB0,0x2CB0}, + {0x2CB2,0x2CB2},{0x2CB4,0x2CB4},{0x2CB6,0x2CB6},{0x2CB8,0x2CB8},{0x2CBA,0x2CBA},{0x2CBC,0x2CBC}, + {0x2CBE,0x2CBE},{0x2CC0,0x2CC0},{0x2CC2,0x2CC2},{0x2CC4,0x2CC4},{0x2CC6,0x2CC6},{0x2CC8,0x2CC8}, + {0x2CCA,0x2CCA},{0x2CCC,0x2CCC},{0x2CCE,0x2CCE},{0x2CD0,0x2CD0},{0x2CD2,0x2CD2},{0x2CD4,0x2CD4}, + {0x2CD6,0x2CD6},{0x2CD8,0x2CD8},{0x2CDA,0x2CDA},{0x2CDC,0x2CDC},{0x2CDE,0x2CDE},{0x2CE0,0x2CE0}, + {0x2CE2,0x2CE2},{0x2CEB,0x2CEB},{0x2CED,0x2CED},{0x2CF2,0x2CF2},{0xA640,0xA640},{0xA642,0xA642}, + {0xA644,0xA644},{0xA646,0xA646},{0xA648,0xA648},{0xA64A,0xA64A},{0xA64C,0xA64C},{0xA64E,0xA64E}, + {0xA650,0xA650},{0xA652,0xA652},{0xA654,0xA654},{0xA656,0xA656},{0xA658,0xA658},{0xA65A,0xA65A}, + {0xA65C,0xA65C},{0xA65E,0xA65E},{0xA660,0xA660},{0xA662,0xA662},{0xA664,0xA664},{0xA666,0xA666}, + {0xA668,0xA668},{0xA66A,0xA66A},{0xA66C,0xA66C},{0xA680,0xA680},{0xA682,0xA682},{0xA684,0xA684}, + {0xA686,0xA686},{0xA688,0xA688},{0xA68A,0xA68A},{0xA68C,0xA68C},{0xA68E,0xA68E},{0xA690,0xA690}, + {0xA692,0xA692},{0xA694,0xA694},{0xA696,0xA696},{0xA698,0xA698},{0xA69A,0xA69A},{0xA722,0xA722}, + {0xA724,0xA724},{0xA726,0xA726},{0xA728,0xA728},{0xA72A,0xA72A},{0xA72C,0xA72C},{0xA72E,0xA72E}, + {0xA732,0xA732},{0xA734,0xA734},{0xA736,0xA736},{0xA738,0xA738},{0xA73A,0xA73A},{0xA73C,0xA73C}, + {0xA73E,0xA73E},{0xA740,0xA740},{0xA742,0xA742},{0xA744,0xA744},{0xA746,0xA746},{0xA748,0xA748}, + {0xA74A,0xA74A},{0xA74C,0xA74C},{0xA74E,0xA74E},{0xA750,0xA750},{0xA752,0xA752},{0xA754,0xA754}, + {0xA756,0xA756},{0xA758,0xA758},{0xA75A,0xA75A},{0xA75C,0xA75C},{0xA75E,0xA75E},{0xA760,0xA760}, + {0xA762,0xA762},{0xA764,0xA764},{0xA766,0xA766},{0xA768,0xA768},{0xA76A,0xA76A},{0xA76C,0xA76C}, + {0xA76E,0xA76E},{0xA779,0xA779},{0xA77B,0xA77B},{0xA77D,0xA77E},{0xA780,0xA780},{0xA782,0xA782}, + {0xA784,0xA784},{0xA786,0xA786},{0xA78B,0xA78B},{0xA78D,0xA78D},{0xA790,0xA790},{0xA792,0xA792}, + {0xA796,0xA796},{0xA798,0xA798},{0xA79A,0xA79A},{0xA79C,0xA79C},{0xA79E,0xA79E},{0xA7A0,0xA7A0}, + {0xA7A2,0xA7A2},{0xA7A4,0xA7A4},{0xA7A6,0xA7A6},{0xA7A8,0xA7A8},{0xA7AA,0xA7AE},{0xA7B0,0xA7B4}, + {0xA7B6,0xA7B6},{0xA7B8,0xA7B8},{0xA7BA,0xA7BA},{0xA7BC,0xA7BC},{0xA7BE,0xA7BE},{0xA7C2,0xA7C2}, + {0xA7C4,0xA7C7},{0xA7C9,0xA7C9},{0xA7F5,0xA7F5},{0xFF21,0xFF3A},{0x10400,0x10427},{0x104B0,0x104D3}, + {0x10C80,0x10CB2},{0x118A0,0x118BF},{0x16E40,0x16E5F},{0x1D400,0x1D419},{0x1D434,0x1D44D},{0x1D468,0x1D481}, + {0x1D49C,0x1D49C},{0x1D49E,0x1D49F},{0x1D4A2,0x1D4A2},{0x1D4A5,0x1D4A6},{0x1D4A9,0x1D4AC},{0x1D4AE,0x1D4B5}, + {0x1D4D0,0x1D4E9},{0x1D504,0x1D505},{0x1D507,0x1D50A},{0x1D50D,0x1D514},{0x1D516,0x1D51C},{0x1D538,0x1D539}, + {0x1D53B,0x1D53E},{0x1D540,0x1D544},{0x1D546,0x1D546},{0x1D54A,0x1D550},{0x1D56C,0x1D585},{0x1D5A0,0x1D5B9}, + {0x1D5D4,0x1D5ED},{0x1D608,0x1D621},{0x1D63C,0x1D655},{0x1D670,0x1D689},{0x1D6A8,0x1D6C0},{0x1D6E2,0x1D6FA}, + {0x1D71C,0x1D734},{0x1D756,0x1D76E},{0x1D790,0x1D7A8},{0x1D7CA,0x1D7CA},{0x1E900,0x1E921}, +}; +static const int uni_U_n = 641; + +static const uint32_t uni_X[][2] = { + {0xAA,0xAA},{0xBA,0xBA},{0x1BB,0x1BB},{0x1C0,0x1C3},{0x294,0x294},{0x2B0,0x2C1}, + {0x2C6,0x2D1},{0x2E0,0x2E4},{0x2EC,0x2EC},{0x2EE,0x2EE},{0x300,0x36F},{0x374,0x374}, + {0x37A,0x37A},{0x483,0x489},{0x559,0x559},{0x591,0x5BD},{0x5BF,0x5BF},{0x5C1,0x5C2}, + {0x5C4,0x5C5},{0x5C7,0x5C7},{0x5D0,0x5EA},{0x5EF,0x5F2},{0x610,0x61A},{0x620,0x65F}, + {0x66E,0x6D3},{0x6D5,0x6DC},{0x6DF,0x6E8},{0x6EA,0x6EF},{0x6FA,0x6FC},{0x6FF,0x6FF}, + {0x710,0x74A},{0x74D,0x7B1},{0x7CA,0x7F5},{0x7FA,0x7FA},{0x7FD,0x7FD},{0x800,0x82D}, + {0x840,0x85B},{0x860,0x86A},{0x8A0,0x8B4},{0x8B6,0x8C7},{0x8D3,0x8E1},{0x8E3,0x963}, + {0x971,0x983},{0x985,0x98C},{0x98F,0x990},{0x993,0x9A8},{0x9AA,0x9B0},{0x9B2,0x9B2}, + {0x9B6,0x9B9},{0x9BC,0x9C4},{0x9C7,0x9C8},{0x9CB,0x9CE},{0x9D7,0x9D7},{0x9DC,0x9DD}, + {0x9DF,0x9E3},{0x9F0,0x9F1},{0x9FC,0x9FC},{0x9FE,0x9FE},{0xA01,0xA03},{0xA05,0xA0A}, + {0xA0F,0xA10},{0xA13,0xA28},{0xA2A,0xA30},{0xA32,0xA33},{0xA35,0xA36},{0xA38,0xA39}, + {0xA3C,0xA3C},{0xA3E,0xA42},{0xA47,0xA48},{0xA4B,0xA4D},{0xA51,0xA51},{0xA59,0xA5C}, + {0xA5E,0xA5E},{0xA70,0xA75},{0xA81,0xA83},{0xA85,0xA8D},{0xA8F,0xA91},{0xA93,0xAA8}, + {0xAAA,0xAB0},{0xAB2,0xAB3},{0xAB5,0xAB9},{0xABC,0xAC5},{0xAC7,0xAC9},{0xACB,0xACD}, + {0xAD0,0xAD0},{0xAE0,0xAE3},{0xAF9,0xAFF},{0xB01,0xB03},{0xB05,0xB0C},{0xB0F,0xB10}, + {0xB13,0xB28},{0xB2A,0xB30},{0xB32,0xB33},{0xB35,0xB39},{0xB3C,0xB44},{0xB47,0xB48}, + {0xB4B,0xB4D},{0xB55,0xB57},{0xB5C,0xB5D},{0xB5F,0xB63},{0xB71,0xB71},{0xB82,0xB83}, + {0xB85,0xB8A},{0xB8E,0xB90},{0xB92,0xB95},{0xB99,0xB9A},{0xB9C,0xB9C},{0xB9E,0xB9F}, + {0xBA3,0xBA4},{0xBA8,0xBAA},{0xBAE,0xBB9},{0xBBE,0xBC2},{0xBC6,0xBC8},{0xBCA,0xBCD}, + {0xBD0,0xBD0},{0xBD7,0xBD7},{0xC00,0xC0C},{0xC0E,0xC10},{0xC12,0xC28},{0xC2A,0xC39}, + {0xC3D,0xC44},{0xC46,0xC48},{0xC4A,0xC4D},{0xC55,0xC56},{0xC58,0xC5A},{0xC60,0xC63}, + {0xC80,0xC83},{0xC85,0xC8C},{0xC8E,0xC90},{0xC92,0xCA8},{0xCAA,0xCB3},{0xCB5,0xCB9}, + {0xCBC,0xCC4},{0xCC6,0xCC8},{0xCCA,0xCCD},{0xCD5,0xCD6},{0xCDE,0xCDE},{0xCE0,0xCE3}, + {0xCF1,0xCF2},{0xD00,0xD0C},{0xD0E,0xD10},{0xD12,0xD44},{0xD46,0xD48},{0xD4A,0xD4E}, + {0xD54,0xD57},{0xD5F,0xD63},{0xD7A,0xD7F},{0xD81,0xD83},{0xD85,0xD96},{0xD9A,0xDB1}, + {0xDB3,0xDBB},{0xDBD,0xDBD},{0xDC0,0xDC6},{0xDCA,0xDCA},{0xDCF,0xDD4},{0xDD6,0xDD6}, + {0xDD8,0xDDF},{0xDF2,0xDF3},{0xE01,0xE3A},{0xE40,0xE4E},{0xE81,0xE82},{0xE84,0xE84}, + {0xE86,0xE8A},{0xE8C,0xEA3},{0xEA5,0xEA5},{0xEA7,0xEBD},{0xEC0,0xEC4},{0xEC6,0xEC6}, + {0xEC8,0xECD},{0xEDC,0xEDF},{0xF00,0xF00},{0xF18,0xF19},{0xF35,0xF35},{0xF37,0xF37}, + {0xF39,0xF39},{0xF3E,0xF47},{0xF49,0xF6C},{0xF71,0xF84},{0xF86,0xF97},{0xF99,0xFBC}, + {0xFC6,0xFC6},{0x1000,0x103F},{0x1050,0x108F},{0x109A,0x109D},{0x10FC,0x10FC},{0x1100,0x1248}, + {0x124A,0x124D},{0x1250,0x1256},{0x1258,0x1258},{0x125A,0x125D},{0x1260,0x1288},{0x128A,0x128D}, + {0x1290,0x12B0},{0x12B2,0x12B5},{0x12B8,0x12BE},{0x12C0,0x12C0},{0x12C2,0x12C5},{0x12C8,0x12D6}, + {0x12D8,0x1310},{0x1312,0x1315},{0x1318,0x135A},{0x135D,0x135F},{0x1380,0x138F},{0x1401,0x166C}, + {0x166F,0x167F},{0x1681,0x169A},{0x16A0,0x16EA},{0x16F1,0x16F8},{0x1700,0x170C},{0x170E,0x1714}, + {0x1720,0x1734},{0x1740,0x1753},{0x1760,0x176C},{0x176E,0x1770},{0x1772,0x1773},{0x1780,0x17D3}, + {0x17D7,0x17D7},{0x17DC,0x17DD},{0x180B,0x180D},{0x1820,0x1878},{0x1880,0x18AA},{0x18B0,0x18F5}, + {0x1900,0x191E},{0x1920,0x192B},{0x1930,0x193B},{0x1950,0x196D},{0x1970,0x1974},{0x1980,0x19AB}, + {0x19B0,0x19C9},{0x1A00,0x1A1B},{0x1A20,0x1A5E},{0x1A60,0x1A7C},{0x1A7F,0x1A7F},{0x1AA7,0x1AA7}, + {0x1AB0,0x1AC0},{0x1B00,0x1B4B},{0x1B6B,0x1B73},{0x1B80,0x1BAF},{0x1BBA,0x1BF3},{0x1C00,0x1C37}, + {0x1C4D,0x1C4F},{0x1C5A,0x1C7D},{0x1CD0,0x1CD2},{0x1CD4,0x1CFA},{0x1D2C,0x1D6A},{0x1D78,0x1D78}, + {0x1D9B,0x1DF9},{0x1DFB,0x1DFF},{0x2071,0x2071},{0x207F,0x207F},{0x2090,0x209C},{0x20D0,0x20F0}, + {0x2135,0x2138},{0x2C7C,0x2C7D},{0x2CEF,0x2CF1},{0x2D30,0x2D67},{0x2D6F,0x2D6F},{0x2D7F,0x2D96}, + {0x2DA0,0x2DA6},{0x2DA8,0x2DAE},{0x2DB0,0x2DB6},{0x2DB8,0x2DBE},{0x2DC0,0x2DC6},{0x2DC8,0x2DCE}, + {0x2DD0,0x2DD6},{0x2DD8,0x2DDE},{0x2DE0,0x2DFF},{0x2E2F,0x2E2F},{0x3005,0x3006},{0x302A,0x302F}, + {0x3031,0x3035},{0x303B,0x303C},{0x3041,0x3096},{0x3099,0x309A},{0x309D,0x309F},{0x30A1,0x30FA}, + {0x30FC,0x30FF},{0x3105,0x312F},{0x3131,0x318E},{0x31A0,0x31BF},{0x31F0,0x31FF},{0x3400,0x4DBF}, + {0x4E00,0x9FFC},{0xA000,0xA48C},{0xA4D0,0xA4FD},{0xA500,0xA60C},{0xA610,0xA61F},{0xA62A,0xA62B}, + {0xA66E,0xA672},{0xA674,0xA67D},{0xA67F,0xA67F},{0xA69C,0xA6E5},{0xA6F0,0xA6F1},{0xA717,0xA71F}, + {0xA770,0xA770},{0xA788,0xA788},{0xA78F,0xA78F},{0xA7F7,0xA7F9},{0xA7FB,0xA827},{0xA82C,0xA82C}, + {0xA840,0xA873},{0xA880,0xA8C5},{0xA8E0,0xA8F7},{0xA8FB,0xA8FB},{0xA8FD,0xA8FF},{0xA90A,0xA92D}, + {0xA930,0xA953},{0xA960,0xA97C},{0xA980,0xA9C0},{0xA9CF,0xA9CF},{0xA9E0,0xA9EF},{0xA9FA,0xA9FE}, + {0xAA00,0xAA36},{0xAA40,0xAA4D},{0xAA60,0xAA76},{0xAA7A,0xAAC2},{0xAADB,0xAADD},{0xAAE0,0xAAEF}, + {0xAAF2,0xAAF6},{0xAB01,0xAB06},{0xAB09,0xAB0E},{0xAB11,0xAB16},{0xAB20,0xAB26},{0xAB28,0xAB2E}, + {0xAB5C,0xAB5F},{0xAB69,0xAB69},{0xABC0,0xABEA},{0xABEC,0xABED},{0xAC00,0xD7A3},{0xD7B0,0xD7C6}, + {0xD7CB,0xD7FB},{0xF900,0xFA6D},{0xFA70,0xFAD9},{0xFB1D,0xFB28},{0xFB2A,0xFB36},{0xFB38,0xFB3C}, + {0xFB3E,0xFB3E},{0xFB40,0xFB41},{0xFB43,0xFB44},{0xFB46,0xFBB1},{0xFBD3,0xFD3D},{0xFD50,0xFD8F}, + {0xFD92,0xFDC7},{0xFDF0,0xFDFB},{0xFE00,0xFE0F},{0xFE20,0xFE2F},{0xFE70,0xFE74},{0xFE76,0xFEFC}, + {0xFF66,0xFFBE},{0xFFC2,0xFFC7},{0xFFCA,0xFFCF},{0xFFD2,0xFFD7},{0xFFDA,0xFFDC},{0x10000,0x1000B}, + {0x1000D,0x10026},{0x10028,0x1003A},{0x1003C,0x1003D},{0x1003F,0x1004D},{0x10050,0x1005D},{0x10080,0x100FA}, + {0x101FD,0x101FD},{0x10280,0x1029C},{0x102A0,0x102D0},{0x102E0,0x102E0},{0x10300,0x1031F},{0x1032D,0x10340}, + {0x10342,0x10349},{0x10350,0x1037A},{0x10380,0x1039D},{0x103A0,0x103C3},{0x103C8,0x103CF},{0x10450,0x1049D}, + {0x10500,0x10527},{0x10530,0x10563},{0x10600,0x10736},{0x10740,0x10755},{0x10760,0x10767},{0x10800,0x10805}, + {0x10808,0x10808},{0x1080A,0x10835},{0x10837,0x10838},{0x1083C,0x1083C},{0x1083F,0x10855},{0x10860,0x10876}, + {0x10880,0x1089E},{0x108E0,0x108F2},{0x108F4,0x108F5},{0x10900,0x10915},{0x10920,0x10939},{0x10980,0x109B7}, + {0x109BE,0x109BF},{0x10A00,0x10A03},{0x10A05,0x10A06},{0x10A0C,0x10A13},{0x10A15,0x10A17},{0x10A19,0x10A35}, + {0x10A38,0x10A3A},{0x10A3F,0x10A3F},{0x10A60,0x10A7C},{0x10A80,0x10A9C},{0x10AC0,0x10AC7},{0x10AC9,0x10AE6}, + {0x10B00,0x10B35},{0x10B40,0x10B55},{0x10B60,0x10B72},{0x10B80,0x10B91},{0x10C00,0x10C48},{0x10D00,0x10D27}, + {0x10E80,0x10EA9},{0x10EAB,0x10EAC},{0x10EB0,0x10EB1},{0x10F00,0x10F1C},{0x10F27,0x10F27},{0x10F30,0x10F50}, + {0x10FB0,0x10FC4},{0x10FE0,0x10FF6},{0x11000,0x11046},{0x1107F,0x110BA},{0x110D0,0x110E8},{0x11100,0x11134}, + {0x11144,0x11147},{0x11150,0x11173},{0x11176,0x11176},{0x11180,0x111C4},{0x111C9,0x111CC},{0x111CE,0x111CF}, + {0x111DA,0x111DA},{0x111DC,0x111DC},{0x11200,0x11211},{0x11213,0x11237},{0x1123E,0x1123E},{0x11280,0x11286}, + {0x11288,0x11288},{0x1128A,0x1128D},{0x1128F,0x1129D},{0x1129F,0x112A8},{0x112B0,0x112EA},{0x11300,0x11303}, + {0x11305,0x1130C},{0x1130F,0x11310},{0x11313,0x11328},{0x1132A,0x11330},{0x11332,0x11333},{0x11335,0x11339}, + {0x1133B,0x11344},{0x11347,0x11348},{0x1134B,0x1134D},{0x11350,0x11350},{0x11357,0x11357},{0x1135D,0x11363}, + {0x11366,0x1136C},{0x11370,0x11374},{0x11400,0x1144A},{0x1145E,0x11461},{0x11480,0x114C5},{0x114C7,0x114C7}, + {0x11580,0x115B5},{0x115B8,0x115C0},{0x115D8,0x115DD},{0x11600,0x11640},{0x11644,0x11644},{0x11680,0x116B8}, + {0x11700,0x1171A},{0x1171D,0x1172B},{0x11800,0x1183A},{0x118FF,0x11906},{0x11909,0x11909},{0x1190C,0x11913}, + {0x11915,0x11916},{0x11918,0x11935},{0x11937,0x11938},{0x1193B,0x11943},{0x119A0,0x119A7},{0x119AA,0x119D7}, + {0x119DA,0x119E1},{0x119E3,0x119E4},{0x11A00,0x11A3E},{0x11A47,0x11A47},{0x11A50,0x11A99},{0x11A9D,0x11A9D}, + {0x11AC0,0x11AF8},{0x11C00,0x11C08},{0x11C0A,0x11C36},{0x11C38,0x11C40},{0x11C72,0x11C8F},{0x11C92,0x11CA7}, + {0x11CA9,0x11CB6},{0x11D00,0x11D06},{0x11D08,0x11D09},{0x11D0B,0x11D36},{0x11D3A,0x11D3A},{0x11D3C,0x11D3D}, + {0x11D3F,0x11D47},{0x11D60,0x11D65},{0x11D67,0x11D68},{0x11D6A,0x11D8E},{0x11D90,0x11D91},{0x11D93,0x11D98}, + {0x11EE0,0x11EF6},{0x11FB0,0x11FB0},{0x12000,0x12399},{0x12480,0x12543},{0x13000,0x1342E},{0x14400,0x14646}, + {0x16800,0x16A38},{0x16A40,0x16A5E},{0x16AD0,0x16AED},{0x16AF0,0x16AF4},{0x16B00,0x16B36},{0x16B40,0x16B43}, + {0x16B63,0x16B77},{0x16B7D,0x16B8F},{0x16F00,0x16F4A},{0x16F4F,0x16F87},{0x16F8F,0x16F9F},{0x16FE0,0x16FE1}, + {0x16FE3,0x16FE4},{0x16FF0,0x16FF1},{0x17000,0x187F7},{0x18800,0x18CD5},{0x18D00,0x18D08},{0x1B000,0x1B11E}, + {0x1B150,0x1B152},{0x1B164,0x1B167},{0x1B170,0x1B2FB},{0x1BC00,0x1BC6A},{0x1BC70,0x1BC7C},{0x1BC80,0x1BC88}, + {0x1BC90,0x1BC99},{0x1BC9D,0x1BC9E},{0x1D165,0x1D169},{0x1D16D,0x1D172},{0x1D17B,0x1D182},{0x1D185,0x1D18B}, + {0x1D1AA,0x1D1AD},{0x1D242,0x1D244},{0x1DA00,0x1DA36},{0x1DA3B,0x1DA6C},{0x1DA75,0x1DA75},{0x1DA84,0x1DA84}, + {0x1DA9B,0x1DA9F},{0x1DAA1,0x1DAAF},{0x1E000,0x1E006},{0x1E008,0x1E018},{0x1E01B,0x1E021},{0x1E023,0x1E024}, + {0x1E026,0x1E02A},{0x1E100,0x1E12C},{0x1E130,0x1E13D},{0x1E14E,0x1E14E},{0x1E2C0,0x1E2EF},{0x1E800,0x1E8C4}, + {0x1E8D0,0x1E8D6},{0x1E944,0x1E94B},{0x1EE00,0x1EE03},{0x1EE05,0x1EE1F},{0x1EE21,0x1EE22},{0x1EE24,0x1EE24}, + {0x1EE27,0x1EE27},{0x1EE29,0x1EE32},{0x1EE34,0x1EE37},{0x1EE39,0x1EE39},{0x1EE3B,0x1EE3B},{0x1EE42,0x1EE42}, + {0x1EE47,0x1EE47},{0x1EE49,0x1EE49},{0x1EE4B,0x1EE4B},{0x1EE4D,0x1EE4F},{0x1EE51,0x1EE52},{0x1EE54,0x1EE54}, + {0x1EE57,0x1EE57},{0x1EE59,0x1EE59},{0x1EE5B,0x1EE5B},{0x1EE5D,0x1EE5D},{0x1EE5F,0x1EE5F},{0x1EE61,0x1EE62}, + {0x1EE64,0x1EE64},{0x1EE67,0x1EE6A},{0x1EE6C,0x1EE72},{0x1EE74,0x1EE77},{0x1EE79,0x1EE7C},{0x1EE7E,0x1EE7E}, + {0x1EE80,0x1EE89},{0x1EE8B,0x1EE9B},{0x1EEA1,0x1EEA3},{0x1EEA5,0x1EEA9},{0x1EEAB,0x1EEBB},{0x20000,0x2A6DD}, + {0x2A700,0x2B734},{0x2B740,0x2B81D},{0x2B820,0x2CEA1},{0x2CEB0,0x2EBE0},{0x2F800,0x2FA1D},{0x30000,0x3134A}, + {0xE0100,0xE01EF}, +}; +static const int uni_X_n = 595; + +static inline int is_U(uint32_t c){ return uni_in(uni_U,uni_U_n,c); } +static inline int is_X(uint32_t c){ return uni_in(uni_X,uni_X_n,c); } +#endif From 9e2d41567c5e0902da1ef5a3daaff9a957d83997 Mon Sep 17 00:00:00 2001 From: Nicholas Beerbower Date: Thu, 16 Jul 2026 15:17:21 -0400 Subject: [PATCH 02/71] tests: o200k tokenizer coverage, no model download MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit tests/tok_o200k_tiny.json is a synthetic byte-level BPE (274 vocab, a few KB) whose Split regex is the o200k pattern; expected ids in tok_o200k_cases.txt were generated by HF tokenizers on that same file. test_tok_o200k (in TEST_BINS) scores 40/40 encode + 40/40 round-trip: case-transition splits, contractions, digit groups, the [\r\n/]* tail, whitespace branches, CJK/Greek/Cyrillic, added-token atomicity. The cl100k path is untouched by construction — dispatch requires \p{Lu} in the tokenizer's own Split pattern, which cl100k lacks — and stays covered by the GLM oracle (verified on this branch: 32/32). Co-Authored-By: Claude Fable 5 --- c/Makefile | 5 ++- c/tests/test_tok_o200k.c | 61 +++++++++++++++++++++++++++++++++++++ c/tests/tok_o200k_cases.txt | 40 ++++++++++++++++++++++++ c/tests/tok_o200k_tiny.json | 1 + 4 files changed, 106 insertions(+), 1 deletion(-) create mode 100644 c/tests/test_tok_o200k.c create mode 100644 c/tests/tok_o200k_cases.txt create mode 100644 c/tests/tok_o200k_tiny.json diff --git a/c/Makefile b/c/Makefile index 27902cb..9887c74 100644 --- a/c/Makefile +++ b/c/Makefile @@ -166,7 +166,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_tok_o200k$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -293,6 +293,9 @@ iobench$(EXE): iobench.c compat.h tests/test_json$(EXE): tests/test_json.c json.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +tests/test_tok_o200k$(EXE): tests/test_tok_o200k.c tok.h tok_unicode.h tok_unicode_o200k.h json.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + tests/test_st$(EXE): tests/test_st.c st.h json.h compat.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) diff --git a/c/tests/test_tok_o200k.c b/c/tests/test_tok_o200k.c new file mode 100644 index 0000000..b86bda9 --- /dev/null +++ b/c/tests/test_tok_o200k.c @@ -0,0 +1,61 @@ +/* o200k pre-tokenizer validation against HF-tokenizers-generated expectations. + * Self-contained for the test-c harness: loads tests/tok_o200k_tiny.json (a + * synthetic byte-level BPE whose Split regex is the o200k pattern — a few KB, + * no model download) and scores tests/tok_o200k_cases.txt, whose expected ids + * were produced by HF `tokenizers` on the same file. Guards the case-aware + * letter matcher, contractions, digit groups, the [\r\n/]* punctuation tail, + * whitespace branches, and added-token atomicity; round-trips every case. + * The cl100k path is untouched by construction (dispatch requires \p{Lu} in + * the tokenizer's own Split pattern) and stays covered by the GLM oracle. */ +#define _GNU_SOURCE +#include "../tok.h" + +int main(void) { + Tok T; + tok_load(&T, "tests/tok_o200k_tiny.json"); + if (!T.o200k) { fprintf(stderr, "test_tok_o200k: o200k pattern not detected\n"); return 1; } + FILE *f = fopen("tests/tok_o200k_cases.txt", "rb"); + if (!f) { perror("tests/tok_o200k_cases.txt"); return 1; } + char *line = NULL; size_t cap = 0; ssize_t nr; + int pass = 0, tot = 0, dpass = 0; + while ((nr = getline(&line, &cap, f)) >= 0) { + if (nr > 0 && line[nr-1] == '\n') line[--nr] = 0; + if (nr == 0) continue; + char *tab = strchr(line, '\t'); if (!tab) continue; + *tab = 0; + const char *text = line, *idstr = tab + 1; + char tbuf[4096]; int tn = 0; + for (const char *q = text; *q && tn < 4095; q++) { + if (q[0]=='\\' && q[1]=='n') { tbuf[tn++]='\n'; q++; } + else if (q[0]=='\\' && q[1]=='t') { tbuf[tn++]='\t'; q++; } + else if (q[0]=='\\' && q[1]=='r') { tbuf[tn++]='\r'; q++; } + else if (q[0]=='\\' && q[1]=='\\') { tbuf[tn++]='\\'; q++; } + else tbuf[tn++] = *q; + } + tbuf[tn] = 0; + int exp[512], ne = 0; + for (const char *q = idstr; *q; ) { + while (*q == ',' || *q == ' ') q++; + if (!*q) break; + exp[ne++] = atoi(q); + while (*q && *q != ',') q++; + } + int got[512]; int ng = tok_encode(&T, tbuf, tn, got, 512); + int ok = (ng == ne); + for (int i = 0; i < ng && ok; i++) ok = (got[i] == exp[i]); + tot++; if (ok) pass++; + char dec[8192]; int dn = tok_decode(&T, got, ng, dec, 8191); + int drt = (dn == tn) && !memcmp(dec, tbuf, tn); + if (drt) dpass++; + if (!ok || !drt) { + fprintf(stderr, "MISMATCH text=%s\n exp(%d):", text, ne); + for (int i = 0; i < ne; i++) fprintf(stderr, " %d", exp[i]); + fprintf(stderr, "\n got(%d):", ng); + for (int i = 0; i < ng; i++) fprintf(stderr, " %d", got[i]); + fprintf(stderr, "\n decode_ok=%d\n", drt); + } + } + fclose(f); + printf("test_tok_o200k: ENCODE %d/%d DECODE %d/%d\n", pass, tot, dpass, tot); + return (pass == tot && dpass == tot) ? 0 : 2; +} diff --git a/c/tests/tok_o200k_cases.txt b/c/tests/tok_o200k_cases.txt new file mode 100644 index 0000000..38e37fa --- /dev/null +++ b/c/tests/tok_o200k_cases.txt @@ -0,0 +1,40 @@ +hello world 259,32,119,111,114,263 +HelloWorld 72,101,257,111,262 +XMLHttpRequest 88,77,76,72,116,116,112,82,101,113,117,101,115,116 +helloWORLDhello 259,87,79,82,76,68,259 +dog's 100,111,103,270 +DON'T 68,79,78,39,84 +don't 100,111,110,39,116 +I'll've 73,39,257,39,118,101 +O'Brien 79,39,66,114,105,101,110 +the theatre 265,32,116,256,97,116,114,101 +12345 268,52,53 +a1b22c333d4444 97,49,98,50,50,99,51,51,51,100,52,52,52,52 +3.14 51,46,49,52 +http://x.com/a/b 104,116,116,112,58,47,47,120,46,99,111,109,47,97,47,98 +path/to/file 112,97,264,47,116,111,47,102,105,108,101 +a//b///c 97,47,47,98,47,47,47,99 +!!\r\n//x 33,33,13,10,47,47,120 +one\ntwo\r\nthree 111,110,101,10,116,119,111,13,10,264,114,101,101 + \n x 32,32,10,32,32,120 + 32,32,32 +a b c 97,32,32,98,32,32,32,99 +tab\there 116,269,9,256,114,101 +Café 67,97,102,101,204,129 +naiveBayes 110,97,105,118,101,66,97,121,101,115 +Éclair 195,137,99,108,97,105,114 +北京大学 229,140,151,228,186,172,229,164,167,229,173,166 +ΑΒαβ 206,145,206,146,206,177,206,178 +Иван 208,152,208,178,208,176,208,189 +ẞßscharf 225,186,158,195,159,115,99,104,97,114,102 +i̇stanbul 105,204,135,115,116,97,110,98,117,108 +hello<|endoftext|>world 259,274,119,111,114,263 +<|message_user|>hi 275,104,105 +mixedCASEand123 109,105,120,101,100,67,65,83,69,97,110,100,268 +'s 270 + 's 32,39,115 +A 65 +aB 97,66 +Ab 65,98 +AB 65,66 +ab 269 diff --git a/c/tests/tok_o200k_tiny.json b/c/tests/tok_o200k_tiny.json new file mode 100644 index 0000000..829301e --- /dev/null +++ b/c/tests/tok_o200k_tiny.json @@ -0,0 +1 @@ +{"version": "1.0", "truncation": null, "padding": null, "added_tokens": [{"id": 274, "content": "<|endoftext|>", "single_word": false, "lstrip": false, "rstrip": false, "normalized": false, "special": true}, {"id": 275, "content": "<|message_user|>", "single_word": false, "lstrip": false, "rstrip": false, "normalized": false, "special": true}], "normalizer": null, "pre_tokenizer": {"type": "Sequence", "pretokenizers": [{"type": "Split", "pattern": {"Regex": "[^\\r\\n\\p{L}\\p{N}]?[\\p{Lu}\\p{Lt}\\p{Lm}\\p{Lo}\\p{M}]*[\\p{Ll}\\p{Lm}\\p{Lo}\\p{M}]+(?i:'s|'t|'re|'ve|'m|'ll|'d)?|[^\\r\\n\\p{L}\\p{N}]?[\\p{Lu}\\p{Lt}\\p{Lm}\\p{Lo}\\p{M}]+[\\p{Ll}\\p{Lm}\\p{Lo}\\p{M}]*(?i:'s|'t|'re|'ve|'m|'ll|'d)?|\\p{N}{1,3}| ?[^\\s\\p{L}\\p{N}]+[\\r\\n/]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+"}, "behavior": "Isolated", "invert": false}, {"type": "ByteLevel", "add_prefix_space": false, "trim_offsets": true, "use_regex": false}]}, "post_processor": null, "decoder": {"type": "ByteLevel", "add_prefix_space": true, "trim_offsets": true, "use_regex": true}, "model": {"type": "BPE", "dropout": null, "unk_token": null, "continuing_subword_prefix": null, "end_of_word_suffix": null, "fuse_unk": false, "byte_fallback": false, "ignore_merges": true, "vocab": {"Ā": 0, "ā": 1, "Ă": 2, "ă": 3, "Ą": 4, "ą": 5, "Ć": 6, "ć": 7, "Ĉ": 8, "ĉ": 9, "Ċ": 10, "ċ": 11, "Č": 12, "č": 13, "Ď": 14, "ď": 15, "Đ": 16, "đ": 17, "Ē": 18, "ē": 19, "Ĕ": 20, "ĕ": 21, "Ė": 22, "ė": 23, "Ę": 24, "ę": 25, "Ě": 26, "ě": 27, "Ĝ": 28, "ĝ": 29, "Ğ": 30, "ğ": 31, "Ġ": 32, "!": 33, "\"": 34, "#": 35, "$": 36, "%": 37, "&": 38, "'": 39, "(": 40, ")": 41, "*": 42, "+": 43, ",": 44, "-": 45, ".": 46, "/": 47, "0": 48, "1": 49, "2": 50, "3": 51, "4": 52, "5": 53, "6": 54, "7": 55, "8": 56, "9": 57, ":": 58, ";": 59, "<": 60, "=": 61, ">": 62, "?": 63, "@": 64, "A": 65, "B": 66, "C": 67, "D": 68, "E": 69, "F": 70, "G": 71, "H": 72, "I": 73, "J": 74, "K": 75, "L": 76, "M": 77, "N": 78, "O": 79, "P": 80, "Q": 81, "R": 82, "S": 83, "T": 84, "U": 85, "V": 86, "W": 87, "X": 88, "Y": 89, "Z": 90, "[": 91, "\\": 92, "]": 93, "^": 94, "_": 95, "`": 96, "a": 97, "b": 98, "c": 99, "d": 100, "e": 101, "f": 102, "g": 103, "h": 104, "i": 105, "j": 106, "k": 107, "l": 108, "m": 109, "n": 110, "o": 111, "p": 112, "q": 113, "r": 114, "s": 115, "t": 116, "u": 117, "v": 118, "w": 119, "x": 120, "y": 121, "z": 122, "{": 123, "|": 124, "}": 125, "~": 126, "ġ": 127, "Ģ": 128, "ģ": 129, "Ĥ": 130, "ĥ": 131, "Ħ": 132, "ħ": 133, "Ĩ": 134, "ĩ": 135, "Ī": 136, "ī": 137, "Ĭ": 138, "ĭ": 139, "Į": 140, "į": 141, "İ": 142, "ı": 143, "IJ": 144, "ij": 145, "Ĵ": 146, "ĵ": 147, "Ķ": 148, "ķ": 149, "ĸ": 150, "Ĺ": 151, "ĺ": 152, "Ļ": 153, "ļ": 154, "Ľ": 155, "ľ": 156, "Ŀ": 157, "ŀ": 158, "Ł": 159, "ł": 160, "¡": 161, "¢": 162, "£": 163, "¤": 164, "¥": 165, "¦": 166, "§": 167, "¨": 168, "©": 169, "ª": 170, "«": 171, "¬": 172, "Ń": 173, "®": 174, "¯": 175, "°": 176, "±": 177, "²": 178, "³": 179, "´": 180, "µ": 181, "¶": 182, "·": 183, "¸": 184, "¹": 185, "º": 186, "»": 187, "¼": 188, "½": 189, "¾": 190, "¿": 191, "À": 192, "Á": 193, "Â": 194, "Ã": 195, "Ä": 196, "Å": 197, "Æ": 198, "Ç": 199, "È": 200, "É": 201, "Ê": 202, "Ë": 203, "Ì": 204, "Í": 205, "Î": 206, "Ï": 207, "Ð": 208, "Ñ": 209, "Ò": 210, "Ó": 211, "Ô": 212, "Õ": 213, "Ö": 214, "×": 215, "Ø": 216, "Ù": 217, "Ú": 218, "Û": 219, "Ü": 220, "Ý": 221, "Þ": 222, "ß": 223, "à": 224, "á": 225, "â": 226, "ã": 227, "ä": 228, "å": 229, "æ": 230, "ç": 231, "è": 232, "é": 233, "ê": 234, "ë": 235, "ì": 236, "í": 237, "î": 238, "ï": 239, "ð": 240, "ñ": 241, "ò": 242, "ó": 243, "ô": 244, "õ": 245, "ö": 246, "÷": 247, "ø": 248, "ù": 249, "ú": 250, "û": 251, "ü": 252, "ý": 253, "þ": 254, "ÿ": 255, "he": 256, "ll": 257, "hell": 258, "hello": 259, "Wo": 260, "Wor": 261, "World": 262, "ld": 263, "th": 264, "the": 265, "Ġthe": 266, "12": 267, "123": 268, "ab": 269, "'s": 270, "Ġa": 271, "./": 272, "ĊĊ": 273}, "merges": [["h", "e"], ["l", "l"], ["he", "ll"], ["hell", "o"], ["W", "o"], ["Wo", "r"], ["Wor", "ld"], ["l", "d"], ["t", "h"], ["th", "e"], ["Ġ", "the"], ["1", "2"], ["12", "3"], ["a", "b"], ["'", "s"], ["Ġ", "a"], [".", "/"], ["Ċ", "Ċ"]]}} \ No newline at end of file From 5ad4d540ab7427e6ca134ce88fe93e3d9f869456 Mon Sep 17 00:00:00 2001 From: Nicholas Beerbower Date: Thu, 16 Jul 2026 16:33:22 -0400 Subject: [PATCH 03/71] test_tok_o200k: fgets instead of getline for the windows job MinGW's UCRT has no getline; fixed-buffer fgets with CRLF trimming reads the same case file everywhere. Co-Authored-By: Claude Fable 5 --- c/tests/test_tok_o200k.c | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/c/tests/test_tok_o200k.c b/c/tests/test_tok_o200k.c index b86bda9..fe62890 100644 --- a/c/tests/test_tok_o200k.c +++ b/c/tests/test_tok_o200k.c @@ -16,10 +16,13 @@ int main(void) { if (!T.o200k) { fprintf(stderr, "test_tok_o200k: o200k pattern not detected\n"); return 1; } FILE *f = fopen("tests/tok_o200k_cases.txt", "rb"); if (!f) { perror("tests/tok_o200k_cases.txt"); return 1; } - char *line = NULL; size_t cap = 0; ssize_t nr; + /* fgets, not getline: MinGW's UCRT lacks getline and this must run on + * the windows job. Case lines are short; 8 KB is generous. */ + char line[8192]; int pass = 0, tot = 0, dpass = 0; - while ((nr = getline(&line, &cap, f)) >= 0) { - if (nr > 0 && line[nr-1] == '\n') line[--nr] = 0; + while (fgets(line, sizeof(line), f)) { + size_t nr = strlen(line); + while (nr > 0 && (line[nr-1] == '\n' || line[nr-1] == '\r')) line[--nr] = 0; if (nr == 0) continue; char *tab = strchr(line, '\t'); if (!tab) continue; *tab = 0; From 288edd71903a222b79472b327640cfae844e1010 Mon Sep 17 00:00:00 2001 From: noobdev-ph Date: Fri, 17 Jul 2026 12:39:59 +0800 Subject: [PATCH 04/71] fix: GPU backend failure-path hardening + tests MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Three vendor-neutral fixes to backend_cuda.cu, each with test coverage: 1. Upload check-order: a cached device tensor is now usable when the caller's host pointers are stale or NULL. CUDA_RELEASE_HOST slots null their host pointers after upload; the current engine reaches those tensors through direct handles (coli_cuda_expert_mlp etc.), but any caller going through coli_cuda_matmul/tensor_upload with a cached tensor — as matmul_qt does for QT tensors — hits the !weights check before the cached-tensor branch and fails spuriously. This hardens the API contract rather than fixing a measured regression; the contract is pinned by a 64x sustained-reuse test. 2. Sticky runtime error (real bug, test-caught): a failed allocation left the last-error state set, so the NEXT healthy launch's cudaGetLastError() check reported 'out of memory' and disabled a perfectly good tensor. cuda_ok() now consumes the error on the failure path; regression-covered. 3. COLI_GPU_FAIL_AFTER=N test hook: every GPU compute entry point (19 total: matmul, expert mlp/group, shared mlp, attention ops, pipe ops) reports failure after N successful calls, so the engine's CPU fallbacks and expert_host_ensure rematerialization can be exercised end-to-end without real hardware faults. Unset = zero effect; uploads/queries are never gated. Validated on GLM-5.2: total failure (N=0) completes coherently with every released expert rematerialized. Tests (run via make cuda-test on any CUDA GPU; vendor-neutral source): 64x sustained matmul reuse after host pointers are freed; upload from a scribbled-and-freed temporary; five graceful upload-failure cases with stats-integrity assertions; healthy-launch-after-failed-alloc (the sticky-error regression); fault-hook on/off restore. Verified on AMD RX 9070 XT via the companion HIP PR's compat header (same test source); a make cuda-test run on NVIDIA hardware would complete the matrix. --- c/backend_cuda.cu | 45 +++++++++++++++++++++++++++++++---- c/tests/test_backend_cuda.cu | 46 ++++++++++++++++++++++++++++++++++++ 2 files changed, 87 insertions(+), 4 deletions(-) diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu index 9ce142f..69a224f 100644 --- a/c/backend_cuda.cu +++ b/c/backend_cuda.cu @@ -45,6 +45,8 @@ static std::mutex g_group_stats_mu; static int cuda_ok(cudaError_t err, const char *what) { if (err == cudaSuccess) return 1; std::fprintf(stderr, "[CUDA] %s: %s\n", what, cudaGetErrorString(err)); + (void)cudaGetLastError(); /* consume the sticky error: a failed call must + not poison the next launch's error check */ return 0; } @@ -453,14 +455,20 @@ extern "C" void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64 extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor, const void *weights, const float *scales, int fmt, int I, int O, int device) { - DeviceContext *ctx = find_ctx(device); - if (!tensor || !weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0; - size_t rb = row_bytes(fmt, I); - if (!rb || (fmt && !scales)) return 0; + if (!tensor) return 0; if (*tensor) { + /* Cached device copy: usable even when the caller's host pointers are + * gone. CUDA_RELEASE_HOST slots null their host pointers after upload, + * and with the old order (!weights checked first) every later matmul + * on such a slot failed here — the GPU tier silently never computed + * for host-released slab experts. */ ColiCudaTensor *t = *tensor; return t->fmt == fmt && t->I == I && t->O == O && t->device == device; } + DeviceContext *ctx = find_ctx(device); + if (!weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0; + size_t rb = row_bytes(fmt, I); + if (!rb || (fmt && !scales)) return 0; ColiCudaTensor *t = static_cast(std::calloc(1, sizeof(*t))); if (!t) return 0; t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O; @@ -502,10 +510,21 @@ extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor, (size_t)tensor->O*sizeof(float),cudaMemcpyHostToDevice),"scale refresh"); } +/* Test hook: COLI_GPU_FAIL_AFTER=N makes every GPU COMPUTE entry point report + * failure after N successful calls (N=0: every call fails), exercising the + * engine's CPU fallbacks and host-rematerialization end-to-end without real + * hardware faults. Uploads/queries are not gated. Unset: no effect. */ +static long g_gpu_calls; +static int fault_injected(void) { + const char *fa = std::getenv("COLI_GPU_FAIL_AFTER"); + return fa && g_gpu_calls++ >= std::atol(fa); +} + extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor, float *y, const float *x, const void *weights, const float *scales, int fmt, int S, int I, int O, int device) { + if (fault_injected()) return 0; if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device)) return 0; ColiCudaTensor *t = *tensor; DeviceContext *ctx = find_ctx(t->device); @@ -524,6 +543,7 @@ extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor, extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up, ColiCudaTensor *down, float *y, const float *x, int S) { + if (fault_injected()) return 0; if (!gate || !up || !down || !x || !y || S < 1 || gate->device != up->device || gate->device != down->device || gate->I != up->I || gate->O != up->O || @@ -552,6 +572,7 @@ extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up, extern "C" int coli_cuda_shared_mlp_w4a16(ColiCudaTensor *gate,ColiCudaTensor *up, ColiCudaTensor *down,float *y,const float *x,int S){ + if (fault_injected()) return 0; if(!gate||!up||!down||!x||!y||S<1||gate->fmt!=2||up->fmt!=2||down->fmt!=2|| gate->device!=up->device||gate->device!=down->device||gate->I!=up->I|| gate->O!=up->O||down->I!=gate->O||down->O!=gate->I)return 0; @@ -583,6 +604,7 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates, ColiCudaTensor *const *downs, const int *rows, int count, float *y, const float *x) { + if (fault_injected()) return 0; if (!gates || !ups || !downs || !rows || !x || !y || count < 1) return 0; ColiCudaTensor *first=gates[0]; if (!first) return 0; @@ -708,6 +730,7 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates, extern "C" int coli_cuda_attention_absorb(ColiCudaTensor *w,float *ctx,const float *q, const float *latent,const float *rope,int H,int Q, int R,int V,int K,int T,float scale){ + if (fault_injected()) return 0; if(!w||!ctx||!q||!latent||!rope||H<1||Q<1||R<1||V<1||K<1||K>512||T<1||T>4096|| w->I!=K||w->O!=H*(Q+V))return 0; DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0; @@ -761,12 +784,14 @@ static int attention_absorb_batch_run(ColiCudaTensor *w,ColiCudaTensor *proj,flo extern "C" int coli_cuda_attention_absorb_batch(ColiCudaTensor *w,float *ctx,const float *q, const float *latent,const float *rope,int S,int H,int Q,int R,int V,int K,int T, float scale){ + if (fault_injected()) return 0; return attention_absorb_batch_run(w,nullptr,ctx,q,latent,rope,S,H,Q,R,V,K,T,scale); } extern "C" int coli_cuda_attention_project_batch(ColiCudaTensor *w,ColiCudaTensor *proj, float *out,const float *q,const float *latent,const float *rope,int S,int H,int Q, int R,int V,int K,int T,float scale){ + if (fault_injected()) return 0; return attention_absorb_batch_run(w,proj,out,q,latent,rope,S,H,Q,R,V,K,T,scale); } @@ -870,6 +895,7 @@ extern "C" int coli_cuda_pipe_download(int device,const void *src,void *dst,size } extern "C" int coli_cuda_pipe_rmsnorm(int device,float *y_dev,const float *x_dev, const float *w_dev,int S,int D,float eps){ + if (fault_injected()) return 0; DeviceContext *ctx=find_ctx(device); if(S<1||D<1||!select_ctx(ctx)) return 0; pipe_rmsnorm_rows<<>>(y_dev,x_dev,w_dev,D,eps,D,D); @@ -878,6 +904,7 @@ extern "C" int coli_cuda_pipe_rmsnorm(int device,float *y_dev,const float *x_dev extern "C" int coli_cuda_pipe_rmsnorm_s(int device,float *y_dev,const float *x_dev, const float *w_dev,int S,int D,float eps, int xstride,int ystride){ + if (fault_injected()) return 0; DeviceContext *ctx=find_ctx(device); if(S<1||D<1||xstride>>(y_dev,x_dev,w_dev,D,eps,xstride,ystride); @@ -886,6 +913,7 @@ extern "C" int coli_cuda_pipe_rmsnorm_s(int device,float *y_dev,const float *x_d extern "C" int coli_cuda_pipe_rope(int device,float *v_dev,const int *pos_dev, int rows,int stride,int offset,int R,int heads, float theta){ + if (fault_injected()) return 0; DeviceContext *ctx=find_ctx(device); if(rows<1||R<2||R>256||heads<1||!select_ctx(ctx)) return 0; pipe_rope_rows<<>>(v_dev,pos_dev,0,stride,offset,R,heads,theta); @@ -893,6 +921,7 @@ extern "C" int coli_cuda_pipe_rope(int device,float *v_dev,const int *pos_dev, } extern "C" int coli_cuda_pipe_rope_base(int device,float *v_dev,int pos_base,int rows, int stride,int offset,int R,int heads,float theta){ + if (fault_injected()) return 0; DeviceContext *ctx=find_ctx(device); if(rows<1||R<2||R>256||heads<1||!select_ctx(ctx)) return 0; pipe_rope_rows<<>>(v_dev,NULL,pos_base,stride,offset,R,heads,theta); @@ -910,6 +939,7 @@ extern "C" int coli_cuda_pipe_copy2d(int device,float *dst,int dpitch,const floa extern "C" int coli_cuda_attention_project_batch_dev(ColiCudaTensor *w,ColiCudaTensor *proj, float *out,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale){ + if (fault_injected()) return 0; if(!w||!proj||!out||!q_dev||!latent_dev||!rope_dev||S<1||H<1||Q<1||R<1||V<1|| K<1||K>512||T8192||w->I!=K||w->O!=H*(Q+V)|| proj->device!=w->device||proj->I!=H*V)return 0; @@ -931,17 +961,20 @@ extern "C" int coli_cuda_attention_project_batch_dev(ColiCudaTensor *w,ColiCudaT } extern "C" int coli_cuda_pipe_silu_mul(int device,float *gate_dev,const float *up_dev, size_t n){ + if (fault_injected()) return 0; DeviceContext *ctx=find_ctx(device); if(!n||!select_ctx(ctx)) return 0; silu_mul<<<(unsigned)((n+255)/256),256>>>(gate_dev,up_dev,n); return cuda_ok(cudaGetLastError(),"pipe silu mul"); } extern "C" int coli_cuda_pipe_add(int device,float *x_dev,const float *t_dev,size_t n){ + if (fault_injected()) return 0; DeviceContext *ctx=find_ctx(device); if(!n||!select_ctx(ctx)) return 0; pipe_add_n<<<(unsigned)((n+255)/256),256>>>(x_dev,t_dev,n); return cuda_ok(cudaGetLastError(),"pipe add"); } extern "C" int coli_cuda_pipe_rows_add(int device,float *x_dev,const float *partial_dev, const int *rows_dev,int nrows,int D){ + if (fault_injected()) return 0; DeviceContext *ctx=find_ctx(device); if(nrows<1||D<1||!select_ctx(ctx)) return 0; pipe_rows_add<<>>(x_dev,partial_dev,rows_dev,D); return cuda_ok(cudaGetLastError(),"pipe rows add"); @@ -950,6 +983,7 @@ extern "C" int coli_cuda_pipe_rows_add(int device,float *x_dev,const float *part * coli_cuda_matmul, zero host transfers. */ extern "C" int coli_cuda_pipe_gemm(ColiCudaTensor *t,float *y_dev,const float *x_dev, int S){ + if (fault_injected()) return 0; if(!t||S<1) return 0; DeviceContext *ctx=find_ctx(t->device); if(!select_ctx(ctx)) return 0; dim3 grid((unsigned)t->O,(unsigned)S); @@ -969,6 +1003,7 @@ extern "C" int coli_cuda_pipe_peer_copy(int dst_dev,float *dst,int src_dev, extern "C" int coli_cuda_attention_project_batch_dev_out(ColiCudaTensor *w,ColiCudaTensor *proj, float *out_dev,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale){ + if (fault_injected()) return 0; if(!w||!proj||!out_dev||!q_dev||!latent_dev||!rope_dev||S<1||H<1||Q<1||R<1||V<1|| K<1||K>512||T8192||w->I!=K||w->O!=H*(Q+V)|| proj->device!=w->device||proj->I!=H*V)return 0; @@ -990,6 +1025,7 @@ extern "C" int coli_cuda_attention_project_batch_dev_out(ColiCudaTensor *w,ColiC extern "C" int coli_cuda_attention_absorb_batch_dev(ColiCudaTensor *w,float *ctx_dev, const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale){ + if (fault_injected()) return 0; if(!w||!ctx_dev||!q_dev||!latent_dev||!rope_dev||S<1||H<1||Q<1||R<1||V<1|| K<1||K>512||T8192||w->I!=K||w->O!=H*(Q+V))return 0; DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0; @@ -1004,6 +1040,7 @@ extern "C" int coli_cuda_attention_absorb_batch_dev(ColiCudaTensor *w,float *ctx extern "C" int coli_cuda_attention_absorb_kvdev(ColiCudaTensor *w,float *ctx,const float *q, const float *latent_dev,const float *rope_dev,int H,int Q,int R,int V,int K,int T, float scale){ + if (fault_injected()) return 0; if(!w||!ctx||!q||!latent_dev||!rope_dev||H<1||Q<1||R<1||V<1||K<1||K>512||T<1||T>8192|| w->I!=K||w->O!=H*(Q+V))return 0; DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0; diff --git a/c/tests/test_backend_cuda.cu b/c/tests/test_backend_cuda.cu index 5550bde..259fe6b 100644 --- a/c/tests/test_backend_cuda.cu +++ b/c/tests/test_backend_cuda.cu @@ -50,6 +50,52 @@ int main(int argc, char **argv) { if (coli_cuda_tensor_upload(&t8, q8, s8, 1, 5, 2, d0)) return 1; if (ndev > 1 && coli_cuda_tensor_upload(&t8, q8, s8, 1, 4, 2, d1)) return 1; if (!coli_cuda_matmul(&t8, got, x, q8, s8, 1, 2, 4, 2, d0) || !close_enough(got, want8, 4)) return 1; + /* Cached tensor must stay callable without live host pointers + * (CUDA_RELEASE_HOST slots null theirs after upload) — including + * SUSTAINED reuse, not just the first call. */ + for (int rep = 0; rep < 64; rep++) + if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) || + !close_enough(got, want8, 4)) return 1; + /* A tensor uploaded from a TEMPORARY host buffer must survive the buffer + * being scribbled and freed (the release-host lifecycle). */ + { + int8_t *tmpw = static_cast(std::malloc(8)); + float *tmps = static_cast(std::malloc(2 * sizeof(float))); + if (!tmpw || !tmps) return 2; + for (int i = 0; i < 8; i++) tmpw[i] = q8[i]; + tmps[0] = s8[0]; tmps[1] = s8[1]; + ColiCudaTensor *tt = nullptr; + if (!coli_cuda_tensor_upload(&tt, tmpw, tmps, 1, 4, 2, d0)) return 1; + for (int i = 0; i < 8; i++) tmpw[i] = 99; + std::free(tmpw); std::free(tmps); + if (!coli_cuda_matmul(&tt, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) || + !close_enough(got, want8, 4)) return 1; + coli_cuda_tensor_free(tt); + } + /* Upload failures must be graceful and must not corrupt accounting — + * and must not poison LATER healthy launches (sticky-error regression). */ + { + size_t c0 = 0, b0 = 0, c1 = 0, b1 = 0; + coli_cuda_stats(-1, &c0, &b0); + ColiCudaTensor *bad = nullptr; + if (coli_cuda_tensor_upload(&bad, q8, s8, 1, 4, 2, 9999)) return 1; + if (coli_cuda_tensor_upload(&bad, q8, s8, 7, 4, 2, d0)) return 1; + if (coli_cuda_tensor_upload(&bad, q8, nullptr, 1, 4, 2, d0)) return 1; + if (coli_cuda_tensor_upload(&bad, nullptr, s8, 1, 4, 2, d0)) return 1; + if (coli_cuda_tensor_upload(&bad, q8, s8, 1, 1 << 20, 1 << 24, d0)) return 1; /* ~16 TB */ + if (bad) return 1; + coli_cuda_stats(-1, &c1, &b1); + if (c0 != c1 || b0 != b1) return 1; + /* healthy launch immediately after the failed allocation */ + if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) || + !close_enough(got, want8, 4)) return 1; + } + /* Fault injection hook: on/off, restores cleanly. */ + if (setenv("COLI_GPU_FAIL_AFTER", "0", 1)) return 2; + if (coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0)) return 1; + if (unsetenv("COLI_GPU_FAIL_AFTER")) return 2; + if (!coli_cuda_matmul(&t8, got, x, nullptr, nullptr, 1, 2, 4, 2, d0) || + !close_enough(got, want8, 4)) return 1; const int8_t q8b[8]={-1,-2,-3,-4, 1,-2,3,-4}; const float s8b[2]={1.f,.5f},want8b[4]={10.f,15.f,-3.f,-2.5f}; if(!coli_cuda_tensor_update(t8,q8b,s8b)|| From dc196633f1b168a3e6801f35daabacac2f4955ba Mon Sep 17 00:00:00 2001 From: ebootheee Date: Thu, 16 Jul 2026 23:20:16 -0600 Subject: [PATCH 05/71] pipe: blocking pipe_wait (COLI_PIPE_BLOCK=1) + PIPE_WORKERS implies PIPE=1 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit pipe_wait's sched_yield spin storms the scheduler for the full 0.5-3ms of each in-flight expert read; behind COLI_PIPE_BLOCK=1 it parks on a condvar instead (~5us wake, no lost-wakeup: workers store ready with release BEFORE taking the mutex to broadcast, and the waiter re-checks under the lock). Default OFF = byte-identical spin. Pthread pool only: the URING backend has no waiter spin to replace. Setting PIPE_WORKERS>0 in the env without PIPE now implies PIPE=1 with a stderr note: sizing the pool declares the intent to use it (a full benchmark campaign ran with PIPE_WORKERS=16 and the pipe silently off). The implication fires only when the platform default left the pipe off (no-op on _WIN32 where PIPE defaults to 1), only on a positive value (PIPE_WORKERS=0/empty does not enable a clamped 1-worker pipe), and an explicit PIPE=0 always wins. The rule lives in pipe_workers_imply_pipe() so the table is unit-testable. tests/test_pipe_block.c (in TEST_BINS, all platforms): pins the implication table, and drives the pool through 200 generations under each waiter against an on-disk expert fixture — spin as control, condvar arm alternating parked and fast-path waits — verifying identical slot bytes. Measured on the spin side (2x5090 + Gen5 NVMe, GLM-5.2 744B int4, CPU decode): 1.98 -> 2.16 tok/s at 192 tokens (expert-disk service 44.6s -> 33.5s). On Metal/GPU decode an M5 Pro A/B (PR thread) measured no change, consistent with the mechanism: the win is freeing the core the spinner was stealing from the CPU matmul team. Co-Authored-By: Claude Fable 5 --- c/Makefile | 5 +- c/glm.c | 39 +++++++++ c/tests/test_pipe_block.c | 169 ++++++++++++++++++++++++++++++++++++++ 3 files changed, 212 insertions(+), 1 deletion(-) create mode 100644 c/tests/test_pipe_block.c diff --git a/c/Makefile b/c/Makefile index 27902cb..db94a9a 100644 --- a/c/Makefile +++ b/c/Makefile @@ -166,7 +166,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_pipe_block$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -329,6 +329,9 @@ tests/test_compat_direct$(EXE): tests/test_compat_direct.c compat.h tests/test_uring$(EXE): tests/test_uring.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +tests/test_pipe_block$(EXE): tests/test_pipe_block.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + test-c: $(TEST_BINS) $(PYTHON) tools/run_tests.py $(TEST_BINS) diff --git a/c/glm.c b/c/glm.c index 645be12..6bd66f4 100644 --- a/c/glm.c +++ b/c/glm.c @@ -2146,6 +2146,22 @@ static int g_pipe=0; /* PIPE=1: async expert-load pipeline. Default ON for * the matmul. PIPE=0 opts back into the blocking serial path. */ static int g_pipe_nw=8; /* PIPE_WORKERS=n: I/O worker threads (disk-parallel reads) */ static int g_uring=0; /* URING=1: Linux io_uring load/completion backend; implies PIPE */ +static int g_pipe_block=0;/* COLI_PIPE_BLOCK=1: pipe_wait blocca su una condvar invece dello + * spin sched_yield (default OFF = spin byte-identico). EN: a yield + * storm on the main thread fights the OpenMP team for cycles during + * multi-ms loads; the condvar wake costs ~5us against reads that + * cost 0.5-3ms (#159). Pthread pool only: the URING backend has no + * waiter spin to replace. */ +/* PIPE_WORKERS>0 esplicito nell'env implica PIPE=1: dimensionare il pool + * dichiara l'intento di usarlo (una campagna intera l'ha impostato con la + * pipe spenta senza accorgersene). EN: fires ONLY when PIPE is unset in the + * env AND the platform default left the pipe off (on _WIN32 it already + * defaults to 1) AND PIPE_WORKERS parses positive — the internal default of + * 8 does not count, PIPE_WORKERS=0/empty does not, and an explicit PIPE=0 + * always wins. */ +static int pipe_workers_imply_pipe(const char *pipe_env, const char *pw_env, int pipe_now){ + return !pipe_now && !pipe_env && pw_env && atoi(pw_env)>0; +} typedef struct { _Atomic uint64_t cur; /* (gen<<8)|index; gen main-only, index 0..njobs (≤64) */ _Atomic int njobs; /* current batch job count */ @@ -2153,6 +2169,7 @@ typedef struct { _Atomic int layer; /* current batch layer */ _Atomic int ready[64]; /* per-slot load-done flag */ pthread_mutex_t mx; pthread_cond_t cv; /* ONLY for parking/waking idle workers */ + pthread_cond_t cv_done; /* COLI_PIPE_BLOCK: signals ready[] transitions */ Model *m; pthread_t th[16]; int nw; int started; } PipePool; @@ -2177,6 +2194,11 @@ static void *pipe_worker(void *arg){ int eid=atomic_load_explicit(&p->eids[i],memory_order_relaxed); /* AFTER winning CAS */ expert_load(p->m,L,eid,&p->m->ws[i],1); /* needed-now load: fatal on I/O error (matches serial path) */ atomic_store_explicit(&p->ready[i],1,memory_order_release); + if(g_pipe_block){ /* wake a main thread parked in pipe_wait */ + pthread_mutex_lock(&p->mx); + pthread_cond_broadcast(&p->cv_done); + pthread_mutex_unlock(&p->mx); + } } /* CAS failed → another worker advanced index (or gen advanced): re-loop */ } @@ -2194,6 +2216,7 @@ static void pipe_init(Model *m){ g_pp.m=m; g_pp.nw=g_pipe_nw; if(g_pp.nw>16) g_pp.nw=16; if(g_pp.nw<1) g_pp.nw=1; atomic_store(&g_pp.cur,0); atomic_store(&g_pp.njobs,0); pthread_mutex_init(&g_pp.mx,NULL); pthread_cond_init(&g_pp.cv,NULL); + pthread_cond_init(&g_pp.cv_done,NULL); for(int i=0;i PIPE implication table: fires ONLY when + * PIPE is unset in the env AND the platform default left the pipe off AND + * PIPE_WORKERS parses positive (PIPE_WORKERS=0/empty/negative must NOT + * silently enable a clamped 1-worker pipe). */ +#include +#include +#include +#include +#include +#include +#define main coli_glm_main_unused +#include "../glm.c" +#undef main + +static int fail(const char *s){ fprintf(stderr,"FAIL: %s\n",s); return 1; } + +enum { NE=8, LAYER=1 }; /* experts 0..NE-1 on one MoE layer */ +/* per-expert file image: [gate 12][up 12][down 12][gate.qs 12][up.qs 12][down.qs 16] */ +enum { WB=12, QS_G=12, QS_U=12, QS_D=16, ESZ=3*WB+QS_G+QS_U+QS_D }; + +static unsigned char wbyte(int e,int j){ return (unsigned char)(e*31+j+1); } +static float scale(int e,int i){ return (float)(e*8+i)+0.5f; } + +#define TMPF "test_pipe_block.tmp" + +static int write_fixture(void){ + FILE *w=fopen(TMPF,"wb"); if(!w) return fail("create temp"); + for(int e=0;ec.hidden=4; m->c.moe_inter=3; m->ebits=8; + m->S.n=NE*6; m->S.cap=NE*6; m->S.t=calloc(NE*6,sizeof(st_tensor)); + if(!m->S.t) return fail("tensor metadata allocation"); + const char *proj[3]={"gate_proj","up_proj","down_proj"}; + int sbytes[3]={QS_G,QS_U,QS_D}; + for(int e=0;eS.t[e*6+k]=(st_tensor){strdup(name),fd,wo,WB,3,WB}; wo+=WB; + size_t n=strlen(name); memcpy(name+n,".qs",4); + m->S.t[e*6+3+k]=(st_tensor){strdup(name),fd,so,sbytes[k],2,sbytes[k]/4}; so+=sbytes[k]; + } + } + return 0; +} + +static int check_slot(ESlot *s,int e){ + if(s->eid!=e || s->g.fmt!=1 || s->u.fmt!=1 || s->d.fmt!=1){ + fprintf(stderr," slot: eid=%d (want %d) fmt g/u/d=%d/%d/%d (want 1/1/1)\n", + s->eid,e,s->g.fmt,s->u.fmt,s->d.fmt); + return 1; + } + const unsigned char *g=(const unsigned char*)s->g.q8, + *u=(const unsigned char*)s->u.q8, + *d=(const unsigned char*)s->d.q8; /* q8 is int8_t; compare raw bytes */ + for(int j=0;jg.s[i]!=scale(e,i) || s->u.s[i]!=scale(e,3+i)){ + fprintf(stderr," slot e=%d scale %d: g=%g/%g u=%g/%g (got/want)\n",e,i, + (double)s->g.s[i],(double)scale(e,i),(double)s->u.s[i],(double)scale(e,3+i)); + return 1; + } + for(int i=0;i<4;i++) + if(s->d.s[i]!=scale(e,6+i)){ + fprintf(stderr," slot e=%d scale %d: d=%g/%g (got/want)\n",e,i, + (double)s->d.s[i],(double)scale(e,6+i)); + return 1; + } + return 0; +} + +static int run_generations(Model *m,int block,int gens){ + g_pipe_block=block; + for(int gen=0;genws[q],eids[q])) return fail(block?"slot contents (block)":"slot contents (spin)"); + } + } + return 0; +} + +static int test_implication_table(void){ + struct { const char *pipe_env,*pw_env; int pipe_now,want; } T[]={ + {NULL,"4",0,1}, /* pool sized, pipe off, PIPE unset → imply */ + {NULL,"16",0,1}, + {NULL,"0",0,0}, /* PIPE_WORKERS=0 must NOT enable a clamped pipe */ + {NULL,"",0,0}, + {NULL,"-3",0,0}, + {"0","4",0,0}, /* explicit PIPE=0 always wins */ + {"1","4",1,0}, /* explicit PIPE=1: nothing to imply */ + {NULL,"4",1,0}, /* platform default already ON (win32) */ + {NULL,NULL,0,0}, + }; + for(size_t i=0;i",T[i].pw_env?T[i].pw_env:"",T[i].pipe_now); + return 1; + } + return 0; +} + +int main(void){ + if(test_implication_table()) return 1; + + /* Relative to the CWD, like test_compat_direct's TMPF — NOT "/tmp/...": + * the windows job builds native .exe files and "/tmp" is not a Windows + * path. fwrite then reopen read-only: Windows compat has pread, not pwrite. */ + if(write_fixture()) return 1; + int fd=open(TMPF,COMPAT_O_RDONLY); + if(fd<0) return fail("open temp"); + + static Model m; /* zeroed: buffered pread path, no mmap/cuda */ + if(build_fixture(&m,fd)){ close(fd); remove(TMPF); return 1; } + + g_pipe=1; g_pipe_nw=4; + pipe_init(&m); + + /* spin waiter first (control), then the condvar waiter under the same + * dispatch pattern; 200 generations each cycles the gen-tagged cursor + * and alternates parked/fast-path waits. */ + if(run_generations(&m,0,200) || run_generations(&m,1,200)){ close(fd); remove(TMPF); return 1; } + + for(int q=0;q Date: Sun, 19 Jul 2026 10:06:27 +0200 Subject: [PATCH 06/71] flake: remove unnecessary rec --- flake.nix | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/flake.nix b/flake.nix index 584a088..04d3133 100644 --- a/flake.nix +++ b/flake.nix @@ -90,8 +90,7 @@ mainProgram = "glm"; }; }; - in - rec { + in { packages = { default = colibri; inherit colibri; From 73aef4d010366418c74a5b5c2f737c0c530bf7bf Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 10:14:24 +0200 Subject: [PATCH 07/71] nix: configure and apply a formatter This sets the formatter to alejandra. I was trying to stay as close as the original as I could; otherwise we could set it to `nixfmt` which would keep list spacing more similar, but would introduce additional indentation in a few places. --- flake.nix | 39 ++++++++++++++++++++++++--------------- 1 file changed, 24 insertions(+), 15 deletions(-) diff --git a/flake.nix b/flake.nix index 04d3133..a63839a 100644 --- a/flake.nix +++ b/flake.nix @@ -6,20 +6,27 @@ flake-utils.url = "github:numtide/flake-utils"; }; - outputs = { self, nixpkgs, flake-utils }: - flake-utils.lib.eachDefaultSystem (system: - let - pkgs = import nixpkgs { inherit system; }; + outputs = { + self, + nixpkgs, + flake-utils, + }: + flake-utils.lib.eachDefaultSystem ( + system: let + pkgs = import nixpkgs {inherit system;}; # Python with the packages needed by the offline converter tools - pythonEnv = pkgs.python3.withPackages (ps: with ps; [ - torch - safetensors - huggingface-hub - numpy - tokenizers - datasets - ]); + pythonEnv = pkgs.python3.withPackages ( + ps: + with ps; [ + torch + safetensors + huggingface-hub + numpy + tokenizers + datasets + ] + ); colibri = pkgs.stdenv.mkDerivation { pname = "colibri"; @@ -28,7 +35,7 @@ # python3 is needed by checkPhase: `make test-c` shells out to # `python3 tools/run_tests.py` (see c/Makefile, PYTHON ?= python3). - nativeBuildInputs = [ pkgs.makeWrapper pkgs.python3 ]; + nativeBuildInputs = [pkgs.makeWrapper pkgs.python3]; buildInputs = [ pkgs.gcc @@ -107,14 +114,16 @@ }; }; + formatter = (import nixpkgs {inherit system;}).alejandra; + devShells.default = pkgs.mkShell { - inputsFrom = [ colibri ]; + inputsFrom = [colibri]; packages = [ pythonEnv pkgs.gcc pkgs.gnumake - pkgs.clang-tools # clangd / clang-tidy for IDE support + pkgs.clang-tools # clangd / clang-tidy for IDE support pkgs.pkg-config ]; From a6d0a6c4a7e0d43ee139a65517af414eb98d20b1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 10:21:50 +0200 Subject: [PATCH 08/71] nix: use with pkgs in some package lists --- flake.nix | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/flake.nix b/flake.nix index a63839a..96d0625 100644 --- a/flake.nix +++ b/flake.nix @@ -35,11 +35,11 @@ # python3 is needed by checkPhase: `make test-c` shells out to # `python3 tools/run_tests.py` (see c/Makefile, PYTHON ?= python3). - nativeBuildInputs = [pkgs.makeWrapper pkgs.python3]; + nativeBuildInputs = with pkgs; [makeWrapper python3]; - buildInputs = [ - pkgs.gcc - pkgs.gmp + buildInputs = with pkgs; [ + gcc + gmp ]; # Use x86-64-v3 (AVX2) for a portable binary; override with ARCH=native for local builds @@ -119,12 +119,12 @@ devShells.default = pkgs.mkShell { inputsFrom = [colibri]; - packages = [ + packages = with pkgs; [ pythonEnv - pkgs.gcc - pkgs.gnumake - pkgs.clang-tools # clangd / clang-tidy for IDE support - pkgs.pkg-config + gcc + gnumake + clang-tools # clangd / clang-tidy for IDE support + pkg-config ]; shellHook = '' From e93d574e139995efeb921008355e7fece380ac82 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 10:29:46 +0200 Subject: [PATCH 09/71] nix: add the python env to check inputs The tests now require Python so the env should be added to the check inputs, otherwise the build fails. --- flake.nix | 2 ++ 1 file changed, 2 insertions(+) diff --git a/flake.nix b/flake.nix index 96d0625..af249bd 100644 --- a/flake.nix +++ b/flake.nix @@ -42,6 +42,8 @@ gmp ]; + checkInputs = [pythonEnv]; + # Use x86-64-v3 (AVX2) for a portable binary; override with ARCH=native for local builds ARCH = "x86-64-v3"; From bcb984a61f2a8305d543c38af11003bd503a445a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 10:40:19 +0200 Subject: [PATCH 10/71] nix: use getExe instead of hardcoding the path --- flake.nix | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/flake.nix b/flake.nix index af249bd..276dbb6 100644 --- a/flake.nix +++ b/flake.nix @@ -108,11 +108,11 @@ apps = { default = { type = "app"; - program = "${colibri}/bin/glm"; + program = pkgs.lib.getExe colibri; }; coli = { type = "app"; - program = "${colibri}/bin/coli"; + program = pkgs.lib.getExe' colibri "coli"; }; }; From 58fd0e557f23e95827d1be9a8226befbfffefca7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 10:48:00 +0200 Subject: [PATCH 11/71] nix: commit lockfile, ignore result This makes the build actually reproducible by committing a working lockfile to the repo. --- .gitignore | 1 + flake.lock | 6 +++--- 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/.gitignore b/.gitignore index 981e740..6e17a30 100644 --- a/.gitignore +++ b/.gitignore @@ -35,6 +35,7 @@ c/tests/test_schema_gbnf c/tests/test_schema_gbnf.exe c/tests/test_compat_direct c/tests/test_compat_direct.exe +result # oracoli tiny generati (make_glm_oracle.py) e dati benchmark scaricati c/glm_tiny/ diff --git a/flake.lock b/flake.lock index 9b788fa..a89029e 100644 --- a/flake.lock +++ b/flake.lock @@ -20,11 +20,11 @@ }, "nixpkgs": { "locked": { - "lastModified": 1784160687, - "narHash": "sha256-iYL/bixrb6FlHFu/gIuBYzq6c6lM5AAXsXNSWXtIgQc=", + "lastModified": 1784280462, + "narHash": "sha256-DtoqIqM7VkR6NxAkcLpMwmi02USwWb3JdmNGLyhthc0=", "owner": "NixOS", "repo": "nixpkgs", - "rev": "4382ed2b7a6839d4280a9b386db49cbc5907414d", + "rev": "293d6abedf0478e681a4dfcfcb35b30fc796a32f", "type": "github" }, "original": { From da97f0dbf1834b6801f05744e55e73e0b926e481 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 10:59:16 +0200 Subject: [PATCH 12/71] nix: build on darwin with -march=native --- flake.nix | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/flake.nix b/flake.nix index 276dbb6..df83d4c 100644 --- a/flake.nix +++ b/flake.nix @@ -45,7 +45,10 @@ checkInputs = [pythonEnv]; # Use x86-64-v3 (AVX2) for a portable binary; override with ARCH=native for local builds - ARCH = "x86-64-v3"; + ARCH = + if pkgs.stdenv.hostPlatform.isx86_64 + then "x86-64-v3" + else "native"; buildPhase = '' runHook preBuild @@ -95,7 +98,7 @@ description = "Run GLM-5.2 (744B MoE) on a consumer machine with ~25 GB RAM"; homepage = "https://github.com/JustVugg/colibri"; license = licenses.asl20; - platforms = platforms.linux; + platforms = with platforms; linux ++ darwin; mainProgram = "glm"; }; }; From 51fe03f615c42a07f5dd31f0287cf7105d22c2cf Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 11:36:13 +0200 Subject: [PATCH 13/71] nix: set main program to coli This is the higher-level user interface, which should be the main entry point to the binary, not the engine itself. The engine is still available. --- flake.nix | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/flake.nix b/flake.nix index df83d4c..17f8a36 100644 --- a/flake.nix +++ b/flake.nix @@ -99,7 +99,7 @@ homepage = "https://github.com/JustVugg/colibri"; license = licenses.asl20; platforms = with platforms; linux ++ darwin; - mainProgram = "glm"; + mainProgram = "coli"; }; }; in { @@ -113,9 +113,9 @@ type = "app"; program = pkgs.lib.getExe colibri; }; - coli = { + glm = { type = "app"; - program = pkgs.lib.getExe' colibri "coli"; + program = "${colibri}/share/colibri/glm"; }; }; From 4775f28cebae2e127a9e2b883d8e7cf24b1f98bb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 13:50:45 +0200 Subject: [PATCH 14/71] nix: copy missing version.py to the environment --- flake.nix | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/flake.nix b/flake.nix index 17f8a36..3d62a5e 100644 --- a/flake.nix +++ b/flake.nix @@ -68,7 +68,8 @@ cp c/glm $out/lib/colibri/glm cp c/coli $out/lib/colibri/coli chmod +x $out/lib/colibri/coli - cp c/openai_server.py c/resource_plan.py c/doctor.py $out/lib/colibri/ + cp c/openai_server.py c/resource_plan.py c/doctor.py c/version.py \ + $out/lib/colibri/ cp -r c/tools/* $out/lib/colibri/tools/ # $out/bin holds the user-facing entry points. From 40a3596354d025e686e8e1203ea16862c7887afd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Attila=20Ol=C3=A1h?= Date: Sun, 19 Jul 2026 14:09:45 +0200 Subject: [PATCH 15/71] nix: python3 is only needed for checks, not for the actual build --- flake.nix | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/flake.nix b/flake.nix index 3d62a5e..55f95ee 100644 --- a/flake.nix +++ b/flake.nix @@ -33,16 +33,16 @@ version = "1.0"; src = ./.; - # python3 is needed by checkPhase: `make test-c` shells out to - # `python3 tools/run_tests.py` (see c/Makefile, PYTHON ?= python3). - nativeBuildInputs = with pkgs; [makeWrapper python3]; + nativeBuildInputs = with pkgs; [makeWrapper]; buildInputs = with pkgs; [ gcc gmp ]; - checkInputs = [pythonEnv]; + # python3 is needed by checkPhase: `make test-c` shells out to + # `python3 tools/run_tests.py` (see c/Makefile, PYTHON ?= python3). + nativeCheckInputs = with pkgs; [python3]; # Use x86-64-v3 (AVX2) for a portable binary; override with ARCH=native for local builds ARCH = From 8a9a0fca4d9c5b814d42799c97dc8780b3b3ae8c Mon Sep 17 00:00:00 2001 From: ZacharyZcR Date: Sun, 19 Jul 2026 04:50:45 +0800 Subject: [PATCH 16/71] =?UTF-8?q?plan:=20auto-tune=20heuristics=20?= =?UTF-8?q?=E2=80=94=20bottleneck=20classification=20+=20parameter=20decis?= =?UTF-8?q?ions?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Extend resource_plan to classify the hardware into bottleneck regimes (disk / memory / mixed / compute) and derive tuning knobs automatically: MTP: off when compute-bound (42% loss at full residency, #389) or disk-bound with <90% hit (union growth adds reads) PIPE: COLI_CUDA_PIPE=1 single-GPU, =2 multi-GPU, PIPE=1 CPU disk NUMA: selective interleave for GPU hosts, blanket hint for CPU-only PIN: PIN_GB=all when fully resident + no GPU OMP: COLI_NO_OMP_TUNE=1 for Metal (spin steals GPU power) `coli plan` now shows an auto-tune section with each knob and its reason. `environment_for_plan()` applies them via setdefault so explicit user settings always win. plan version stays at 2 (additive fields: bottleneck_class, projected_hit_rate, tune). 7 new tests covering all regimes. --- c/resource_plan.py | 95 ++++++++++++++++++++++++++++++++--- c/tests/test_resource_plan.py | 75 +++++++++++++++++++++++++++ 2 files changed, 163 insertions(+), 7 deletions(-) diff --git a/c/resource_plan.py b/c/resource_plan.py index 5f17a6a..15ce02a 100644 --- a/c/resource_plan.py +++ b/c/resource_plan.py @@ -219,6 +219,54 @@ def cpu_socket_count(): return 1 +def _auto_tune(bottleneck_class, projected_hit, gpus, cpu_sockets, plan_has_metal): + """Derive tuning knobs from the bottleneck classification.""" + tune = {} + has_gpu = bool(gpus) + n_gpu = len(gpus) + + # MTP: costs more than it saves when compute-bound (#389 measured 42% loss) + if bottleneck_class == "compute": + tune["DRAFT"] = {"value": "0", + "reason": "compute-bound: MTP batch overhead exceeds yield"} + elif bottleneck_class == "disk" and projected_hit < 0.90: + tune["DRAFT"] = {"value": "0", + "reason": "low hit rate: MTP widens expert union, adds disk reads"} + # otherwise leave DRAFT unset (engine default: auto) + + # PIPE: resident pipeline mode depends on GPU count + if has_gpu and n_gpu == 1: + tune["COLI_CUDA_PIPE"] = {"value": "1", + "reason": "single GPU: S=1 pipeline gate"} + elif has_gpu and n_gpu > 1: + tune["COLI_CUDA_PIPE"] = {"value": "2", + "reason": "multi-GPU: residual stays on-device across layers"} + elif not has_gpu and bottleneck_class == "disk": + tune["PIPE"] = {"value": "1", + "reason": "overlap disk reads with resident expert compute"} + + # NUMA: selective interleave for GPU hosts, blanket hint for CPU-only + if cpu_sockets > 1 and has_gpu: + tune["COLI_NUMA"] = {"value": "1", + "reason": "multi-socket + GPU: interleave expert slabs, protect DMA buffers"} + elif cpu_sockets > 1 and not has_gpu: + tune["COLI_NUMA"] = {"value": "1", + "reason": "multi-socket CPU-only: interleave expert slabs across nodes"} + tune["_numa_hint"] = "numactl --interleave=all may perform better on CPU-only hosts" + + # OMP: kill hot-thread spin when GPU/Metal owns the power budget + if plan_has_metal: + tune["COLI_NO_OMP_TUNE"] = {"value": "1", + "reason": "Metal: OMP spin-wait steals GPU power budget"} + + # PIN: fully resident if RAM allows and no GPU tier competes + if projected_hit >= 0.99 and not has_gpu: + tune["PIN_GB"] = {"value": "all", + "reason": "enough RAM for full expert residency"} + + return tune + + POLICIES = { "quality": {"preserve_quantization": True, "preserve_router": True}, "balanced": {"preserve_quantization": True, "preserve_router": True}, @@ -290,12 +338,28 @@ def build_plan(model, ram_gb=0, context=4096, gpu_indices=None, vram_gb=0, if cold_bytes: warnings.append("cold expert misses may reach disk; normal decode speed depends on hit rate") + total_expert = info["expert_bytes"] + resident_expert = hot_bytes + warm_bytes + projected_hit = resident_expert / total_expert if total_expert else 1.0 + if cold_bytes: bottleneck = "disk expert misses" - elif warm_bytes: - bottleneck = "CPU expert compute and RAM bandwidth" + bottleneck_class = "disk" + elif warm_bytes and gpus: + bottleneck = "CPU expert tail and GPU compute" + bottleneck_class = "mixed" + elif projected_hit >= 0.99: + if gpus: + bottleneck = "GPU compute and interconnect" + else: + bottleneck = "CPU expert compute (fully resident)" + bottleneck_class = "compute" else: - bottleneck = "GPU compute and interconnect" + bottleneck = "CPU expert compute and RAM bandwidth" + bottleneck_class = "memory" + + tune = _auto_tune(bottleneck_class, projected_hit, gpus, cpu_sockets, + plan_has_metal=False) return { "version": 2, @@ -317,6 +381,9 @@ def build_plan(model, ram_gb=0, context=4096, gpu_indices=None, vram_gb=0, "expert_capacity": vram_experts, "requires_host_backing": False}, }, "expected_bottleneck": bottleneck, + "bottleneck_class": bottleneck_class, + "projected_hit_rate": round(projected_hit, 4), + "tune": tune, "decisions": [ {"target": "VRAM", "reason": "profile-ranked hot experts"}, {"target": "RAM", "reason": "warm experts execute on CPU without quality loss"}, @@ -336,10 +403,11 @@ def environment_for_plan(plan, env=None, cuda_enabled=True): # ("Affinity not supported on this configuration"): non impostarle li'. result.setdefault("OMP_PROC_BIND", "spread") result.setdefault("OMP_PLACES", "cores") - if sys.platform.startswith("linux") and plan["cpu"].get("sockets", 1) > 1: - # Selectively interleave large expert/dense slabs across memory controllers. - # Unlike blanket numactl interleave, this leaves CUDA staging buffers local. - result.setdefault("COLI_NUMA", "1") + tune = plan.get("tune", {}) + for key, entry in tune.items(): + if key.startswith("_"): + continue + result.setdefault(key, entry["value"]) if plan["policy"]["name"] == "balanced": result.setdefault("REPIN", "64") ram = plan["tiers"]["ram"] @@ -386,5 +454,18 @@ def format_plan(plan): else: lines.append("VRAM no NVIDIA device detected · CPU path") lines.append(f"limit {plan['expected_bottleneck']}") + hit = plan.get("projected_hit_rate", 0) + lines.append(f"hit {hit:.0%} projected expert residency") + tune = plan.get("tune", {}) + if tune: + lines.append("") + lines.append("auto-tune:") + for key, entry in tune.items(): + if key.startswith("_"): + continue + lines.append(f" {key}={entry['value']:12s} {entry['reason']}") + hint = tune.get("_numa_hint") + if hint: + lines.append(f" hint: {hint}") lines.extend(f"warn {warning}" for warning in plan["warnings"]) return "\n".join(lines) diff --git a/c/tests/test_resource_plan.py b/c/tests/test_resource_plan.py index b184c80..4ef4706 100644 --- a/c/tests/test_resource_plan.py +++ b/c/tests/test_resource_plan.py @@ -140,6 +140,81 @@ class ResourcePlanTest(unittest.TestCase): plan = build_plan(self.model, available_memory=16 * GB, available_disk=1, gpus=[], physical_cpus=8, cpu_sockets=1) self.assertNotIn("COLI_NUMA", environment_for_plan(plan)) + + def test_auto_tune_mtp_off_when_compute_bound(self): + # Tiny model with 64 GB RAM and no GPU: all experts fit in RAM with no + # warm tier, so the plan classifies as compute-bound. + plan = build_plan(self.model, ram_gb=64, available_memory=64 * GB, + available_disk=100 * GB, gpus=[], physical_cpus=24, + cpu_sockets=2) + # With such a small model fully in RAM and no GPU, bottleneck is compute + self.assertEqual(plan["bottleneck_class"], "compute") + self.assertIn("DRAFT", plan["tune"]) + self.assertEqual(plan["tune"]["DRAFT"]["value"], "0") + env = environment_for_plan(plan) + self.assertEqual(env["DRAFT"], "0") + explicit = environment_for_plan(plan, {"DRAFT": "3"}) + self.assertEqual(explicit["DRAFT"], "3") + + def test_auto_tune_mtp_off_when_disk_low_hit(self): + # Use a model large enough that 8 GB RAM can't hold all experts. + big = tempfile.TemporaryDirectory() + bigmodel = Path(big.name) + (bigmodel / "config.json").write_text(json.dumps({ + "num_hidden_layers": 2, "n_routed_experts": 4, + "kv_lora_rank": 4, "qk_rope_head_dim": 2, + "qk_nope_head_dim": 3, "v_head_dim": 5, "num_attention_heads": 2, + })) + expert_size = 3 * GB # each expert 3 GB → 12 GB total, won't fit in 8 GB budget + write_shard(bigmodel / "out-00000.safetensors", [ + ("model.embed_tokens.weight", 100), + ("model.layers.0.self_attn.q_a_proj.weight", 200), + ]) + for i in range(4): + write_shard(bigmodel / f"out-{i+1:05d}.safetensors", [ + (f"model.layers.1.mlp.experts.{i}.gate_proj.weight", expert_size), + ]) + plan = build_plan(bigmodel, ram_gb=0, available_memory=4 * GB, + available_disk=100 * GB, gpus=[], physical_cpus=8, + cpu_sockets=1) + big.cleanup() + self.assertEqual(plan["bottleneck_class"], "disk") + self.assertLess(plan["projected_hit_rate"], 0.90) + self.assertEqual(plan["tune"]["DRAFT"]["value"], "0") + + def test_auto_tune_pipe_multi_gpu(self): + gpus = [ + {"index": 0, "name": "a", "total_bytes": 32 * GB, "free_bytes": 30 * GB}, + {"index": 1, "name": "b", "total_bytes": 32 * GB, "free_bytes": 30 * GB}, + ] + plan = build_plan(self.model, ram_gb=16, available_memory=32 * GB, + available_disk=1, gpus=gpus, cpu_sockets=2) + self.assertEqual(plan["tune"]["COLI_CUDA_PIPE"]["value"], "2") + env = environment_for_plan(plan) + self.assertEqual(env["COLI_CUDA_PIPE"], "2") + + def test_auto_tune_pipe_single_gpu(self): + gpus = [{"index": 0, "name": "a", "total_bytes": 12 * GB, "free_bytes": 10 * GB}] + plan = build_plan(self.model, ram_gb=16, available_memory=32 * GB, + available_disk=1, gpus=gpus, cpu_sockets=1) + self.assertEqual(plan["tune"]["COLI_CUDA_PIPE"]["value"], "1") + + def test_auto_tune_numa_hint_for_cpu_only(self): + plan = build_plan(self.model, ram_gb=64, available_memory=64 * GB, + available_disk=1, gpus=[], physical_cpus=64, cpu_sockets=2) + self.assertIn("_numa_hint", plan["tune"]) + self.assertIn("numactl", plan["tune"]["_numa_hint"]) + self.assertIn("auto-tune", format_plan(plan)) + + def test_format_plan_shows_tune_and_hit_rate(self): + plan = build_plan(self.model, ram_gb=64, available_memory=64 * GB, + available_disk=100 * GB, gpus=[], physical_cpus=24, + cpu_sockets=1) + text = format_plan(plan) + self.assertIn("hit", text) + self.assertIn("auto-tune", text) + self.assertIn("DRAFT", text) + def test_cpu_binary_does_not_apply_gpu_tier(self): plan = build_plan(self.model, available_memory=16 * GB, available_disk=1, gpus=[{"index": 0, "name": "a", "total_bytes": 8 * GB, From 27b4c4263a74bbfd568f5b91be01b0d5481eac6a Mon Sep 17 00:00:00 2001 From: Davide Quack Date: Sun, 19 Jul 2026 16:49:29 +0200 Subject: [PATCH 17/71] =?UTF-8?q?Added=20Dockerfile=20and=20instructions?= =?UTF-8?q?=20for=20using=20Colibr=C3=AC=20through=20a=20docker=20containe?= =?UTF-8?q?r?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- docker/Dockerfile | 15 ++ docker/README.IT.md | 454 +++++++++++++++++++++++++++++++++++++++++ docker/README.md | 488 ++++++++++++++++++++++++++++++++++++++++++++ 3 files changed, 957 insertions(+) create mode 100644 docker/Dockerfile create mode 100644 docker/README.IT.md create mode 100644 docker/README.md diff --git a/docker/Dockerfile b/docker/Dockerfile new file mode 100644 index 0000000..fa639f9 --- /dev/null +++ b/docker/Dockerfile @@ -0,0 +1,15 @@ +FROM debian:stable-slim + +# We install the necessary packages to download and compile the code +RUN apt-get update && \ + apt-get install -y locales git build-essential python3 python3-pip && \ + rm -rf /var/lib/apt/lists/* && \ + mkdir /app + +# Let's clone the repository +WORKDIR /app +RUN git clone https://github.com/JustVugg/colibri.git . +WORKDIR /app/c +# It's time to compile +RUN ./setup.sh && ./coli build + diff --git a/docker/README.IT.md b/docker/README.IT.md new file mode 100644 index 0000000..43c3e31 --- /dev/null +++ b/docker/README.IT.md @@ -0,0 +1,454 @@ +_Read the read me in [English](README.md)._ + + +# Guida a Colibrì - Motore di Inferenza Locale + +Una guida semplice per eseguire **Colibrì**, un motore di inferenza locale basato su GLM 5.2, senza conoscenze di programmazione. Se hai già Docker installato, sei a buon punto. + +--- + +## 📋 Sommario + +- [Cosa è Colibrì?](#cosa-è-colibrì) +- [Cosa serve](#cosa-serve) + - [Hardware](#hardware) + - [Software](#software) +- [Come iniziare](#come-iniziare) + - [Passo 1: Scarica il modello](#passo-1-scarica-il-modello) + - [Passo 2: Scarica il Dockerfile di Colibrì](#passo-2-scarica-il-dockerfile-di-colibrì) + - [Passo 3: Compila l'immagine Docker](#passo-3-compila-limmagine-docker) + - [Passo 4: Avvia Colibrì](#passo-4-avvia-colibr%C3%AC) + - [Cosa significa quel comando?](#cosa-significa-quel-comando) + - [Usare Colibrì](#usare-colibrì) +- [Entrare in una console Linux dentro il container](#entrare-in-una-console-linux-dentro-il-container) +- [Risoluzione dei problemi](#risoluzione-dei-problemi) +- [Note tecniche](#note-tecniche) +- [Domande frequenti](#domande-frequenti) +- [Supporto e contributi](#supporto-e-contributi) +- [Test sul mio PC](#test-sul-mio-pc) + +--- + +## Cosa è Colibrì? + +Colibrì è un'applicazione che ti permette di eseguire un modello di intelligenza artificiale (GLM 5.2) direttamente sul tuo computer, senza connettersi a server esterni. È possbile anche farlo girare in Docker, che isola l'applicazione dal resto del sistema. + +> **Nota importante**: Il modello è molto grande. Attendi anche diversi minuti per una risposta a una domanda semplice, specialmente con poca RAM. Alla fine di questo readme vedrai il risultato sul mio PC (senza scheda grafica discreta) e arrivo a 0.01 token al secondo. + +--- + +## Cosa serve + +### Hardware + +| Memoria RAM | Funziona? | Note | +|:---:|:---:|---| +| < 16 GB | ❌ No | Memoria insufficiente | +| 24 GB | ⚠️ Forse | Possibile, da testare | +| 32 GB | ✅ Sì | Il minimo (ma vedi sezione memoria su Windows) | +| 48+ GB | ✅ Sì | Meglio | + +Inoltre: un **disco SSD veloce** è essenziale. Colibrì usa il disco come memoria aggiuntiva. Con una scheda grafica NVidia è ancora meglio. + +### Software + +- **Docker Desktop** (Windows, Mac, Linux) — [scarica qui](https://www.docker.com/products/docker-desktop/) +- **Python** (solo se vuoi scaricare il modello da casa tua) + - Windows: [python.org](https://www.python.org) oppure Microsoft Store + - Linux: `apt-get install python3 python3-pip` + - Mac: [python.org](https://www.python.org) oppure Homebrew + +Non serve nessun ambiente di compilazione. Tutto avviene dentro il container Docker. + +--- + +## Come iniziare + +### Passo 1: Scarica il modello + +Il modello GLM 5.2 è circa **360 GB**. Scegli uno di questi metodi: + +#### Metodo A: Con Python (consigliato) + +1. **Installa la libreria per Hugging Face:** + ```bash + python -m pip install -U huggingface_hub[cli] + ``` + Su Linux, usa `python3` al posto di `python`. + +2. **Scarica il modello** (apri il terminale nella cartella dove lo vuoi salvare): + ```bash + hf_download mateogrgic/GLM-5.2-colibri-int4-with-int8-mtp --local-dir . + ``` + + **Esempio**: se vuoi salvarlo in `C:\LLM\models\glm-5.2` (Windows): + - Apri PowerShell in quella cartella + - Copia e incolla il comando sopra + - Attendi (molto) + +#### Metodo B: Senza Python (solo se necessario) + +Se sei su Windows e non riesci con Python: +- Scarica manualmente da [Hugging Face](https://huggingface.co/mateogrgic/GLM-5.2-colibri-int4-with-int8-mtp) +- Decomprimi in una cartella (es. `C:\LLM\models\glm-5.2`) + +--- + +### Passo 2: Scarica il Dockerfile di Colibrì + +1. Vai a: https://github.com/JustVugg/colibri/blob/main/docker/Dockerfile +2. Clicca il pulsante **Download** (icona ⬇️) in alto a destra +3. Salva il file in una cartella (es. `C:\LLM\Colibrì`) + +--- + +### Passo 3: Compila l'immagine Docker + +Apri il terminale (PowerShell su Windows, Terminal su Mac/Linux) **nella cartella dove hai salvato il Dockerfile** e digita: + +**Windows:** +```bash +docker build -t colibri-i . +``` + +**Linux/Mac:** +```bash +sudo docker build -t colibri-i . +``` + +Attendi che finisca (pochi minuti). Se tutto va bene, vedrai: `Successfully tagged colibri-i:latest` + +> **Se vuoi recepire gli aggiornamenti del repository**: Cancella prima l'immagine vecchia con `docker rmi colibri-i` e ricompila. + +--- + +### Passo 4: Avvia Colibrì + +Apri il terminale e digita il comando sottostante (sostituisci `C:\LLM\models\glm-5.2` con il percorso reale del tuo PC): + +**Windows** (PowerShell): +```bash +$MODEL_PATH="C:\LLM\models\glm-5.2" +docker run --rm -it --name colibri-c ` + -v "$MODEL_PATH`:/app/glm-5.2" ` + -e COLI_MODEL=/app/glm-5.2 ` + colibri-i ./coli chat +``` + +**Mac/Linux** (Terminal/Bash): +```bash +MODEL_PATH="/path/to/glm-5.2" +docker run --rm -it --name colibri-c \ + -v "$MODEL_PATH:/app/glm-5.2" \ + -e COLI_MODEL=/app/glm-5.2 \ + colibri-i ./coli chat +``` + +**Esempio per Linux:** +```bash +MODEL_PATH="/home/user/LLM/glm-5.2" +docker run --rm -it --name colibri-c \ + -v "$MODEL_PATH:/app/glm-5.2" \ + -e COLI_MODEL=/app/glm-5.2 \ + colibri-i ./coli chat +``` + +--- + +### Cosa significa quel comando? + +| Parte | Spiegazione | +|-------|-------------| +| `docker run` | Avvia un container | +| `--rm` | Cancella il container quando chiudi | +| `-it` | Modalità interattiva (puoi scrivere e leggere) | +| `-v "PERCORSO:/app/glm-5.2"` | Collega il tuo modello dentro il container | +| `-e COLI_MODEL=/app/glm-5.2` | Dice a Colibrì dove trovare il modello | +| `colibri-i` | Nome dell'immagine Docker | +| `./coli chat` | Avvia Colibrì in modalità chat | + +--- + +### Usare Colibrì + +Una volta avviato, vedrai un prompt come questo: + +``` + ────────────────────────────────────────────────────────── + type and press Enter · Ctrl-C stops the answer · :more continues · :reset clears memory · :q exits +``` + +**Comandi utili:** +- `Scrivi una domanda + Invio` → Ricevi la risposta +- `Ctrl + C` → Interrompi la risposta +- `:reset` → Cancella la memoria della conversazione +- `:q` → Esci + +**Esempio di uso:** +``` +› Quanti abitanti ha la Cina? + +La Cina è attualmente il paese più popoloso al mondo. +La popolazione è di circa 1,41 miliardi di abitanti. +``` + +Il modello capisce **italiano, inglese, cinese e altre lingue**, anche se è ottimizzato per inglese e cinese. + +--- + +## Entrare in una console Linux dentro il container + +Se vuoi esplorare il container come fosse una macchina Linux normale: + +```bash +docker run --rm -it --name colibri-c \ + -v "PERCORSO_MODELLO:/app/glm-5.2" \ + -e COLI_MODEL=/app/glm-5.2 \ + colibri-i /bin/bash +``` + +Ora sei dentro Linux. Digita `exit` per uscire. + +--- + +## Risoluzione dei problemi + +### ❌ "Docker non trovato" + +**Causa**: Docker non è installato o il terminale non lo riconosce. + +**Soluzione**: +1. Reinstalla [Docker Desktop](https://www.docker.com/products/docker-desktop/) +2. Riavvia il computer +3. Apri un nuovo terminale e riprova + +--- + +### ❌ "Out of memory" (memoria insufficiente) o container che si chiude subito + +**Causa**: Il tuo computer non ha abbastanza RAM, oppure su Windows, WSL usa meno memoria di quella disponibile. + +**Soluzione per Windows (WSL):** + +1. Apri PowerShell e controlla la memoria disponibile a WSL: + ```bash + wsl + cat /proc/meminfo | grep MemTotal + exit + ``` + Dividi il numero per 1.073.741.824 (è 1024³) per averlo in GB. + +2. Se WSL usa meno di quello che hai, crea un file di configurazione: + - Apri un editor di testo (Notepad va bene) + - Copia questo: + ```ini + [wsl2] + memory=24GB + processors=12 + swap=16GB + ``` + - Salva il file con il nome: `.wslconfig` (con il punto) + - Posizionalo in: `C:\Users\TuoNomeUtente\` + +3. Riavvia WSL da PowerShell: + ```bash + wsl --shutdown + wsl + ``` + +4. Controlla di nuovo: + ```bash + # cat /proc/meminfo | grep MemTotal + # exit + ``` + +**Soluzione per Mac/Linux**: Aumenta la RAM disponibile a Docker dalle impostazioni di Docker Desktop, oppure aggiungi più RAM al computer. + +--- + +### ❌ La risposta è molto lenta + +**Cause possibili**: +1. Il disco è lento +2. Hai poca RAM +3. Colibrì sta usando il disco come memoria aggiuntiva (normale) + +**Come controllare la velocità del disco:** + +**Windows** (PowerShell da amministratore): +```bash +winsat disk -drive C +``` +Cambia `C` con la lettera del tuo disco. + +**Linux/Mac** (Terminal): +```bash +sudo hdparm -Tt /dev/sda +``` +Cambia `/dev/sda` con il tuo disco (vedi con `lsblk` per Linux). + +Un **SSD NVMe moderno** arriva a 15 GB/sec. Se il tuo è sotto 2-3 GB/sec, è lento. + +--- + +### ❌ "Permission denied" su Linux + +**Causa**: Docker richiede permessi da amministratore. + +**Soluzione - Opzione 1** (rapida): +```bash +sudo docker build -t colibri-i . +sudo docker run ... (come sopra, con sudo davanti) +``` + +**Soluzione - Opzione 2** (permanente): +```bash +sudo usermod -aG docker $USER +# Riavvia il computer +docker run ... (senza sudo) +``` + +--- + +### ❌ "Image not found" o errore durante il build + +**Causa**: Il Dockerfile è corrotto o non nella cartella giusta. + +**Soluzione**: +1. Verifica che il Dockerfile sia nella cartella dove apri il terminale: + ```bash + ls Dockerfile # Mac/Linux + dir Dockerfile # Windows + ``` +2. Riscarica il Dockerfile dal repository GitHub +3. Elimina l'immagine vecchia: `docker rmi colibri-i` +4. Riprova il build + +--- + +### ❌ "hf_download: command not found" + +**Causa**: La libreria Hugging Face non è installata correttamente. + +**Soluzione**: +```bash +pip install -U huggingface_hub[cli] +# oppure su Linux/Mac: +pip3 install -U huggingface_hub[cli] +``` + +Poi riprova il comando `hf_download`. + +--- + +### ❌ Il modello non si scarica (timeout o errori di rete) + +**Cause**: Connessione lenta o instabile, Hugging Face temporaneamente non disponibile. + +**Soluzione**: +1. Attendi e riprova il comando `hf_download` +2. Se continua, scarica manualmente da [qui](https://huggingface.co/mateogrgic/GLM-5.2-colibri-int4-with-int8-mtp) +3. Decomprimi il file ZIP nella cartella desiderata + +--- + +## Note tecniche + +### Perché il disco è importante? + +Colibrì usa il disco come "RAM aggiuntiva" virtuale (paging). Un disco **veloce** è cruciale per prestazioni decenti. + +- **SSD NVMe** (consigliato): 1-15 GB/sec +- **SSD SATA**: 0.5-1 GB/sec +- **Hard disk meccanico**: 0.05-0.1 GB/sec ❌ (troppo lento) + +Se il tuo disco è lento, le risposte saranno molto lente anche con molta RAM. + +--- + +### Configurazione di default consigliata per WLS su Windows + +Se hai **esattamente 32 GB di RAM** e usi Windows, è molto probabile che WLS di default sia settato per non consumare più di 16 GB di RAM. Bisogna aumentare questo limite [Risoluzione dei problemi](#risoluzione-dei-problemi) . Nel mio caso ho adottato questa configurazione: +```ini +[wsl2] +memory=24GB +processors=12 +swap=16GB +``` + +Ovvero, nel mio caso, ho lasciato 8 GB di RAM e 4 CPU a Windows e dato 24 GB e 12 processori a WSL + Linux. + +--- + +## Domande frequenti + +**D: E se ho meno di 32 GB di RAM?** +R: Probabile che non funzioni bene. Puoi provare se hai 24 GB, ma non è garantito. + +**D: Posso aumentare la velocità di risposta?** +R: Sì, in parte: +- Usa un SSD NVMe veloce +- Aumenta la RAM +- Riduci la complessità delle domande +- Usa `:reset` per cancellare la memoria e alleggerire il carico + +**D: Posso usare Colibrì senza Docker?** +R: Colibrì è nato così, ma questa guida assume Docker. Per compilare da sorgente, vedi il repository GitHub. + +**D: Quanta connessione internet mi serve dopo aver scaricato il modello?** +R: Zero. Colibrì funziona completamente offline. + +--- + +## Supporto e contributi + +Se trovi errori o hai suggerimenti per migliorare questa guida, aprici una issue o una pull request sul repository GitHub di Colibrì. + +Buon divertimento! 🐦 + +--- + +## Test sul mio PC + +Nel primo caso ho fatto una domanda in italiano, nel secondo in giapponese, e nel terzo ho rifatto la domanda in giapponese ma ho richiesto una risposta in italiano. + +``` +PS C:\quack\llm\colibri\docker> docker run --rm -it --name colibri-c -v "C:\quack\llm\models\glm-5.2:/app/glm-5.2" -e COLI_MODEL=/app/glm-5.2 colibri-i ./coli chat + + ▄▀▀▀▄ ▄ colibrì v1.0 + ▄▄▄▄▀▀▀▀▄▀▀ tiny engine, immense model + ▀▀▀▀▀▀▀ GLM-5.2 · 744B MoE · int4 · streaming CPU + ▀▀▀▀ chat · glm-5.2 · ram -GB · topp off + ▀ + ────────────────────────────────────────────────────────── + type and press Enter · Ctrl-C stops the answer · :more continues · :reset clears memory · :q exits + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › Quanti abitanti ha la Cina? │ + ╰────────────────────────────────────────────────────────────────────────────────────────────────╯ + + ◆ colibrì + La Cina è attualmente il paese più popoloso al mondo (sebbene, secondo alcune stime recenti, sia stata ormai superata dall'India). + + La popolazione totale della Repubblica Popolare Cinese è di circa 1,41 miliardi di abitanti (dati del 2020-2022 circa). + └─ 76 tok · 0.04 tok/s · hit 3% · RSS 15.9 GB · 2012s + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › 漫画「ワンピース」の主人公の名前を教えてください。名前だけで、それ以上のコメントはありません。 │ + ╰────────────────────────────────────────────────────────────────────────────────────────────────╯ + + ◆ colibrì + ルフィ + └─ 2 tok · 0.01 tok/s · hit 1% · RSS 16.7 GB · 260s + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › 漫画「ワンピース」の主人公の名前を教えてください。名前だけで、それ以上のコメントはありません。イタリア語で返信 │ + ╰────────────────────────────────────────────────────────────────────────────────────────────────╯ + + ◆ colibrì + Il nome del protagonista di One Piece è Monkey D. Luffy. + └─ 14 tok · 0.02 tok/s · hit 2% · RSS 17.3 GB · 593s + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › + ``` \ No newline at end of file diff --git a/docker/README.md b/docker/README.md new file mode 100644 index 0000000..44d8a19 --- /dev/null +++ b/docker/README.md @@ -0,0 +1,488 @@ +_Leggi il leggimi in [Italiano](README.IT.md)._ + + +# Colibrì Guide - Local Inference Engine + +A simple guide to running **Colibrì**, a local inference engine based on GLM 5.2, without needing programming knowledge. If you already have Docker installed, you are well on your way. + +--- + +## 📋 Table of Contents + +* [What is Colibrì?](https://www.google.com/search?q=#what-is-colibr%C3%AC) +* [Requirements](https://www.google.com/search?q=#requirements) + * [Hardware](https://www.google.com/search?q=#hardware) + * [Software](https://www.google.com/search?q=#software) +* [How to get started](https://www.google.com/search?q=#how-to-get-started) + * [Step 1: Download the model](https://www.google.com/search?q=#step-1-download-the-model) + * [Step 2: Download the Colibrì Dockerfile](https://www.google.com/search?q=#step-2-download-the-colibr%C3%AC-dockerfile) + * [Step 3: Build the Docker image](https://www.google.com/search?q=#step-3-build-the-docker-image) + * [Step 4: Start Colibrì](https://www.google.com/search?q=#step-4-start-colibr%C3%AC) + * [What does that command mean?](https://www.google.com/search?q=#what-does-that-command-mean) + * [Using Colibrì](https://www.google.com/search?q=#using-colibr%C3%AC) +* [Entering a Linux console inside the container](https://www.google.com/search?q=#entering-a-linux-console-inside-the-container) +* [Troubleshooting](https://www.google.com/search?q=#troubleshooting) +* [Technical notes](https://www.google.com/search?q=#technical-notes) +* [Frequently Asked Questions](https://www.google.com/search?q=#frequently-asked-questions) +* [Support and contributions](https://www.google.com/search?q=#support-and-contributions) +* [Tests on my PC](https://www.google.com/search?q=#tests-on-my-pc) + +--- + +## What is Colibrì? + +Colibrì is an application that allows you to run an artificial intelligence model (GLM 5.2) directly on your computer, without connecting to external servers. It is also possible to run it in Docker, which isolates the application from the rest of the system. + +> **Important note**: The model is very large. Expect to wait several minutes for an answer to a simple question, especially with low RAM. At the end of this readme, you will see the result on my PC (without a discrete graphics card), and I reach 0.01 tokens per second. + +--- + +## Requirements + +### Hardware + +| RAM Memory | Works? | Notes | +| --- | --- | --- | +| < 16 GB | ❌ No | Insufficient memory | +| 24 GB | ⚠️ Maybe | Possible, needs testing | +| 32 GB | ✅ Yes | The minimum (but see memory section for Windows) | +| 48+ GB | ✅ Yes | Better | + +Additionally: a **fast SSD** is essential. Colibrì uses the disk as additional memory. Having an NVidia graphics card is even better. + +### Software + +* **Docker Desktop** (Windows, Mac, Linux) — [download here](https://www.google.com/search?q=https://www.docker.com/products/docker-desktop/) +* **Python** (only if you want to download the model yourself) +* Windows: [python.org](https://www.google.com/search?q=https://www.python.org) or Microsoft Store +* Linux: `apt-get install python3 python3-pip` +* Mac: [python.org](https://www.google.com/search?q=https://www.python.org) or Homebrew + +No build environment is needed. Everything happens inside the Docker container. + +--- + +## How to get started + +### Step 1: Download the model + +The GLM 5.2 model is approximately **360 GB**. Choose one of these methods: + +#### Method A: Using Python (recommended) + +1. **Install the library for Hugging Face:** +```bash +python -m pip install -U huggingface_hub[cli] +``` + +On Linux, use `python3` instead of `python`. +2. **Download the model** (open the terminal in the folder where you want to save it): +```bash +hf_download mateogrgic/GLM-5.2-colibri-int4-with-int8-mtp --local-dir . +``` + +**Example**: if you want to save it in `C:\LLM\models\glm-5.2` (Windows): +* Open PowerShell in that folder +* Copy and paste the command above +* Wait (a long time) + +#### Method B: Without Python (only if necessary) + +If you are on Windows and cannot get it to work with Python: + +* Download manually from [Hugging Face](https://www.google.com/search?q=https://huggingface.co/mateogrgic/GLM-5.2-colibri-int4-with-int8-mtp) +* Unzip into a folder (e.g., `C:\LLM\models\glm-5.2`) + +--- + +### Step 2: Download the Colibrì Dockerfile + +1. Go to: [https://github.com/JustVugg/colibri/blob/main/docker/Dockerfile](https://www.google.com/search?q=https://github.com/JustVugg/colibri/blob/main/docker/Dockerfile) +2. Click the **Download** button (⬇️ icon) in the top right +3. Save the file in a folder (e.g., `C:\LLM\Colibrì`) + +--- + +### Step 3: Build the Docker image + +Open the terminal (PowerShell on Windows, Terminal on Mac/Linux) **in the folder where you saved the Dockerfile** and type: + +**Windows:** + +```bash +docker build -t colibri-i . +``` + +**Linux/Mac:** + +```bash +sudo docker build -t colibri-i . +``` + +Wait for it to finish (a few minutes). If everything goes well, you will see: `Successfully tagged colibri-i:latest` + +> **If you want to receive repository updates**: First delete the old image with `docker rmi colibri-i` and rebuild. + +--- + +### Step 4: Start Colibrì + +Open the terminal and type the command below (replace `C:\LLM\models\glm-5.2` with the actual path on your PC): + +**Windows** (PowerShell): + +```bash +$MODEL_PATH="C:\LLM\models\glm-5.2" +docker run --rm -it --name colibri-c ` + -v "$MODEL_PATH`:/app/glm-5.2" ` + -e COLI_MODEL=/app/glm-5.2 ` + colibri-i ./coli chat +``` + +**Mac/Linux** (Terminal/Bash): + +```bash +MODEL_PATH="/path/to/glm-5.2" +docker run --rm -it --name colibri-c \ + -v "$MODEL_PATH:/app/glm-5.2" \ + -e COLI_MODEL=/app/glm-5.2 \ + colibri-i ./coli chat +``` + +**Example for Linux:** + +```bash +MODEL_PATH="/home/user/LLM/glm-5.2" +docker run --rm -it --name colibri-c \ + -v "$MODEL_PATH:/app/glm-5.2" \ + -e COLI_MODEL=/app/glm-5.2 \ + colibri-i ./coli chat +``` + +--- + +### What does that command mean? + +| Part | Explanation | +| --- | --- | +| `docker run` | Starts a container | +| `--rm` | Deletes the container when you close it | +| `-it` | Interactive mode (you can write and read) | +| `-v "PATH:/app/glm-5.2"` | Mounts your model inside the container | +| `-e COLI_MODEL=/app/glm-5.2` | Tells Colibrì where to find the model | +| `colibri-i` | Name of the Docker image | +| `./coli chat` | Starts Colibrì in chat mode | + +--- + +### Using Colibrì + +Once started, you will see a prompt like this: + +``` + ────────────────────────────────────────────────────────── + type and press Enter · Ctrl-C stops the answer · :more continues · :reset clears memory · :q exits + +``` + +**Useful commands:** + +* `Write a question + Enter` → Receive the answer +* `Ctrl + C` → Stop the answer +* `:reset` → Clear conversation memory +* `:q` → Exit + +**Usage example:** + +``` +› How many inhabitants does China have? + +China is currently the most populous country in the world. +The population is approximately 1.41 billion people. + +``` + +The model understands **Italian, English, Chinese, and other languages**, although it is optimized for English and Chinese. + +--- + +## Entering a Linux console inside the container + +If you want to explore the container as if it were a normal Linux machine: + +```bash +docker run --rm -it --name colibri-c \ + -v "MODEL_PATH:/app/glm-5.2" \ + -e COLI_MODEL=/app/glm-5.2 \ + colibri-i /bin/bash +``` + +Now you are inside Linux. Type `exit` to leave. + +--- + +## Troubleshooting + +### ❌ "Docker not found" + +**Cause**: Docker is not installed or the terminal does not recognize it. + +**Solution**: + +1. Reinstall [Docker Desktop](https://www.google.com/search?q=https://www.docker.com/products/docker-desktop/) +2. Restart your computer +3. Open a new terminal and try again + +--- + +### ❌ "Out of memory" or container closes immediately + +**Cause**: Your computer does not have enough RAM, or on Windows, WSL is using less memory than available. + +**Solution for Windows (WSL):** + +1. Open PowerShell and check the memory available to WSL: +```bash +wsl +cat /proc/meminfo | grep MemTotal +exit +``` + + +Divide the number by 1,073,741,824 (which is 1024³) to get it in GB. +2. If WSL uses less than what you have, create a configuration file: +* Open a text editor (Notepad is fine) +* Copy this: +```ini +[wsl2] +memory=24GB +processors=12 +swap=16GB + +``` + +* Save the file with the name: `.wslconfig` (with the dot) +* Place it in: `C:\Users\YourUsername\` + +3. Restart WSL from PowerShell: +```bash +wsl --shutdown +wsl +``` + +4. Check again: +```bash +# cat /proc/meminfo | grep MemTotal +# exit +``` + +**Solution for Mac/Linux**: Increase the RAM available to Docker from the Docker Desktop settings, or add more RAM to the computer. + +--- + +### ❌ The answer is very slow + +**Possible causes**: + +1. The disk is slow +2. You have low RAM +3. Colibrì is using the disk as additional memory (normal) + +**How to check disk speed:** + +**Windows** (PowerShell as administrator): + +```bash +winsat disk -drive C +``` + +Change `C` with your disk letter. + +**Linux/Mac** (Terminal): + +```bash +sudo hdparm -Tt /dev/sda +``` + +Change `/dev/sda` with your disk (check with `lsblk` on Linux). + +A **modern NVMe SSD** reaches 15 GB/sec. If yours is under 2-3 GB/sec, it is slow. + +--- + +### ❌ "Permission denied" on Linux + +**Cause**: Docker requires administrator permissions. + +**Solution - Option 1** (quick): + +```bash +sudo docker build -t colibri-i . +sudo docker run ... (as above, with sudo in front) +``` + +**Solution - Option 2** (permanent): + +```bash +sudo usermod -aG docker $USER +# Restart the computer +docker run ... (without sudo) +``` +--- + +### ❌ "Image not found" or error during build + +**Cause**: The Dockerfile is corrupted or not in the right folder. + +**Solution**: + +1. Verify that the Dockerfile is in the folder where you open the terminal: +```bash +ls Dockerfile # Mac/Linux +dir Dockerfile # Windows +``` + +2. Redownload the Dockerfile from the GitHub repository +3. Delete the old image: `docker rmi colibri-i` +4. Retry the build + +--- + +### ❌ "hf_download: command not found" + +**Cause**: The Hugging Face library is not installed correctly. + +**Solution**: + +```bash +pip install -U huggingface_hub[cli] +# or on Linux/Mac: +pip3 install -U huggingface_hub[cli] +``` + +Then retry the `hf_download` command. + +--- + +### ❌ The model does not download (timeout or network errors) + +**Causes**: Slow or unstable connection, Hugging Face temporarily unavailable. + +**Solution**: + +1. Wait and retry the `hf_download` command +2. If it continues, download manually from [here](https://www.google.com/search?q=https://huggingface.co/mateogrgic/GLM-5.2-colibri-int4-with-int8-mtp) +3. Unzip the ZIP file into the desired folder + +--- + +## Technical notes + +### Why is the disk important? + +Colibrì uses the disk as "additional virtual RAM" (paging). A **fast** disk is crucial for decent performance. + +* **NVMe SSD** (recommended): 1-15 GB/sec +* **SATA SSD**: 0.5-1 GB/sec +* **Rotational hard disk**: 0.05-0.1 GB/sec ❌ (too slow) + +If your disk is slow, the answers will be very slow even with a lot of RAM. + +--- + +### Recommended default configuration for WLS on Windows + +If you have **exactly 32 GB of RAM** and are using Windows, it is very likely that WLS by default is set to consume no more than 16 GB of RAM. We need to increase this limit [Troubleshooting](https://www.google.com/search?q=#troubleshooting) . In my case I adopted this configuration: + +```ini +[wsl2] +memory=24GB +processors=12 +swap=16GB + +``` + +That is, in my case, I left 8 GB of RAM and 4 CPUs to Windows and gave 24 GB and 12 processors to WSL + Linux. + +--- + +## Frequently Asked Questions + +**Q: What if I have less than 32 GB of RAM?** + +A: It is likely that it will not work well. You can try if you have 24 GB, but it is not guaranteed. + +**Q: Can I increase the response speed?** + +A: Yes, partly: + +* Use a fast NVMe SSD +* Increase RAM +* Reduce the complexity of questions +* Use `:reset` to clear memory and lighten the load + +**Q: Can I use Colibrì without Docker?** + +A: Colibrì was born that way, but this guide assumes Docker. To build from source, see the GitHub repository. + +**Q: How much internet connection do I need after downloading the model?** + +A: Zero. Colibrì works completely offline. + +--- + +## Support and contributions + +If you find errors or have suggestions for improving this guide, open an issue or a pull request on the Colibrì GitHub repository. + +Have fun! 🐦 + +--- + +## Tests on my PC + +In the first case, I asked a question in Italian; in the second, in Japanese; and in the third, I repeated the question in Japanese but requested an answer in Italian. + +``` +PS C:\quack\llm\colibri\docker> docker run --rm -it --name colibri-c -v "C:\quack\llm\models\glm-5.2:/app/glm-5.2" -e COLI_MODEL=/app/glm-5.2 colibri-i ./coli chat + + ▄▀▀▀▄ ▄ colibrì v1.0 + ▄▄▄▄▀▀▀▀▄▀▀ tiny engine, immense model + ▀▀▀▀▀▀▀ GLM-5.2 · 744B MoE · int4 · streaming CPU + ▀▀▀▀ chat · glm-5.2 · ram -GB · topp off + ▀ + ────────────────────────────────────────────────────────── + type and press Enter · Ctrl-C stops the answer · :more continues · :reset clears memory · :q exits + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › Quanti abitanti ha la Cina? │ + ╰────────────────────────────────────────────────────────────────────────────────────────────────╯ + + ◆ colibrì + La Cina è attualmente il paese più popoloso al mondo (sebbene, secondo alcune stime recenti, sia stata ormai superata dall'India). + + La popolazione totale della Repubblica Popolare Cinese è di circa 1,41 miliardi di abitanti (dati del 2020-2022 circa). + └─ 76 tok · 0.04 tok/s · hit 3% · RSS 15.9 GB · 2012s + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › 漫画「ワンピース」の主人公の名前を教えてください。名前だけで、それ以上のコメントはありません。 │ + ╰────────────────────────────────────────────────────────────────────────────────────────────────╯ + + ◆ colibrì + ルフィ + └─ 2 tok · 0.01 tok/s · hit 1% · RSS 16.7 GB · 260s + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › 漫画「ワンピース」の主人公の名前を教えてください。名前だけで、それ以上のコメントはありません。イタリア語で返信 │ + ╰────────────────────────────────────────────────────────────────────────────────────────────────╯ + + ◆ colibrì + Il nome del protagonista di One Piece è Monkey D. Luffy. + └─ 14 tok · 0.02 tok/s · hit 2% · RSS 17.3 GB · 593s + + ╭────────────────────────────────────────────────────────────────────────────────────────────────╮ + │ › + +``` + +[source: 1] \ No newline at end of file From 4586d33c605a6ea6580b9d9b94b71d72037d0ba7 Mon Sep 17 00:00:00 2001 From: Nicholas Beerbower Date: Sun, 19 Jul 2026 10:54:01 -0400 Subject: [PATCH 18/71] fix: MTP-head probe uses the last expert by index, not a hardcoded 255 The has_mtp completeness probe checked for `mlp.experts.255.down_proj.weight`, which only exists when n_routed_experts == 256. REAP-pruned checkpoints (and any MoE with a different expert count) have fewer experts, so the probe spuriously reported has_mtp=0 and disabled MTP speculative decode even though the head was present and complete. Probe `mlp.experts.` instead. Co-Authored-By: Claude Opus 4.8 --- c/colibri.c | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/c/colibri.c b/c/colibri.c index 03eb878..9a81daa 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -860,12 +860,17 @@ static void model_init(Model *m, const char *snap, int cap, int ebits, int dbits "self_attn.q_a_proj.weight","self_attn.q_b_proj.weight","self_attn.kv_a_proj_with_mqa.weight", "self_attn.kv_b_proj.weight","self_attn.o_proj.weight","mlp.gate.weight", "mlp.shared_experts.gate_proj.weight","mlp.shared_experts.down_proj.weight", - "mlp.experts.0.gate_proj.weight","mlp.experts.255.down_proj.weight"}; + "mlp.experts.0.gate_proj.weight"}; char mn[256]; m->has_mtp=1; for(unsigned q=0;qn_layers,req[q]); if(!st_has(&m->S,mn)){ m->has_mtp=0; break; } } + /* probe the LAST expert by index, not a fixed 255: REAP-pruned + * checkpoints have n_routed_experts < 256 and the MTP set stays complete, + * so a hardcoded expert.255 would spuriously report has_mtp=0 on them. */ + snprintf(mn,sizeof(mn),"model.layers.%d.mlp.experts.%d.down_proj.weight",c->n_layers,c->n_experts-1); + if(!st_has(&m->S,mn)) m->has_mtp=0; if(getenv("MTP") && atoi(getenv("MTP"))==0) m->has_mtp=0; if(m->has_mtp){ int i=c->n_layers; Layer *l=&m->mtpL; From ab55f4900c0696b878a179c1733218e974ceccf2 Mon Sep 17 00:00:00 2001 From: ZacharyZcR Date: Fri, 17 Jul 2026 15:32:09 +0800 Subject: [PATCH 19/71] =?UTF-8?q?cuda:=20COLI=5FGROUP=5FASYNC=3D1=20?= =?UTF-8?q?=E2=80=94=20async=20expert-group=20issue/take=20with=20CPU/GPU?= =?UTF-8?q?=20overlap=20at=20decode=20(opt-in,=20+6-8%=20measured)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- c/backend_cuda.cu | 69 ++++++++++++++++++++++++ c/backend_cuda.h | 10 ++++ c/colibri.c | 134 ++++++++++++++++++++++++++++++++++++++++++++++ 3 files changed, 213 insertions(+) diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu index 413e53b..49e60bc 100644 --- a/c/backend_cuda.cu +++ b/c/backend_cuda.cu @@ -38,6 +38,7 @@ typedef struct { cudaStream_t stream; void *group_desc; size_t group_desc_cap; size_t tensor_count, tensor_bytes; + int group_pending; size_t group_pending_bytes; /* async expert-group in flight (Inc.4) */ } DeviceContext; typedef struct { @@ -757,6 +758,74 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates, return 1; } +/* ---- Async expert group (Inc.4): issue/take split of coli_cuda_expert_group ---- + * The measured cost of the sync call at decode is ~0.45 ms/call of HOST-side wait + * (stream sync + staging), vs ~0.18 ms of actual GPU work — 70% tax, paid ~5x per + * layer because a token's 8 experts scatter across devices. issue() stages and + * launches on the device stream and returns immediately; take() syncs and hands + * back the pinned result rows. One issue may be outstanding per device; moe() + * takes at each layer end, which also orders the next layer's reuse of the ctx + * scratch buffers. Small batches only (decode/spec): bigger totals keep the sync + * path with its TC variants. Numerics are the sync path's small-batch kernels, + * so greedy output is byte-identical by construction. */ +extern "C" int coli_cuda_expert_group_issue(ColiCudaTensor *const *gates, + ColiCudaTensor *const *ups, + ColiCudaTensor *const *downs, + const int *rows, int count, + const float *x) { + if (!gates || !ups || !downs || !rows || !x || count < 1 || count > 64) return 0; + ColiCudaTensor *first=gates[0]; + if (!first) return 0; + int device=first->device,D=first->I,I=first->O,total=0; + GroupDesc host[64]; + for(int c=0;cdevice!=device||u->device!=device||d->device!=device|| + g->I!=D||u->I!=D||g->O!=I||u->O!=I||d->I!=I||d->O!=D) return 0; + host[c]={g->weights,u->weights,d->weights,g->scales,u->scales,d->scales, + g->fmt,u->fmt,d->fmt,rows[c],total}; + total+=rows[c]; + } + if(total>8) return 0; /* decode-scale only */ + DeviceContext *ctx=find_ctx(device); if(!ctx||ctx->group_pending||!select_ctx(ctx)) return 0; + size_t xb=(size_t)total*D*sizeof(float), ib=(size_t)total*I*sizeof(float); + if(!reserve(&ctx->x,&ctx->x_cap,xb)||!reserve(&ctx->y,&ctx->y_cap,xb)|| + !reserve(&ctx->gate,&ctx->gate_cap,ib)||!reserve(&ctx->up,&ctx->up_cap,ib)|| + !reserve_pinned(&ctx->host_x,&ctx->host_x_cap,xb)|| + !reserve_pinned(&ctx->host_y,&ctx->host_y_cap,xb)) return 0; + std::memcpy(ctx->host_x,x,xb); + if(!cuda_ok(cudaMemcpyAsync(ctx->x,ctx->host_x,xb,cudaMemcpyHostToDevice,ctx->stream), + "expert group issue upload")) return 0; + for(int c=0;cgate+(size_t)host[c].offset*I,*u16=ctx->up+(size_t)host[c].offset*I; + float *x16=ctx->x+(size_t)host[c].offset*D,*y16=ctx->y+(size_t)host[c].offset*D; + quant_matmul<<stream>>>(g16,x16, + host[c].g,host[c].gs,host[c].gf,r,D,I,row_bytes(host[c].gf,D)); + quant_matmul<<stream>>>(u16,x16, + host[c].u,host[c].us,host[c].uf,r,D,I,row_bytes(host[c].uf,D)); + silu_mul<<<(unsigned)(((size_t)r*I+255)/256),256,0,ctx->stream>>>(g16,u16,(size_t)r*I); + quant_matmul<<stream>>>(y16,g16, + host[c].d,host[c].ds,host[c].df,r,I,D,row_bytes(host[c].df,I)); + } + if(!cuda_ok(cudaGetLastError(),"expert group issue launch")|| + !cuda_ok(cudaMemcpyAsync(ctx->host_y,ctx->y,xb,cudaMemcpyDeviceToHost,ctx->stream), + "expert group issue download")) return 0; + ctx->group_pending=1; ctx->group_pending_bytes=xb; + { std::lock_guard lock(g_group_stats_mu); + g_group_calls++; g_group_experts+=(uint64_t)count; g_group_rows+=(uint64_t)total; } + return 1; +} + +extern "C" const float *coli_cuda_expert_group_take(int device) { + DeviceContext *ctx=find_ctx(device); + if(!ctx||!ctx->group_pending) return nullptr; + ctx->group_pending=0; + if(!select_ctx(ctx)) return nullptr; + if(!cuda_ok(cudaStreamSynchronize(ctx->stream),"expert group take")) return nullptr; + return ctx->host_y; +} + extern "C" int coli_cuda_attention_absorb(ColiCudaTensor *w,float *ctx,const float *q, const float *latent,const float *rope,int H,int Q, diff --git a/c/backend_cuda.h b/c/backend_cuda.h index 63c1f11..bceb0bc 100644 --- a/c/backend_cuda.h +++ b/c/backend_cuda.h @@ -67,6 +67,16 @@ COLI_CUDA_DLLEXPORT int coli_cuda_shared_mlp_w4a16(ColiCudaTensor *gate, ColiCud /* Packed group of same-shaped experts. Inputs and outputs contain sum(rows) * consecutive [D] rows in call order. */ +/* Async issue/take split of the group call below (Inc.4): issue launches on the + * device stream and returns; take syncs and returns the pinned result rows (valid + * until the next issue on that device). Small totals only (<=8 rows); one + * outstanding issue per device. */ +COLI_CUDA_DLLEXPORT int coli_cuda_expert_group_issue(ColiCudaTensor *const *gates, + ColiCudaTensor *const *ups, + ColiCudaTensor *const *downs, + const int *rows, int count, const float *x); +COLI_CUDA_DLLEXPORT const float *coli_cuda_expert_group_take(int device); + COLI_CUDA_DLLEXPORT int coli_cuda_expert_group(ColiCudaTensor *const *gates, ColiCudaTensor *const *ups, ColiCudaTensor *const *downs, diff --git a/c/colibri.c b/c/colibri.c index 03eb878..33ce4e4 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2516,6 +2516,11 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int } #ifdef COLI_CUDA ESlot *group_e[64]; int group_n[64]; int ngroup=0; + /* Inc.4 overlap stash: pass-1 packing kept for the take phase after the CPU loop */ + ESlot *eg_e[64]; int eg_n[64], eg_row[64][4], eg_npg=0; float eg_w[64][4]; + int dev_nc0[COLI_CUDA_MAX_DEVICES], dev_off0[COLI_CUDA_MAX_DEVICES], + dev_total0[COLI_CUDA_MAX_DEVICES], dev_which0[COLI_CUDA_MAX_DEVICES][64]; + memset(dev_nc0,0,sizeof(dev_nc0)); (void)eg_npg; (void)dev_total0; (void)dev_off0; #endif #ifdef COLI_METAL if(g_metal_enabled){ @@ -2544,9 +2549,82 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int free(mxg); free(mrows); free(mrw); } #undef MB_BUILD +#endif +#ifdef COLI_CUDA + /* Inc.4 pass 1: collect the VRAM-resident experts' groups and ISSUE them async + * BEFORE the CPU loop below, so the GPU computes its share while the CPU works + * through the RAM-tier/miss rows — t_emm becomes max(cpu, gpu) instead of the + * sum. Only resident experts are collected (misses are never cuda_eligible), so + * no pipe_wait is needed here; the CPU loop keeps its own waits. Any issue + * failure drops the layer back to the collect-in-loop + sync-group path. */ + int early_issued=0, done_j[64]={0}; + { + static int g_group_async2=-1; + if(g_group_async2<0) g_group_async2=getenv("COLI_GROUP_ASYNC")?atoi(getenv("COLI_GROUP_ASYNC")):0; + if(!metal_done && g_group_async2 && group_enabled && S<=4 && g_cuda_enabled && + g_cuda_ndev>0 && !omp_in_parallel()){ + ESlot *pg_e[64]; int pg_n[64], pg_j[64], npg=0; + int prow[64][4]; float pw[64][4]; + for(int j=0;jg.cuda_eligible&&e->u.cuda_eligible&&e->d.cuda_eligible)) continue; + int nr=0; + for(int s=0;sg.cuda_device==g_cuda_devices[di]) pd_total[di]+=pg_n[q]; + for(int di=1;dig.cuda_device==device){ + int nc=pd_nc[di]++; ESlot *e=pg_e[q]; + pd_g[di][nc]=e->g.cuda; pd_u[di][nc]=e->u.cuda; pd_d[di][nc]=e->d.cuda; + pd_rows[di][nc]=pg_n[q]; pd_which[di][nc]=q; + for(int r=0;rt_emm+=now_s()-tg0; + for(int q=0;qgpu_expert_calls++; + } + } else { + for(int di=0;di=1 routing invariant. @@ -2593,6 +2671,35 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int m->cpu_expert_rows+=(uint64_t)nr;} } #ifdef COLI_CUDA + /* Inc.4 take phase: the CPU loop above ran while the GPU computed the issued + * groups — collect them now. A failed device recomputes its experts on the CPU + * (expert_host_ensure reloads slabs released by CUDA_RELEASE_HOST). */ + if(early_issued){ + double tg1=now_s(); + for(int di=0;dig,&e->u,nr); + for(int64_t z=0;z<(int64_t)nr*I;z++) gg[z]=siluf(gg[z])*uu[z]; + matmul_qt(hh,gg,&e->d,nr); + for(int r=0;rt_emm+=now_s()-tg1; + } ColiCudaTensor *dev_g[COLI_CUDA_MAX_DEVICES][64],*dev_u[COLI_CUDA_MAX_DEVICES][64]; ColiCudaTensor *dev_d[COLI_CUDA_MAX_DEVICES][64]; int dev_rows[COLI_CUDA_MAX_DEVICES][64],dev_which[COLI_CUDA_MAX_DEVICES][64]; @@ -2614,6 +2721,33 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int } } double tg=now_s(); + /* Inc.4: at decode scale, issue every device's group WITHOUT syncing, then take + * them all — one stream sync per device per layer instead of a full staged + * round-trip per call (measured: ~70% of the sync call is host-side wait). + * Any issue failure drains what was issued and the whole layer falls back to + * the sync path below, which recomputes from group_x (idempotent). */ + int async_done=0; + static int g_group_async=-1; + if(g_group_async<0) g_group_async=getenv("COLI_GROUP_ASYNC")?atoi(getenv("COLI_GROUP_ASYNC")):0; + if(g_group_async && S<=4 && g_cuda_ndev>0){ + int issued[COLI_CUDA_MAX_DEVICES]={0}, all=1; + for(int di=0;di1) schedule(static) for(int di=0;di Date: Fri, 17 Jul 2026 15:41:38 +0800 Subject: [PATCH 20/71] cuda: cache layernorm weights on-device in pipe_layer_sparse (kill 152 sync H2D/token) + overlap-window profiling counters --- c/colibri.c | 32 +++++++++++++++++++++++++------- 1 file changed, 25 insertions(+), 7 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index 33ce4e4..af99ff9 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -163,6 +163,7 @@ typedef struct { KVState *kv; ESlot **ecache; int *ecn; int ecap; /* LRU expert per-layer */ float **kv_dev_L, **kv_dev_R; int *kv_dev_valid; /* ombra KV su device (decode) */ + float **ln_dev; /* in_ln/post_ln cached on device: [layer*2+{0,1}] (Inc.4) */ ESlot ws[64]; /* working set del layer corrente (load paralleli) */ ESlot **pin; int *npin; /* HOT-STORE: expert pinnati in RAM (mai evicted) */ uint32_t **eusage; /* contatori persistenti (per STATS/PIN) */ @@ -242,6 +243,7 @@ static int qt_cuda_update(QT *t){ t->fmt==1?(const void*)t->q8:(const void*)t->q4; return coli_cuda_tensor_update(t->cuda,weights,t->s); } +static double g_ovl_issue,g_ovl_cpu,g_ovl_take,g_ovl_mark; /* Inc.4 overlap-window split (OVL report) */ static void cuda_stats_print(void){ size_t n=0,b=0; coli_cuda_stats(-1,&n,&b); fprintf(stderr,"[CUDA] resident set: %zu tensors, %.2f GB VRAM\n",n,b/1e9); @@ -257,6 +259,9 @@ static void cuda_stats_print(void){ getenv("COLI_CUDA_PROFILE")?"; timing sotto":""); if(calls&&getenv("COLI_CUDA_PROFILE")) fprintf(stderr, "[CUDA] expert groups timing: H2D %.1f ms | kernel %.1f ms | D2H %.1f ms\n",h2d,kernel,d2h); + if(g_ovl_issue+g_ovl_cpu+g_ovl_take>0) fprintf(stderr, + "[CUDA] overlap window: pack+issue %.2fs | cpu-rows %.2fs | take(sync+acc) %.2fs\n", + g_ovl_issue,g_ovl_cpu,g_ovl_take); } static int parse_cuda_devices(const char *list, int *out){ if(!list||!*list) return 0; @@ -2597,10 +2602,11 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int for(int di=0;dit_emm+=now_s()-tg1; + m->t_emm+=now_s()-tg1; g_ovl_take+=now_s()-tg1; } ColiCudaTensor *dev_g[COLI_CUDA_MAX_DEVICES][64],*dev_u[COLI_CUDA_MAX_DEVICES][64]; ColiCudaTensor *dev_d[COLI_CUDA_MAX_DEVICES][64]; @@ -3194,17 +3201,28 @@ static int pipe_layer_sparse(Model *m, Layer *l, int li, float *x_dev, int S, in if(!l->sh_gate.cuda_eligible||!l->sh_up.cuda_eligible||!l->sh_down.cuda_eligible|| !qt_cuda_upload(&l->sh_gate)||!qt_cuda_upload(&l->sh_up)||!qt_cuda_upload(&l->sh_down)|| l->sh_gate.cuda_device!=dev||l->sh_up.cuda_device!=dev||l->sh_down.cuda_device!=dev) return 0; - float *w_in =coli_cuda_pipe_scratch(dev,8,(size_t)D*4); - float *w_post=coli_cuda_pipe_scratch(dev,9,(size_t)D*4); + /* Inc.4: the layernorm weights are constants — upload once per layer and keep them + * on the layer's device, instead of two synchronous 24 KB uploads per layer per + * token (152 sync H2D/token measured on the profile). */ + if(!m->ln_dev) m->ln_dev=calloc((size_t)(c->n_layers+1)*2,sizeof(float*)); + float *w_in=m->ln_dev[(size_t)li*2], *w_post=m->ln_dev[(size_t)li*2+1]; + if(!w_in){ + w_in=coli_cuda_pipe_alloc(dev,(size_t)D*4); + if(!w_in||!coli_cuda_pipe_upload(dev,w_in,l->in_ln,(size_t)D*4)) return 0; + m->ln_dev[(size_t)li*2]=w_in; + } + if(!w_post){ + w_post=coli_cuda_pipe_alloc(dev,(size_t)D*4); + if(!w_post||!coli_cuda_pipe_upload(dev,w_post,l->post_ln,(size_t)D*4)) return 0; + m->ln_dev[(size_t)li*2+1]=w_post; + } float *nrm_d=coli_cuda_pipe_scratch(dev,10,xb); float *y_d =coli_cuda_pipe_scratch(dev,11,xb); float *sg_d =coli_cuda_pipe_scratch(dev,12,(size_t)S*sI*4); float *su_d =coli_cuda_pipe_scratch(dev,13,(size_t)S*sI*4); float *snap =coli_cuda_pipe_scratch(dev,14,xb); - if(!w_in||!w_post||!nrm_d||!y_d||!sg_d||!su_d||!snap) return 0; + if(!nrm_d||!y_d||!sg_d||!su_d||!snap) return 0; if(!coli_cuda_pipe_peer_copy(dev,snap,dev,x_dev,xb)) return 0; /* snapshot per il fallback */ - if(!coli_cuda_pipe_upload(dev,w_in,l->in_ln,(size_t)D*4)|| - !coli_cuda_pipe_upload(dev,w_post,l->post_ln,(size_t)D*4)) return 0; double ta=now_s(); if(!coli_cuda_pipe_rmsnorm(dev,nrm_d,x_dev,w_in,S,D,c->eps)) return 0; /* DSA: i layer con indexer FULL cachano k_idx dalla nrm pre-attention (CPU, piccolo) */ From 7d01d21023a0996623a8f831cd7ab361881aeca4 Mon Sep 17 00:00:00 2001 From: ZacharyZcR Date: Sat, 18 Jul 2026 18:44:27 +0800 Subject: [PATCH 21/71] fix Windows async expert loader --- c/backend_loader.c | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/c/backend_loader.c b/c/backend_loader.c index eedbd50..9a83498 100644 --- a/c/backend_loader.c +++ b/c/backend_loader.c @@ -41,6 +41,11 @@ typedef int (*fn_expert_mlp)(ColiCudaTensor *gate, ColiCudaTensor *up typedef int (*fn_expert_group)(ColiCudaTensor *const *gates, ColiCudaTensor *const *ups, ColiCudaTensor *const *downs, const int *rows, int count, float *y, const float *x); +typedef int (*fn_expert_group_issue)(ColiCudaTensor *const *gates, + ColiCudaTensor *const *ups, + ColiCudaTensor *const *downs, + const int *rows, int count, const float *x); +typedef const float * (*fn_expert_group_take)(int device); typedef int (*fn_attention_absorb)(ColiCudaTensor *kv_b, float *ctx, const float *q, const float *latent, const float *rope, int H, int Q, int R, int V, int K, int T, float attention_scale); @@ -100,6 +105,8 @@ static struct { fn_group_stats group_stats; fn_expert_mlp expert_mlp; fn_expert_group expert_group; + fn_expert_group_issue expert_group_issue; + fn_expert_group_take expert_group_take; fn_attention_absorb attention_absorb; fn_tensor_upload tensor_upload; fn_matmul matmul; @@ -194,6 +201,8 @@ static int coli_cuda_load(void){ RESOLVE(group_stats, fn_group_stats) RESOLVE(expert_mlp, fn_expert_mlp) RESOLVE(expert_group, fn_expert_group) + RESOLVE(expert_group_issue, fn_expert_group_issue) + RESOLVE(expert_group_take, fn_expert_group_take) RESOLVE(attention_absorb, fn_attention_absorb) RESOLVE(tensor_upload, fn_tensor_upload) RESOLVE(matmul, fn_matmul) @@ -289,6 +298,19 @@ int coli_cuda_expert_group(ColiCudaTensor *const *gates, ColiCudaTensor *const * return g_cuda.expert_group(gates, ups, downs, rows, count, y, x); } +int coli_cuda_expert_group_issue(ColiCudaTensor *const *gates, + ColiCudaTensor *const *ups, + ColiCudaTensor *const *downs, + const int *rows, int count, const float *x){ + if(!g_cuda.available) return 0; + return g_cuda.expert_group_issue(gates, ups, downs, rows, count, x); +} + +const float *coli_cuda_expert_group_take(int device){ + if(!g_cuda.available) return NULL; + return g_cuda.expert_group_take(device); +} + int coli_cuda_attention_absorb(ColiCudaTensor *kv_b, float *ctx, const float *q, const float *latent, const float *rope, int H, int Q, int R, int V, int K, int T, float attention_scale){ From 3e45074bf27bcb7fe75dfbdcba07f31da1848198 Mon Sep 17 00:00:00 2001 From: ZacharyZcR Date: Sun, 19 Jul 2026 23:31:37 +0800 Subject: [PATCH 22/71] =?UTF-8?q?site:=20official=20website=20=E2=80=94=20?= =?UTF-8?q?animated=20demo,=20expert=20brain,=203-D=20atlas,=20Pages=20dep?= =?UTF-8?q?loy?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A zero-build static site under site/, deployed to GitHub Pages by Actions: - hero with the pixel hummingbird, key numbers, CTA - 'watch it think': a chat replay paced at measured decode speeds (6x5090 / 128GB CPU / 5070 Ti / 25GB floor), with a live tok/s meter and the full 19,456-expert grid — colour = tier, brightness = heat, routed experts flash white per token - the expert atlas as a draggable 3-D galaxy (measured-affinity clusters) - three-tier explainer and the measured hardware ladder - single HTML file, no dependencies, no build step; custom domain later is just a site/CNAME + DNS --- .github/workflows/site.yml | 36 +++ site/colibri-icon.svg | 50 ++++ site/colibri.svg | 55 +++++ site/index.html | 456 +++++++++++++++++++++++++++++++++++++ 4 files changed, 597 insertions(+) create mode 100644 .github/workflows/site.yml create mode 100644 site/colibri-icon.svg create mode 100644 site/colibri.svg create mode 100644 site/index.html diff --git a/.github/workflows/site.yml b/.github/workflows/site.yml new file mode 100644 index 0000000..160cd61 --- /dev/null +++ b/.github/workflows/site.yml @@ -0,0 +1,36 @@ +name: Deploy website + +# Publishes site/ to GitHub Pages. One-time repo setup: +# Settings → Pages → Build and deployment → Source: "GitHub Actions". +# Custom domain later: add site/CNAME with the bare domain, point DNS +# (A/AAAA to GitHub Pages IPs or CNAME to .github.io), done. + +on: + push: + branches: [main] + paths: ['site/**', '.github/workflows/site.yml'] + workflow_dispatch: + +permissions: + contents: read + pages: write + id-token: write + +concurrency: + group: pages + cancel-in-progress: true + +jobs: + deploy: + runs-on: ubuntu-latest + environment: + name: github-pages + url: ${{ steps.deployment.outputs.page_url }} + steps: + - uses: actions/checkout@v4 + - uses: actions/configure-pages@v5 + - uses: actions/upload-pages-artifact@v3 + with: + path: site + - id: deployment + uses: actions/deploy-pages@v4 diff --git a/site/colibri-icon.svg b/site/colibri-icon.svg new file mode 100644 index 0000000..76838b9 --- /dev/null +++ b/site/colibri-icon.svg @@ -0,0 +1,50 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + diff --git a/site/colibri.svg b/site/colibri.svg new file mode 100644 index 0000000..4aed0bb --- /dev/null +++ b/site/colibri.svg @@ -0,0 +1,55 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + colibrì + tiny engine, immense model + GLM-5.2 · 744B MoE · int4 · streaming CPU + diff --git a/site/index.html b/site/index.html new file mode 100644 index 0000000..3972690 --- /dev/null +++ b/site/index.html @@ -0,0 +1,456 @@ + + + + + +colibrì — tiny engine, immense model + + + + + + + + + + + + + +
+
+ +

tiny engine, immense model

+

Run GLM-5.2 — a 744-billion-parameter Mixture-of-Experts — on your own machine. + Pure C, zero dependencies, experts streamed from disk. From a 25 GB dev box to a + six-GPU workstation: same engine, same container, the hardware only changes where the experts live.

+
+ ★ GitHub + Quick start + Benchmarks +
+
+
744B
parameters
+
19,456
experts
+
~11 GB
routed / token
+
0
dependencies
+
25 GB
min RAM
+
+
+ +
+

# watch it think

+

A replay of the engine at work, paced to measured decode speeds from real hardware. + Every token routes through 8 experts in each of 75 MoE layers — the grid on the right is all + 19,456 experts: colour is the storage tier, brightness is routing heat, and every expert + a token touches flashes white.

+
+
+
+
+ coli chat —
+
+
+
+
+ the brain — 75 MoE layers × 256 experts + MTP
+
+ +
+
0.0 tok/s
decode
+
— s
ttft (measured)
+
— %
resident hit
+
+
+ VRAM tier + RAM tier + disk tier + routed now +
+
+
+
+

Simulation replays a fixed transcript at each profile's measured decode rate; TTFT shown is the measured value, not a live wait.

+
+ +
+
+

# the expert atlas

+

Routing affinity is measured, not assumed. Characterised experts cluster by what + they actually fire on — poetry, law, Chinese, SQL, code, mathematics… Position below is measured + routing affinity, not a learned embedding. Drag to spin.

+
+
+ +
+
drag to spin · scroll past to release
+
+
+ +
+

# how it fits

+

A 744B MoE activates only ~40B parameters per token — and only ~11 GB of those + change from token to token (the routed experts). colibrì keeps the dense weights resident and tiers the + 19,456 experts (~19 MB each, int4) across three levels by measured heat:

+
+

VRAM

+
hottest experts · grouped GPU matmuls
+

With CUDA or Metal, the hottest experts live on-device and are batched into grouped kernels. + Six RTX 5090s hold the entire routed set: disk reads drop to zero.

+

RAM

+
warm experts · pinned & wired
+

The warm set is pinned in RAM (PIN_GB), guided by + recorded usage stats, and served by AVX-512/NEON int4 kernels with an LRU cache behind it.

+

disk

+
cold tail · streamed on demand
+

Everything else streams from NVMe with io_uring, prefetched a layer ahead by the router. + This is how 744B fits a 25 GB machine at all — the proven floor.

+
+
token → router picks 8 of 256 experts per layer → engine gathers them from + VRAM / RAM / disk → int4 matmuls → next token. Output is validated token-exact against the + reference transformers implementation.
+
+ +
+

# the ladder

+

Same engine, same int4 container — measured decode on real machines, + from the proven floor to full residency.

+ + + + + + + + + + + + +
hardwaredecodewhere the experts live
6× RTX 5090 · full residency5.8–6.8 tok/s
all in VRAM+RAM · disk 0
128 GB CPU-only desktop~1.8 tok/s
hot set pinned in RAM
single RTX 5070 Ti laptop-class1.07 tok/s
GPU-resident pipeline
25 GB dev box · cold0.05–0.1 tok/s
streamed from NVMe
+

Full tables, methodology and quality ablations: docs/benchmarks.md

+
+
+ +
+ 🐦 colibrì — MIT license + GitHub + API + Tuning + the hummingbird: tiny, fast, precise. +
+ + + + From bca4e95d92615c94d9618a6e4d8811d2aea42bda Mon Sep 17 00:00:00 2001 From: ZacharyZcR Date: Mon, 20 Jul 2026 00:08:28 +0800 Subject: [PATCH 23/71] =?UTF-8?q?site:=20redesign=20=E2=80=94=20measured?= =?UTF-8?q?=20atlas=20everywhere,=20vision,=20algorithm,=20models?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - hero: the measured expert atlas as a full-screen slowly-turning backdrop - atlas rebuilt on real data: canonical atlas v1 (721 canonical + 637 gate-sensitive specialists, 10 measured categories) embedded inline; position IS the measured affinity vector, colour = top topic - demo: third panel 'the atlas, live' — token routing flashes measured specialists for the active topic in both the brain grid and the galaxy; added SQL and Chinese-poetry turns so cluster shifts are visible - profiles/ladder updated to current community numbers (#82 NUMA 9.0-9.2, #389 Xeon 1TB 5.42, #387 M5 Max 2.0, #161 GB10 3.33, #120 1.23); unpublished TTFT/hit values shown as em-dash, never invented - new sections: vision manifesto (run/study/improve), 'A JIT, but for weights' algorithm explainer with tier stack, models roadmap (Kimi K2 / Qwen3 MoE / MiniMax planned) + open-weights acknowledgements, contribute cards - visual pass: numbered sections, gradient type, glass panels, fixed nav --- site/index.html | 724 ++++++++++++++++++++++++++++++++---------------- 1 file changed, 483 insertions(+), 241 deletions(-) diff --git a/site/index.html b/site/index.html index 3972690..dd6dac1 100644 --- a/site/index.html +++ b/site/index.html @@ -4,64 +4,124 @@ colibrì — tiny engine, immense model - + - + - -