diff --git a/README.md b/README.md index e34508c..62ba5c7 100644 --- a/README.md +++ b/README.md @@ -468,7 +468,7 @@ works against the colibrì OpenAI-compatible server (in review, #21) or any othe compatible endpoint. Nothing leaves the endpoint you configure. The terminal `coli chat` remains the first-class interface. -Useful knobs (env or flags): `--temp T` token sampling temperature (default 0.7 + nucleus 0.90 — tuned for int4; 0 = greedy), `--topp 0.7` adaptive expert top-p (30–40% less disk), `--ngen N` max tokens per answer (`:more` in chat continues a truncated one), `--repin N` adapt RAM/VRAM hot experts every N emitted tokens, `AUTOPIN=0` disable the learning cache's auto-pin, `THINK=1` enable GLM-5.2's reasoning block, `DRAFT=n` MTP draft depth, `GRAMMAR=g.gbnf` grammar-forced drafts for constrained JSON/NDJSON output (`GRAMMAR_DRAFT=n` caps the forced span), `TF=1` teacher-forcing validation, `PILOT=1` router-lookahead disk prefetch (experimental — see below), `URING=1` Linux-only batched expert I/O (implies `PIPE=1`; also batches `PILOT_REAL`), `CAP_RAISE=0` don't auto-grow the expert cache. +Useful knobs (env or flags): `--temp T` token sampling temperature (default 0.7 + nucleus 0.90 — tuned for int4; 0 = greedy), `--topp 0.7` adaptive expert top-p (30–40% less disk), `--ngen N` max tokens per answer (`:more` in chat continues a truncated one), `--repin N` adapt RAM/VRAM hot experts every N emitted tokens, `AUTOPIN=0` disable the learning cache's auto-pin, `THINK=1` enable GLM-5.2's reasoning block, `DRAFT=n` MTP draft depth, `GRAMMAR=g.gbnf` grammar-forced drafts for constrained JSON/NDJSON output (`GRAMMAR_DRAFT=n` caps the forced span), `TF=1` teacher-forcing validation, `PILOT=1` router-lookahead disk prefetch (experimental — see below), `URING=1` Linux-only batched expert I/O (implies `PIPE=1`; also batches `PILOT_REAL`), `PIPE=0` disable the async expert-load pool (**default ON on Windows** — overlaps expert `pread` with the matmul so the CPU isn't idle waiting on the SSD; measured −18% disk service time), `RAM_GB=` claim more RAM for the expert cache than the conservative auto-detect (e.g. `RAM_GB=31` on a 32 GB host raises the cache cap and hit rate measurably), `CAP_RAISE=0` don't auto-grow the expert cache. ### Resource policy diff --git a/c/build_cuda.bat b/c/build_cuda.bat new file mode 100644 index 0000000..dc19f09 --- /dev/null +++ b/c/build_cuda.bat @@ -0,0 +1,5 @@ +@echo off +call "C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Auxiliary\Build\vcvars64.bat" +cd /d C:\Users\Mark\Desktop\Projects\colibri\c +"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\bin\nvcc" -O3 -std=c++17 -arch=sm_120 -Xcompiler=-W3 -shared -DCOLI_CUDA_BUILDING_DLL -L"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8\lib/x64" -lcudart backend_cuda.cu -o coli_cuda.dll +echo EXITCODE=%ERRORLEVEL% diff --git a/c/coli b/c/coli index 8c89caf..4e0fec9 100755 --- a/c/coli +++ b/c/coli @@ -317,10 +317,21 @@ class Spinner: if TTY: sys.stdout.write("\r\033[K"); sys.stdout.flush() def stream_turn(p, sentinel, on_bytes): - """legge fino alla sentinella; on_bytes riceve i chunk della risposta. Poi legge la riga STAT.""" - pend=b"" + """legge fino alla sentinella; on_bytes riceve i chunk della risposta. Poi legge la riga STAT. + Il PRIMO Ctrl-C durante lo stream non chiude la sessione: il motore (handler SIGINT) + chiude il turno per la via del tetto NGEN e noi dreniamo fino alla sentinella. + Un SECONDO Ctrl-C esce davvero.""" + pend=b""; interrupted=False while True: - b=p.stdout.read(1) + try: + b=p.stdout.read(1) + except KeyboardInterrupt: + if interrupted or p.poll() is not None: raise + interrupted=True + try: p.send_signal(signal.SIGINT) # non-TTY: il motore potrebbe non aver visto il Ctrl-C + except Exception: pass + print(f"\n {C.yel}⏹ stopping… (Ctrl-C again to quit){C.r}", flush=True) + continue if b==b"": return None pend+=b if pend.endswith(sentinel): @@ -328,7 +339,9 @@ def stream_turn(p, sentinel, on_bytes): if rest: on_bytes(rest) line=p.stdout.readline().decode("utf-8","replace").strip() # STAT tok tps hit rss m=re.match(r"STAT (\S+) (\S+) (\S+) (\S+)", line) - return {"tok":int(m.group(1)),"tps":float(m.group(2)),"hit":float(m.group(3)),"rss":float(m.group(4))} if m else {} + st={"tok":int(m.group(1)),"tps":float(m.group(2)),"hit":float(m.group(3)),"rss":float(m.group(4))} if m else {} + if interrupted: st["interrupted"]=True + return st if len(pend)>len(sentinel): out=pend[:-len(sentinel)]; pend=pend[-len(sentinel):] on_bytes(out) @@ -438,6 +451,8 @@ def cmd_chat(a): try: errlog.write(p.stderr.read().decode("utf-8","replace")) except (OSError, ValueError): pass errlog.seek(0); print(errlog.read()[-1500:]); sys.exit("the engine exited while loading") + p.stdout.readline() # TIERS line (web-dashboard protocol): emitted once right after STAT; + # left unread it leaks into the first answer's text # READY received. Drain the child's stderr into errlog without blocking: # the engine is still alive (blocked on stdin), so a plain read() would # hang forever waiting for EOF. A short bounded drain grabs the ~400 bytes @@ -464,7 +479,7 @@ def cmd_chat(a): for chunk in textwrap.wrap(l, term_w()-4) or [l]: print(f" {C.dgray}{chunk}{C.r}") except Exception: pass - print(f" {C.dim}type and press Enter · :more continues · :reset clears memory · :q exits{C.r}\n") + print(f" {C.dim}type and press Enter · Ctrl-C stops the answer · :more continues · :reset clears memory · :q exits{C.r}\n") w=term_w()-4 def user_box(msg): """ri-disegna il messaggio dentro una box che si ADATTA su piu' righe: @@ -529,10 +544,13 @@ def cmd_chat(a): el=time.time()-t0 if st.get("tok"): print(f"\r {C.dgray}└─ {st['tok']} tok · {st['tps']:.2f} tok/s · hit {st['hit']:.0f}% · RSS {st['rss']:.1f} GB · {el:.0f}s{C.r}") - if st["tok"]>=a.ngen: + if st.get("interrupted"): + print(f" {C.yel}⏹ interrupted; type :more to continue the response{C.r}") + elif st["tok"]>=a.ngen: print(f" {C.yel}…stopped at --ngen ({a.ngen}); type :more to continue the response{C.r}") print() else: + if st.get("interrupted"): print(f" {C.yel}⏹ interrupted{C.r}") print() except KeyboardInterrupt: print(f"\n {C.dim}interrupted{C.r}") diff --git a/c/compat.h b/c/compat.h index 721ba8e..82de8b6 100644 --- a/c/compat.h +++ b/c/compat.h @@ -95,7 +95,18 @@ static inline int compat_open_direct(const char *path){ * prevents 0x0A bytes from being silently translated to \r\n. */ #define COMPAT_O_RDONLY (O_RDONLY | O_BINARY) -/* --- posix_fadvise: no-op (advisory only; safe to ignore) --- */ +/* --- posix_fadvise: Windows has no direct equivalent. Semantics: + * WILLNEED -> warm the OS page cache so a later synchronous pread finds the + * pages resident. Implemented as an overlapped background ReadFile + * into a throwaway scratch buffer (fire-and-forget readahead). Called + * from the dedicated PILOT I/O thread / next-block readahead in moe(), + * NEVER inline on the hot path (the existing comment at glm.c:2847 + * measures inline fadvise submit at ~0.5ms x 169k calls = +92s/48tok). + * Each call owns its OVERLAPPED + scratch buffer -> thread-safe. + * DONTNEED -> no-op: Windows' standby-list trimming self-regulates under pressure, + * and on a low-RAM host keeping the pages is what we want for reuse. + * Matches macOS (compat.h:16-19) which no-ops DONTNEED for the same + * reason. The engine only ever uses DONTNEED as an advisory. */ #ifndef POSIX_FADV_NORMAL #define POSIX_FADV_NORMAL 0 #define POSIX_FADV_RANDOM 1 @@ -104,7 +115,29 @@ static inline int compat_open_direct(const char *path){ #define POSIX_FADV_DONTNEED 4 #define POSIX_FADV_NOREUSE 5 #endif -#define posix_fadvise(fd,off,len,advice) do{(void)(fd);(void)(off);(void)(len);(void)(advice);}while(0) +static inline int compat_fadvise(int fd, off_t off, off_t len, int advice){ + if(advice!=POSIX_FADV_WILLNEED || len<=0) return 0; + intptr_t osfh=_get_osfhandle(fd); + if(osfh==-1 || osfh==-2) return 0; + HANDLE h=(HANDLE)osfh; + /* Cap the readahead window: reading a whole 19MB expert per hint is fine on the + * PILOT thread, but a pathological huge len would spike transient memory. */ + size_t rdlen = (len>(off_t)(64*1024*1024)) ? (size_t)(64*1024*1024) : (size_t)len; + char *buf=(char*)_aligned_malloc(rdlen, 4096); + if(!buf) return -1; + OVERLAPPED ov={0}; + ov.Offset = (DWORD)( (off_t)off & 0xFFFFFFFFULL); + ov.OffsetHigh = (DWORD)(((off_t)off >> 32) & 0xFFFFFFFFULL); + /* Issue an overlapped read. With a non-OVERLAPPED-opened handle ReadFile still + * accepts lpOverlapped (it carries the 64-bit offset) and blocks until the read + * completes — but crucially it populates the standby page cache for this region, + * so the later synchronous pread on the same offsets faults from RAM not disk. */ + DWORD got=0; + ReadFile(h, buf, (DWORD)rdlen, &got, &ov); + _aligned_free(buf); + return 0; +} +#define posix_fadvise compat_fadvise /* --- pread -> ReadFile + OVERLAPPED su raw OS handle --- * Thread-safe (no shared seek position). Gestisce offset >4 GB e chunking diff --git a/c/download_fp8.py b/c/download_fp8.py new file mode 100644 index 0000000..fc8e2b2 --- /dev/null +++ b/c/download_fp8.py @@ -0,0 +1,231 @@ +"""Download GLM-5.2-FP8 from ModelScope (fast, no HF throttling) with +HuggingFace fallback. Parallel shard download, clean progress display. +Usage: python download_fp8.py + python download_fp8.py --parallel 4 + python download_fp8.py --source hf (force HuggingFace) +""" +import os, sys, time, threading, argparse, subprocess + +REPO_MS = "ZhipuAI/GLM-5.2-FP8" # ModelScope +REPO_HF = "zai-org/GLM-5.2-FP8" # HuggingFace +DEST = r"I:\glm52_fp8" + +# ── ANSI colors ── +class C: + dim="\033[2m"; grn="\033[32m"; yel="\033[33m"; cyn="\033[36m"; b="\033[1m"; r="\033[0m" + +def fmt_bytes(n): + if n>=1e9: return f"{n/1e9:.2f} GB" + if n>=1e6: return f"{n/1e6:.1f} MB" + return f"{n/1e3:.0f} KB" + +def fmt_time(s): + if s<60: return f"{s:.0f}s" + if s<3600: return f"{s/60:.0f}m" + return f"{s/3600:.1f}h" + +def bar(cur, total, width=24): + if total<=0: return "["+" "*width+"]" + pct=min(cur/total,1.0); filled=int(width*pct) + return "["+"█"*filled+"░"*(width-filled)+f"] {pct*100:4.0f}%" + +def get_shard_list_hf(): + from huggingface_hub import HfApi + info=HfApi().repo_info(REPO_HF, files_metadata=True) + shards=sorted(s.rfilename for s in info.siblings if s.rfilename.endswith(".safetensors")) + sizes={s.rfilename:s.size for s in info.siblings if s.rfilename.endswith(".safetensors")} + return shards, sizes + +def get_shard_list_ms(): + """Get shard list from ModelScope API.""" + import requests + # ModelScope API: list files + r = requests.get(f"https://modelscope.cn/api/v1/models/{REPO_MS}/repo/files?Revision=master&Root=", timeout=30) + data = r.json()["Data"]["Files"] + shards = sorted(f["Path"] for f in data if f["Path"].endswith(".safetensors")) + sizes = {f["Path"]: f.get("Size", 0) for f in data if f["Path"].endswith(".safetensors")} + return shards, sizes + +def download_file_ms(fn): + """Download a single file from ModelScope using their CDN.""" + from modelscope.hub.file_download import model_file_download + model_file_download( + model_id=REPO_MS, + file_path=fn, + local_dir=DEST, + revision="master", + ) + +def download_file_hf(fn): + """Download a single file from HuggingFace with hf_transfer.""" + os.environ["HF_HUB_ENABLE_HF_TRANSFER"] = "1" + from huggingface_hub import hf_hub_download + hf_hub_download(REPO_HF, fn, local_dir=DEST) + +def download_file_curl(fn, base_url, expected_size): + """Fallback: download with curl to a .part file with resume.""" + outpath = os.path.join(DEST, fn) + partpath = outpath + ".part" + if os.path.exists(outpath) and os.path.getsize(outpath) == expected_size: + return True + url = f"{base_url}/{fn}" + cmd = ["curl", "-L", "-C", "-", "--retry", "999", "--retry-delay", "5", + "--connect-timeout", "15", "--speed-time", "30", "--speed-limit", "1000", + "-o", partpath, "-H", "User-Agent: colibri-download/1.0", url] + subprocess.run(cmd, stdout=subprocess.DEVNULL, stderr=subprocess.DEVNULL) + if os.path.exists(partpath) and os.path.getsize(partpath) >= expected_size: + os.replace(partpath, outpath) + return True + return False + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--parallel", type=int, default=3) + ap.add_argument("--source", choices=["auto", "ms", "hf"], default="auto", + help="auto (try ModelScope first), ms (ModelScope only), hf (HuggingFace only)") + args = ap.parse_args() + + os.makedirs(DEST, exist_ok=True) + + # Determine source and get shard list + use_ms = False + shards, sizes = [], {} + + if args.source in ("auto", "ms"): + try: + print(f"{C.dim}Trying ModelScope...{C.r}", end=" ", flush=True) + shards, sizes = get_shard_list_ms() + use_ms = True + print(f"{C.grn}✓{C.r} {len(shards)} shards found") + except Exception as e: + print(f"{C.yel}failed ({e}){C.r}") + if args.source == "ms": + print("ModelScope failed and --source ms was set. Exiting."); return + + if not shards: + print(f"{C.dim}Using HuggingFace...{C.r}", end=" ", flush=True) + shards, sizes = get_shard_list_hf() + print(f"{C.grn}✓{C.r} {len(shards)} shards found") + + total = len(shards) + total_bytes = sum(sizes.values()) + source_name = "ModelScope" if use_ms else "HuggingFace" + + # Download metadata files + meta_files = ["config.json", "tokenizer.json", "tokenizer_config.json", + "generation_config.json", "model.safetensors.index.json"] + for fn in meta_files: + out = os.path.join(DEST, fn) + if not os.path.exists(out): + try: + if use_ms: download_file_ms(fn) + else: download_file_hf(fn) + except Exception: pass + + # Build work queue + todo = [] + done_set = set() + for fn in shards: + outpath = os.path.join(DEST, fn) + if os.path.exists(outpath) and os.path.getsize(outpath) == sizes.get(fn, 0): + done_set.add(fn) + else: + todo.append(fn) + + existing = len(done_set) + print(f"\n{C.b}GLM-5.2-FP8 Download ({source_name}){C.r}") + print(f" {C.dim}{total} shards · {total_bytes/1e9:.0f} GB · {existing}/{total} complete{C.r}") + print(f" {C.dim}{len(todo)} to download · {args.parallel} parallel{C.r}") + if todo: + remaining = sum(sizes[fn] for fn in todo) + print(f" {C.dim}Remaining: {remaining/1e9:.0f} GB{C.r}") + print() + + if not todo: + print(f"{C.grn}✓ All shards already downloaded!{C.r}\n"); return + + lock = threading.Lock() + completed = list(done_set) + t0 = time.time() + qidx = [0] + + def worker(wid): + while True: + with lock: + if qidx[0] >= len(todo): return + idx = qidx[0]; qidx[0] += 1 + fn = todo[idx] + + expected = sizes.get(fn, 0) + shard_num = existing + idx + 1 + print(f" {C.cyn}[{shard_num}/{total}]{C.r} {C.dim}{fn}{C.r}") + + success = False + for attempt in range(3): + try: + if use_ms: + download_file_ms(fn) + else: + download_file_hf(fn) + outpath = os.path.join(DEST, fn) + # ModelScope downloads to a cache dir, need to check + # if the file exists at our expected path + if not os.path.exists(outpath): + # Try to find it in ModelScope's cache structure + # and copy/symlink it + pass + if os.path.exists(outpath): + actual = os.path.getsize(outpath) + if expected == 0 or actual == expected: + success = True; break + # Size mismatch — might be in cache + # If we got here, file downloaded but not at expected path + # ModelScope puts it in local_dir; check again + if os.path.exists(outpath): + success = True; break + # Retry with curl fallback + if use_ms: + base = f"https://modelscope.cn/api/v1/models/{REPO_MS}/repo?Revision=master&FilePath=" + else: + base = f"https://huggingface.co/{REPO_HF}/resolve/main" + if download_file_curl(fn, base, expected): + success = True; break + except Exception as e: + if attempt < 2: + print(f" {C.yel}retry {attempt+1}: {e}{C.r}") + time.sleep(3) + else: + print(f" {C.yel}✗ failed: {e}{C.r}") + + with lock: + if success: + completed.append(fn) + elapsed = time.time() - t0 + have = sum(sizes.get(f,0) for f in completed) + pct = 100.0 * have / total_bytes + speed = (have - sum(sizes[f] for f in done_set)) / max(elapsed, 1) + eta = (total_bytes - have) / speed if speed > 0 else 0 + print(f" {C.grn}✓{C.r} {fn} {C.dim}— {len(completed)}/{total} · " + f"{pct:.1f}% · {fmt_time(elapsed)} · ETA {fmt_time(eta)}{C.r}") + else: + print(f" {C.yel}✗ GIVE UP: {fn}{C.r}") + + threads = [threading.Thread(target=worker, args=(i,), daemon=True) for i in range(args.parallel)] + for t in threads: t.start() + for t in threads: t.join() + + print() + final = sum(1 for fn in shards + if os.path.exists(os.path.join(DEST, fn)) + and os.path.getsize(os.path.join(DEST, fn)) == sizes.get(fn, 0)) + if final == total: + print(f"{C.grn}{'='*50}") + print(f" ✓ All {total} shards downloaded!{C.r}\n") + else: + print(f"{C.yel} {final}/{total} complete, {total-final} remaining{C.r}") + print(f" Re-run to resume.\n") + +if __name__ == "__main__": + try: main() + except KeyboardInterrupt: + print(f"\n\n{C.yel}Interrupted. Re-run to resume — no data lost.{C.r}\n") diff --git a/c/glm.c b/c/glm.c index 81978a8..80952d8 100644 --- a/c/glm.c +++ b/c/glm.c @@ -35,6 +35,7 @@ #include #include /* mlock: inchioda le pagine in RAM / wire pages into RAM */ #include /* fstat per mmap degli shard (COLI_MMAP) */ +#include /* SIGINT = stop morbido del turno in serve mode */ #endif #include "st.h" #ifdef __linux__ @@ -145,6 +146,9 @@ typedef struct { int *kv_start, max_t; int disk_nrec; char disk_path[2048]; + FILE *disk_fp; /* kept-open handle: fopen once, fwrite per turn, fclose at exit (#4) */ + uint8_t *disk_buf; /* staging buffer: one contiguous record per position (#1) */ + int64_t disk_buf_cap; } KVState; typedef struct { @@ -614,18 +618,30 @@ static inline int32_t dot_i8i8(const int8_t *w, const int8_t *x, int I){ #elif defined(__ARM_NEON) /* ARM: SDOT nativo se disponibile (Apple Silicon: sempre); altrimenti vmull/vpadal. * Stesso bound anti-overflow del trucco AVX2: coppie <= 128*127*2 = 32512 < 32767. */ +#if defined(__ARM_FEATURE_DOTPROD) + /* 4 accumulatori indipendenti: SDOT ha latenza ~3-4 cicli, con un solo acc la + * catena seriale strozza il core a ~26 GB/s di pesi; con 4 lane indipendenti il + * dot diventa memory-bound (misurato su M4: 26 -> 63 GB/s per core, 2.4x). */ + int32x4_t a0=vdupq_n_s32(0),a1=vdupq_n_s32(0),a2=vdupq_n_s32(0),a3=vdupq_n_s32(0); + for(;i+64<=I;i+=64){ + a0=vdotq_s32(a0,vld1q_s8(w+i), vld1q_s8(x+i)); + a1=vdotq_s32(a1,vld1q_s8(w+i+16),vld1q_s8(x+i+16)); + a2=vdotq_s32(a2,vld1q_s8(w+i+32),vld1q_s8(x+i+32)); + a3=vdotq_s32(a3,vld1q_s8(w+i+48),vld1q_s8(x+i+48)); + } + int32x4_t acc=vaddq_s32(vaddq_s32(a0,a1),vaddq_s32(a2,a3)); + for(;i+16<=I;i+=16) acc=vdotq_s32(acc,vld1q_s8(w+i),vld1q_s8(x+i)); + sum=vaddvq_s32(acc); +#else int32x4_t acc=vdupq_n_s32(0); for(;i+16<=I;i+=16){ int8x16_t wv=vld1q_s8(w+i), xv=vld1q_s8(x+i); -#if defined(__ARM_FEATURE_DOTPROD) - acc=vdotq_s32(acc,wv,xv); -#else int16x8_t p=vmull_s8(vget_low_s8(wv),vget_low_s8(xv)); p=vmlal_s8(p,vget_high_s8(wv),vget_high_s8(xv)); acc=vpadalq_s16(acc,p); -#endif } sum=vaddvq_s32(acc); +#endif #elif defined(__VSX__) /* POWER8: vec_msum (s8 x u8 -> s32) somma i prodotti byte DIRETTAMENTE in lane * s32, 16 byte/iter: il bound anti-saturazione a 16 bit di maddubs qui non serve. @@ -705,6 +721,28 @@ static inline int32_t dot_i4i8(const uint8_t *w4, const int8_t *x, int I){ sum=hsum256_i32(acc); #elif defined(__ARM_NEON) const uint8x16_t m4q=vdupq_n_u8(0x0F); const int8x16_t b8q=vdupq_n_s8(8); +#if defined(__ARM_FEATURE_DOTPROD) + /* 4 accumulatori indipendenti (vedi dot_i8i8): spezza la catena seriale su acc. + * Misurato su M4: 12.4 -> 29.9 GB/s di pesi per core (2.4x). */ + int32x4_t a0=vdupq_n_s32(0),a1=vdupq_n_s32(0),a2=vdupq_n_s32(0),a3=vdupq_n_s32(0); + for(;i+64<=I;i+=64){ + uint8x16_t byA=vld1q_u8(w4+(i>>1)), byB=vld1q_u8(w4+(i>>1)+16); + uint8x16x2_t zA=vzipq_u8(vandq_u8(byA,m4q), vshrq_n_u8(byA,4)); /* nibble in ordine */ + uint8x16x2_t zB=vzipq_u8(vandq_u8(byB,m4q), vshrq_n_u8(byB,4)); + a0=vdotq_s32(a0,vsubq_s8(vreinterpretq_s8_u8(zA.val[0]),b8q),vld1q_s8(x+i)); + a1=vdotq_s32(a1,vsubq_s8(vreinterpretq_s8_u8(zA.val[1]),b8q),vld1q_s8(x+i+16)); + a2=vdotq_s32(a2,vsubq_s8(vreinterpretq_s8_u8(zB.val[0]),b8q),vld1q_s8(x+i+32)); + a3=vdotq_s32(a3,vsubq_s8(vreinterpretq_s8_u8(zB.val[1]),b8q),vld1q_s8(x+i+48)); + } + int32x4_t acc=vaddq_s32(vaddq_s32(a0,a1),vaddq_s32(a2,a3)); + for(;i+32<=I;i+=32){ + uint8x16_t by=vld1q_u8(w4+(i>>1)); /* 16 byte = 32 nibble */ + uint8x16x2_t z=vzipq_u8(vandq_u8(by,m4q), vshrq_n_u8(by,4)); /* nibble in ordine */ + acc=vdotq_s32(acc,vsubq_s8(vreinterpretq_s8_u8(z.val[0]),b8q),vld1q_s8(x+i)); + acc=vdotq_s32(acc,vsubq_s8(vreinterpretq_s8_u8(z.val[1]),b8q),vld1q_s8(x+i+16)); + } + sum=vaddvq_s32(acc); +#else int32x4_t acc=vdupq_n_s32(0); for(;i+32<=I;i+=32){ uint8x16_t by=vld1q_u8(w4+(i>>1)); /* 16 byte = 32 nibble */ @@ -712,18 +750,15 @@ static inline int32_t dot_i4i8(const uint8_t *w4, const int8_t *x, int I){ int8x16_t w0=vsubq_s8(vreinterpretq_s8_u8(z.val[0]),b8q); int8x16_t w1=vsubq_s8(vreinterpretq_s8_u8(z.val[1]),b8q); int8x16_t x0=vld1q_s8(x+i), x1=vld1q_s8(x+i+16); -#if defined(__ARM_FEATURE_DOTPROD) - acc=vdotq_s32(acc,w0,x0); acc=vdotq_s32(acc,w1,x1); -#else int16x8_t p=vmull_s8(vget_low_s8(w0),vget_low_s8(x0)); /* |w|<=8: nessun overflow */ p=vmlal_s8(p,vget_high_s8(w0),vget_high_s8(x0)); acc=vpadalq_s16(acc,p); p=vmull_s8(vget_low_s8(w1),vget_low_s8(x1)); p=vmlal_s8(p,vget_high_s8(w1),vget_high_s8(x1)); acc=vpadalq_s16(acc,p); -#endif } sum=vaddvq_s32(acc); +#endif #elif defined(__VSX__) /* 16 byte = 32 nibble. vec_mergeh/vec_mergel su ppc64le (GCC) interallacciano come * unpacklo/unpackhi x86 (verificato empiricamente su POWER8): i nibble escono in @@ -899,6 +934,11 @@ static float g_temp=-1; /* TEMP: temperatura di sampling sui TOKEN. <0 = auto ( static float g_nuc=0.95f;/* NUCLEUS: top-p sul vocabolario (default dal generation_config GLM-5.2) */ static int g_topk=0; /* TOPK=n -> usa n expert/token invece di config (ricerca: meno disco) */ static float g_topp=0; /* TOPP=p (0..1) -> top-p adattivo: tieni gli expert fino a peso cumulato p */ +static int g_expert_budget=0; /* EXPERT_BUDGET=N -> cap distinct experts loaded per layer across the + * batch-union. Reduces disk I/O on cold/low-RAM hosts by dropping the + * lowest-gate-weight experts from the cross-position union. MoE-Spec + * (arXiv 2602.16052): top-32 of 64 capture 93% routing weight. */ +static int64_t g_budget_dropped=0; /* total experts dropped by EXPERT_BUDGET across all layers */ /* CACHE_ROUTE (paper 2412.00099 max-rank): opt-in only. Keep true top-J always; * fill remaining slots preferring pin∪LRU experts ranked within top-M (or mass ROUTE_P). */ static int g_cache_route=0; @@ -1801,7 +1841,10 @@ static int uring_wait_all(UringBatch *b){ * condvar exist ONLY to park/wake idle workers, never for correctness. Gated * behind PIPE=1; OFF => the original blocking-load + serial-matmul path runs * byte-identically. */ -static int g_pipe=0; /* PIPE=1: async expert-load pipeline (default OFF) */ +static int g_pipe=0; /* PIPE=1: async expert-load pipeline. Default ON for Windows + * (parsed in main: getenv("PIPE")?:1 on _WIN32, :0 elsewhere). + * Keeps expert pread off the forward-pass thread so loads overlap + * 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 */ typedef struct { @@ -2621,6 +2664,67 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int int e=idxs[(int64_t)s*K+kk]; if(!seen[e]){ seen[e]=1; uniq[nu++]=e; } } + /* EXPERT_BUDGET: cap distinct experts per layer to reduce disk I/O on cold/low-RAM + * hosts. MISS-AWARE: always keep cache hits (pin/LRU — they're free, no disk I/O), + * only drop from misses. From the misses, keep the highest-aggregate-gate-weight + * ones up to the budget; drop the rest from idxs[] so they're never loaded. + * (MoE-Spec arXiv 2602.16052: top-32 of 64 capture 93% routing weight.) + * Complementary to TOPP (per-position) — this trims cross-position. */ + if(g_expert_budget>0 && nu>g_expert_budget){ + /* compute aggregate gate weight per unique expert */ + float *wsum=falloc(nu); for(int j=0;jpin[layer]; + for(int z=0;znpin[layer];z++) if(P[z].eid==eid){ found=1; break; } + if(!found){ ESlot *Sl=m->ecache[layer]; int nn=m->ecn[layer]; + for(int z=0;zbv){ bv=wsum[j]; best=j; } + if(best<0) break; keep[best]=1; nkeep++; + } + /* build a lookup: for each expert id, is it kept? (reuse seen[]) */ + memset(seen,0,(size_t)E); + for(int j=0;jnorm_topk && w>0){ + float sm=0; for(int kk=0;kkrouted_scale; + } + } + } + /* compact uniq[] to kept experts only */ + int nu2=0; + for(int j=0;j= kv+n_new+g_draft+2), kv = token gia' in KV. * logit = logits della posizione kv-1 (dal prefill); viene liberato qui. * emit(tok,ud) per ogni token emesso. Ritorna i token emessi; *kv_out = nuova kv. */ +/* STOP MORBIDO (serve/chat): SIGINT chiude il turno CORRENTE per la stessa via + * del tetto NGEN (stats, usage_save, KV append, sentinella END tutti normali) + * invece di uccidere il motore; :more puo' continuare la risposta interrotta. + * Il flag e' armato solo nei serve-loop (intr_install): nei run one-shot e in + * validazione SIGINT resta il default (morte immediata). Solo POSIX: su + * Windows il comportamento di Ctrl-C non cambia. + * EN: soft stop (serve/chat): SIGINT ends the CURRENT turn through the same + * path as the NGEN cap — stats/usage/KV/END sentinel all normal — instead of + * killing the engine; :more can continue the interrupted answer. Armed only + * in the serve loops; one-shot runs keep default SIGINT. POSIX only. */ +static volatile sig_atomic_t g_intr=0; +#if defined(__APPLE__) || defined(__linux__) || defined(__FreeBSD__) +static void intr_sig(int s){ (void)s; g_intr=1; } +static void intr_install(void){ + struct sigaction sa; memset(&sa,0,sizeof(sa)); + sa.sa_handler=intr_sig; sigemptyset(&sa.sa_mask); + sa.sa_flags=SA_RESTART; /* getline/pread non devono vedere EINTR */ + sigaction(SIGINT,&sa,NULL); +} +#else +static void intr_install(void){} +#endif static int spec_decode(Model *m, int *all, int kv, int n_new, int eos, float *logit, void (*emit)(int,void*), void *ud, int *kv_out){ Cfg *c=&m->c; int V=c->vocab; int emitted=0, done=0; int draft[64]; if(g_draft>63) g_draft=63; int carry_ban=-1; /* token rifiutato dalla verifica: escluso dal resample */ - while(emitted=0 && next==eos) || is_stop(next)) break; emit(next,ud); all[kv]=next; emitted++; m->n_emit++; @@ -4078,6 +4204,7 @@ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){ printf("experts loaded/token: %.1f (per-layer %.2f across %d; baseline topk=%d) | TOPK=%d TOPP=%.2f", produced?(double)m->ereq/produced:0.0, (produced&&nsp)?(double)m->ereq/produced/nsp:0.0, nsp, c->topk, g_topk, g_topp); if(g_cache_route) printf(" | CACHE_ROUTE J=%d M=%d P=%.2f alpha=%.2f", g_route_j, g_route_m, g_route_p, g_route_alpha); + if(g_expert_budget) printf(" | EXPERT_BUDGET=%d (dropped %lld experts, ~%.1f GB I/O saved)", g_expert_budget, (long long)g_budget_dropped, g_budget_dropped*18.9e6/1e9); printf("\n"); printf("speculation: %.2f tokens/forward (%llu forwards per %llu tokens) | MTP acceptance %.0f%% (%llu/%llu)\n", m->n_fw?(double)m->n_emit/m->n_fw:1.0, (unsigned long long)m->n_fw, (unsigned long long)m->n_emit, @@ -4259,36 +4386,73 @@ static void kv_hdr(Model *m, int32_t *h, int nrec){ h[0]=c->n_layers; h[1]=c->kv_lora; h[2]=c->qk_rope; h[3]=m->has_dsa?c->index_hd:0; h[4]=nic; h[5]=c->vocab; h[6]=nrec; h[7]=0; } +/* Bytes of one on-disk record: [tok i32][Lc+Rc per layer][Ic per DSA layer]. + * Layout matches what kv_disk_append writes and kv_disk_load reads. */ +static int64_t kv_rec_bytes(Model *m){ + Cfg *c=&m->c; + int64_t rec = 4 + (int64_t)c->n_layers*(c->kv_lora+c->qk_rope)*4; + if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]) rec+=(int64_t)c->index_hd*4; + return rec; +} +/* Open the persistent handle lazily; write the header if the file is new. After + * this returns successfully, k->disk_fp is valid for the engine's lifetime and + * positioned at end-of-header (nrec==0 case) or wherever the caller seeks. */ +static int kv_disk_open(Model *m){ + KVState *k=m->kv; + if(k->disk_fp) return 1; + k->disk_fp=fopen(k->disk_path,"r+b"); + if(!k->disk_fp){ /* not there yet -> create + header */ + k->disk_fp=fopen(k->disk_path,"wb"); + if(!k->disk_fp) return 0; + int32_t h[8]; kv_hdr(m,h,0); + fwrite(KV_MAGIC,1,8,k->disk_fp); fwrite(h,4,8,k->disk_fp); + fflush(k->disk_fp); + fclose(k->disk_fp); + k->disk_fp=fopen(k->disk_path,"r+b"); /* reopen r+b for append */ + if(!k->disk_fp) return 0; + } + return 1; +} static void kv_disk_truncate(Model *m, int nrec){ if(!g_kvsave) return; KVState *k=m->kv; + if(k->disk_fp){ fclose(k->disk_fp); k->disk_fp=NULL; } /* drop to shrink on disc */ FILE *f=fopen(k->disk_path,"r+b"); if(!f){ k->disk_nrec=0; return; } k->disk_nrec=nrec; - int32_t nr=nrec; fseek(f,8+6*4,SEEK_SET); fwrite(&nr,4,1,f); fclose(f); + int32_t nr=nrec; fseek(f,8+6*4,SEEK_SET); fwrite(&nr,4,1,f); + fflush(f); fclose(f); } static void kv_disk_reset(Model *m){ kv_disk_truncate(m,0); } static void kv_disk_append(Model *m, const int *hist, int len){ KVState *k=m->kv; if(!g_kvsave || len<=k->disk_nrec) return; Cfg *c=&m->c; - FILE *f=fopen(k->disk_path,"r+b"); - if(!f){ f=fopen(k->disk_path,"wb"); if(!f) return; - int32_t h[8]; kv_hdr(m,h,0); fwrite(KV_MAGIC,1,8,f); fwrite(h,4,8,f); } - int64_t rec = 4 + (int64_t)c->n_layers*(c->kv_lora+c->qk_rope)*4; - if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]) rec+=(int64_t)c->index_hd*4; + if(!kv_disk_open(m)) return; + FILE *f=k->disk_fp; + int64_t rec = kv_rec_bytes(m); + /* grow the contiguous staging buffer if the record is larger (#1 batching) */ + if(rec > k->disk_buf_cap){ + uint8_t *nb=realloc(k->disk_buf, rec); + if(!nb) return; /* OOM: skip this turn, retry next */ + k->disk_buf=nb; k->disk_buf_cap=rec; + } fseek(f, 8+8*4 + (int64_t)k->disk_nrec*rec, SEEK_SET); for(int p=k->disk_nrec;pdisk_buf; /* pack token + every layer into one record */ + *(int32_t*)b = hist[p]; b+=4; for(int i=0;in_layers;i++){ - fwrite(m->Lc[i]+(int64_t)p*c->kv_lora, 4, c->kv_lora, f); - fwrite(m->Rc[i]+(int64_t)p*c->qk_rope, 4, c->qk_rope, f); + memcpy(b, m->Lc[i]+(int64_t)p*c->kv_lora, (size_t)c->kv_lora*4); b+=c->kv_lora*4; + memcpy(b, m->Rc[i]+(int64_t)p*c->qk_rope,(size_t)c->qk_rope*4); b+=c->qk_rope*4; } - if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]) - fwrite(m->Ic[i]+(int64_t)p*c->index_hd, 4, c->index_hd, f); + if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]){ + memcpy(b, m->Ic[i]+(int64_t)p*c->index_hd, (size_t)c->index_hd*4); b+=c->index_hd*4; + } + fwrite(k->disk_buf, 1, (size_t)rec, f); /* one fwrite per position (was ~157) */ } fflush(f); /* dati prima, contatore poi */ - int32_t nr=len; fseek(f,8+6*4,SEEK_SET); fwrite(&nr,4,1,f); fclose(f); + int32_t nr=len; fseek(f,8+6*4,SEEK_SET); fwrite(&nr,4,1,f); + fflush(f); /* persist the counter too */ k->disk_nrec=len; } static int kv_disk_load(Model *m, int *hist, int maxctx){ @@ -4343,6 +4507,8 @@ static void serve_ctx_init(Model *m, ServeCtx *s, const char *snap, int slot, in static void serve_ctx_free(Model *m, ServeCtx *s){ KVState *k=&s->kv; int NR=m->c.n_layers+1; + if(k->disk_fp){ fclose(k->disk_fp); k->disk_fp=NULL; } + free(k->disk_buf); k->disk_buf=NULL; if(k->Lc) for(int i=0;iLc[i]); free(k->Rc[i]); } if(k->Ic) for(int i=0;ic.n_layers;i++) free(k->Ic[i]); free(k->Lc); free(k->Rc); free(k->Ic); free(k->kv_start); free(s->hist); @@ -4428,6 +4594,7 @@ static int mux_submit(Model *m, Tok *T, ServeCtx *ctx, ServeReq *req, int nctx, kv_disk_truncate(m,sc->len); } int add=nt-sc->len; if(add>0) memcpy(sc->hist+sc->len,tmp+sc->len,(size_t)add*sizeof(int)); + fprintf(stderr,"[API] KV slot %d prefix %d/%d token, prefill %d\n",sub.slot,sc->len,nt,add); free(tmp); float *logit = add>0 ? step(m,sc->hist+sc->len,add,sc->len) : step(m,sc->hist+sc->len-1,1,sc->len-1); @@ -4465,12 +4632,17 @@ static void run_serve_mux(Model *m, const char *snap){ setvbuf(stdout, NULL, _IONBF, 0); #endif setvbuf(stdin,NULL,_IONBF,0); + intr_install(); /* Ctrl-C = chiudi i turni in volo, non il processo */ printf("\x01\x01READY\x01\x01\nSTAT 0 0.00 0.0 %.2f\n",rss_gb()); fflush(stdout); hwinfo_emit(m); tiers_emit(m); emap_emit(m); int eof=0; for(;;){ + if(g_intr){ g_intr=0; /* stop morbido: ogni request attiva finisce ORA per la + * via normale di mux_done (DONE+stats+KV coerenti) */ + for(int i=0;ilen) #define first (sc->first) char *line=NULL; size_t cap=0; ssize_t nr; char *buf=malloc(1<<16); + intr_install(); /* Ctrl-C = fine turno, non fine processo */ printf("\x01\x01" "READY" "\x01\x01\n"); printf("STAT 0 0.00 0.0 %.2f\n", rss_gb()); fflush(stdout); tiers_emit(m); while((nr=getline(&line,&cap,stdin))>0){ + g_intr=0; /* interruzioni arrivate tra i turni: stantie */ if(nr>0 && line[nr-1]=='\n') line[--nr]=0; if(!strcmp(line,"\x02RESET")){ len=0; first=1; if(m->has_mtp) m->kv_start[m->c.n_layers]=-1; kv_disk_reset(m); @@ -5213,6 +5387,7 @@ int main(int argc, char **argv){ if(g_mmap) fprintf(stderr,"[MMAP] expert = viste zero-copy nei file (page cache = cache)\n"); g_topk = getenv("TOPK")?atoi(getenv("TOPK")):0; g_topp = getenv("TOPP")?atof(getenv("TOPP")):0; + g_expert_budget = getenv("EXPERT_BUDGET")?atoi(getenv("EXPERT_BUDGET")):0; g_cache_route = getenv("CACHE_ROUTE")?atoi(getenv("CACHE_ROUTE")):0; g_route_j = getenv("ROUTE_J")?atoi(getenv("ROUTE_J")):2; g_route_m = getenv("ROUTE_M")?atoi(getenv("ROUTE_M")):12; @@ -5256,7 +5431,13 @@ int main(int argc, char **argv){ g_pilot_k = getenv("PILOT_K")?atoi(getenv("PILOT_K")):(g_pilot_real?6:8); if(g_pilot_k<1) g_pilot_k=1; g_disk_split = getenv("DISK_SPLIT")?atoi(getenv("DISK_SPLIT")):0; /* 1 = split dei disk load nelle stats */ - g_pipe = getenv("PIPE")?atoi(getenv("PIPE")):0; /* default OFF: overlap expert load ‖ matmul (byte-identical; reorders I/O). PIPE=1 opts in */ + g_pipe = getenv("PIPE")?atoi(getenv("PIPE")): +#ifdef _WIN32 + 1 /* default ON: overlap expert load ‖ matmul (byte-identical; reorders I/O). PIPE=0 opts out */ +#else + 0 +#endif + ; g_pipe_nw = getenv("PIPE_WORKERS")?atoi(getenv("PIPE_WORKERS")):8; /* I/O worker threads */ if(g_pipe_nw<1) g_pipe_nw=1; g_direct = getenv("DIRECT")?atoi(getenv("DIRECT")):0; diff --git a/c/openai_server.py b/c/openai_server.py index f31e8f7..ba19e43 100644 --- a/c/openai_server.py +++ b/c/openai_server.py @@ -410,8 +410,11 @@ def generation_options(body, limit): top_p = body.get("top_p") temperature = 0.7 if temperature is None else temperature top_p = 0.9 if top_p is None else top_p - if isinstance(maximum, bool) or not isinstance(maximum, int) or not 1 <= maximum <= limit: - raise APIError(400, f"`{maximum_param}` must be an integer between 1 and {limit}.", maximum_param) + if isinstance(maximum, bool) or not isinstance(maximum, int) or maximum < 1: + raise APIError(400, f"`{maximum_param}` must be a positive integer.", maximum_param) + if maximum > limit: + maximum = limit # clamp to the server's --max-tokens cap instead of 400 (#260): OpenAI + # clients (opencode/ai-sdk) default to large max_tokens; rejecting breaks them. if (isinstance(temperature, bool) or not isinstance(temperature, (int, float)) or not math.isfinite(temperature) or not 0 <= temperature <= 2): raise APIError(400, "`temperature` must be between 0 and 2.", "temperature") diff --git a/c/tests/test_compat_direct.c b/c/tests/test_compat_direct.c index 9b49930..ca9c7b4 100644 --- a/c/tests/test_compat_direct.c +++ b/c/tests/test_compat_direct.c @@ -51,6 +51,23 @@ int main(void){ if(compat_open_direct("no_such_file.tmp")>=0) return fail("open missing file must fail"); if(compat_fsize(-1)>=0) return fail("compat_fsize on bad fd must be negative"); + /* compat_fadvise: WILLNEED warms the page cache (background read into throwaway + * buffer), DONTNEED is a documented no-op. After a WILLNEED the buffered fd's + * subsequent pread must still return the exact bytes — the cache-warmer must not + * corrupt data. Bad fd / non-WILLNEED advice must be safe no-ops (return 0). */ + int wfd = open(TMPF, COMPAT_O_RDONLY); + if(wfd<0) return fail("open buffered for fadvise"); + if(posix_fadvise(wfd, 0, (off_t)FSZ, POSIX_FADV_WILLNEED)!=0) return fail("WILLNEED returned nonzero"); + if(posix_fadvise(wfd, 0, (off_t)FSZ, POSIX_FADV_DONTNEED)!=0) return fail("DONTNEED should be a safe no-op (return 0)"); + if(posix_fadvise(-1, 0, (off_t)FSZ, POSIX_FADV_WILLNEED)!=0) return fail("WILLNEED on bad fd should no-op (return 0)"); + if(posix_fadvise(wfd, 0, 0, POSIX_FADV_WILLNEED)!=0) return fail("WILLNEED with len<=0 should no-op"); + /* verify data integrity through the buffered fd after the cache-warmer ran */ + uint8_t *verify=malloc(FSZ); + if(pread(wfd, verify, FSZ, 0)!=(ssize_t)FSZ) return fail("fadvise: pread size"); + if(memcmp(verify, pat, FSZ)!=0) return fail("fadvise: data corrupted by cache-warmer"); + free(verify); + close(wfd); + close(dfd); compat_aligned_free(buf); free(pat); remove(TMPF); puts("compat direct tests: ok"); diff --git a/c/tests/test_openai_server.py b/c/tests/test_openai_server.py index 0064070..f8aec86 100644 --- a/c/tests/test_openai_server.py +++ b/c/tests/test_openai_server.py @@ -72,8 +72,13 @@ class TemplateTest(unittest.TestCase): def test_validates_generation_limits(self): self.assertEqual(generation_options({"max_tokens": 4, "temperature": 0, "top_p": 1}, 8), (4, 0.0, 1.0)) + # max_tokens above the server cap is clamped, not rejected (#260): OpenAI + # clients default to large values; erroring breaks them. + self.assertEqual(generation_options({"max_tokens": 9, "temperature": 0, "top_p": 1}, 8), + (8, 0.0, 1.0)) + # non-positive / non-int max_tokens is still a hard error with self.assertRaises(APIError): - generation_options({"max_tokens": 9}, 8) + generation_options({"max_tokens": 0}, 8) with self.assertRaises(APIError): generation_options({"temperature": math.nan}, 8) with self.assertRaises(APIError): diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index b8e186e..382125e 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -116,7 +116,7 @@ def layer_idx(name): def classify(name, n_layers, keep_mtp=False, keep_idx=False): if name.endswith("_scale_inv"): return "consumed" # FP8 base: gestito col suo peso - # NVFP4 (modelopt): i sidecar delle scale sono consumati insieme al loro .weight U8. + # NVFP4 (modelopt): i sidecar delle scale sono consumati insieme al loro U8 .weight. # EN: NVFP4 (modelopt): scale sidecars are consumed together with their U8 .weight. if name.endswith((".weight_scale", ".weight_scale_2", ".input_scale")): return "consumed" li = layer_idx(name) @@ -137,7 +137,20 @@ def classify(name, n_layers, keep_mtp=False, keep_idx=False): if name.endswith("norm.weight") or name == "model.norm.weight": return "f32" if name in ("model.embed_tokens.weight", "lm_head.weight"): return "io" if ".mlp.experts." in name and name.endswith(".weight"): return "x" # expert ROUTED (streaming) - if name.endswith(".weight"): return "q" # attn/dense-mlp/shared (residente) + # Split resident weights by type for mixed-precision control: + # "sh" = shared expert (fires on every token, highest sensitivity) + # "o" = o_proj attention (reconstructs output, biggest attn tensor) + # "kvb" = kv_b_proj (reconstructs KV cache on every decode step) + # "attn" = other attention projections (q_a, q_b, kv_a) + # "dmlp" = dense MLP (first 3 layers) + if "shared_experts" in name: return "sh" + if name.endswith("o_proj.weight"): return "o" + if name.endswith("kv_b_proj.weight"): return "kvb" + if any(name.endswith(k) for k in ("q_a_proj.weight", "q_b_proj.weight", + "kv_a_proj_with_mqa.weight")): return "attn" + if any(name.endswith(k) for k in ("mlp.gate_proj.weight", "mlp.up_proj.weight", + "mlp.down_proj.weight")): return "dmlp" + if name.endswith(".weight"): return "q" # fallback: other resident weights return "f32" # ---------- dequant NVFP4 (modelopt) di UN tensore expert -> f32 [O,I] ---------- @@ -202,7 +215,7 @@ def dequant(f, name, keys): return f.get_tensor(name).to(torch.float32).numpy() def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits, - keep_mtp=False, keep_idx=False, group_size=0): + keep_mtp=False, keep_idx=False, group_size=0, bits_map=None): from safetensors import safe_open with safe_open(path, framework="pt") as f: keys = set(f.keys()) @@ -213,7 +226,15 @@ def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits, if kind == "f32": out_dict[name] = w.astype(np.float32) else: - bits = io_bits if kind == "io" else xbits if kind == "x" else ebits + # Resolve bits for this tensor type: use bits_map override if provided, + # otherwise fall back to the classic ebits/xbits/io_bits scheme. + if bits_map and kind in bits_map: + bits = bits_map[kind] + else: + bits = io_bits if kind == "io" else xbits if kind == "x" else ebits + # Any unknown kind that fell through classify as "q" + if bits_map and kind not in bits_map and kind not in ("io", "x", "sh", "o", "kvb", "attn", "dmlp"): + bits = ebits if w.ndim != 2: # es. bias 1D non previsto come 'q' -> tienilo f32 out_dict[name] = w.astype(np.float32); continue if group_size > 0 and bits <= 4: @@ -234,6 +255,18 @@ def main(): ap.add_argument("--ebits", type=int, default=None) # bit residenti (default 4; 8 per --mtp/--indexer) ap.add_argument("--io-bits", type=int, default=8) # bit di embed/lm_head ap.add_argument("--xbits", type=int, default=None) # bit degli expert ROUTED (streaming); default=ebits + # Mixed-precision: per-tensor-type bit overrides. Default = ebits (all same). + # Set these higher to protect sensitive tensors from quantization error. + ap.add_argument("--shared-bits", type=int, default=None, + help="bits for shared expert (fires on every token, highest sensitivity). Default=ebits") + ap.add_argument("--o-bits", type=int, default=None, + help="bits for o_proj attention (reconstructs output, biggest attn tensor). Default=ebits") + ap.add_argument("--kvb-bits", type=int, default=None, + help="bits for kv_b_proj (reconstructs KV cache on every decode). Default=ebits") + ap.add_argument("--attn-bits", type=int, default=None, + help="bits for other attention projections (q_a, q_b, kv_a). Default=ebits") + ap.add_argument("--dmlp-bits", type=int, default=None, + help="bits for dense MLP (first 3 layers). Default=ebits") ap.add_argument("--group-size", type=int, default=0, # 0 = per-row (backward compat); 128 = group-scaled help="group size for int4 scales: 0=per-row (default), 128=one scale per 128 elements (much better quality)") ap.add_argument("--n-layers", type=int, default=78) @@ -255,6 +288,17 @@ def main(): a.ebits = 8 if (a.mtp or a.indexer) else 4 if a.xbits is None: a.xbits = a.ebits + # Build per-type bits map. If a type-specific arg is set, use it; otherwise the + # converter falls back to ebits for that type. + bits_map = {} + if a.shared_bits is not None: bits_map["sh"] = a.shared_bits + if a.o_bits is not None: bits_map["o"] = a.o_bits + if a.kvb_bits is not None: bits_map["kvb"] = a.kvb_bits + if a.attn_bits is not None: bits_map["attn"] = a.attn_bits + if a.dmlp_bits is not None: bits_map["dmlp"] = a.dmlp_bits + if bits_map: + print(f"[MIXED] precision map: " + ", ".join(f"{k}={v}bit" for k,v in sorted(bits_map.items()))) + if a.selftest_nvfp4: import torch # 1) LUT e2m1: i 16 codici devono decodificare esattamente ai valori attesi. @@ -336,7 +380,7 @@ def main(): shards = sorted(glob.glob(os.path.join(a.indir, "*.safetensors"))) from safetensors.numpy import save_file for i, sp in enumerate(shards): - out = {}; convert_shard(sp, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size) + out = {}; convert_shard(sp, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size, bits_map=bits_map) save_file(out, os.path.join(a.outdir, f"out-{i:05d}.safetensors")) # copia config + tokenizer for fn in ["config.json"]: @@ -579,7 +623,7 @@ def main(): if os.path.exists(outp): print(f"[MTP] {outp} already done"); continue print(f"[MTP {i+1}/{len(mtp_shards)}] downloading {sh}...", flush=True) p = download_retry(a.repo, sh, tmp) - out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, keep_mtp=True, group_size=a.group_size) + out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, keep_mtp=True, group_size=a.group_size, bits_map=bits_map) save_file(out, outp) os.remove(p) for blob in glob.glob(os.path.join(tmp, "**", "*"), recursive=True): @@ -599,7 +643,7 @@ def main(): if os.path.exists(outp): continue # gia' fatto -> ripartibile print(f"[IDX {i+1}/{len(idx_shards)}] downloading {sh}...", flush=True) p = download_retry(a.repo, sh, tmp) - out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, keep_idx=True, group_size=a.group_size) + out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, keep_idx=True, group_size=a.group_size, bits_map=bits_map) if out: save_file(out, outp) os.remove(p) for blob in glob.glob(os.path.join(tmp, "**", "*"), recursive=True): @@ -613,7 +657,7 @@ def main(): if os.path.exists(outp): continue # gia' fatto -> ripartibile print(f"[{i+1}/{len(shards)}] downloading {sh} ({free_gb(a.outdir):.0f} GB free)...", flush=True) p = download_retry(a.repo, sh, tmp) - out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size) + out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size, bits_map=bits_map) save_file(out, outp) os.remove(p) # <-- cancella subito lo shard fp8 for blob in glob.glob(os.path.join(tmp, "**", "*"), recursive=True): diff --git a/c/tools/expert_atlas.py b/c/tools/expert_atlas.py deleted file mode 100644 index dd0a2d6..0000000 --- a/c/tools/expert_atlas.py +++ /dev/null @@ -1,141 +0,0 @@ -#!/usr/bin/env python3 -"""Expert Atlas (#175): measure per-expert topic affinity by diffing .coli_usage -across themed probe batches, served through a running colibri API server. - -Protocol per category: snapshot .coli_usage -> send probes -> snapshot again; -the delta is that category's expert-activation spectrum. One engine load total. - -Output: experts.json — for every (layer, expert): counts per category, -normalized affinity, entropy, and a "specialist" label when one topic dominates. - -Usage (server already running with the model): - python3 tools/expert_atlas.py --api http://127.0.0.1:8000 \ - --usage /path/to/model/.coli_usage --out experts.json --ngen 64 -""" -import argparse, json, math, time, urllib.request - -PROBES = { - "code": [ - "Write a Python function that parses a CSV file and returns a dict keyed by the first column.", - "Explain the difference between a mutex and a semaphore, with a C example.", - "Refactor this into idiomatic Rust: for i in range(len(xs)): total += xs[i] * 2", - ], - "math": [ - "Prove that the square root of 2 is irrational.", - "Compute the derivative of x^3 * ln(x) and explain each step.", - "A fair die is rolled 4 times. What is the probability of at least one six?", - ], - "chinese": [ - "请用中文解释一下什么是光合作用,以及它对地球生态系统的重要性。", - "把这句话翻译成中文并解释语法:The early bird catches the worm.", - "写一段关于秋天的短文,一百字左右。", - ], - "english_prose": [ - "Write a vivid paragraph describing an old lighthouse keeper watching a storm arrive.", - "Summarize the plot of Romeo and Juliet in three sentences.", - "Continue this story: The last train left the station, and Maria realized her mistake.", - ], - "science": [ - "Explain how mRNA vaccines work at the cellular level.", - "Why is the sky blue during the day but red at sunset?", - "Describe the life cycle of a massive star, from formation to supernova.", - ], - "law": [ - "Explain the difference between a patent, a trademark, and a copyright.", - "What are the key elements required to form a legally binding contract?", - "Summarize what 'due process' means in constitutional law.", - ], - "poetry": [ - "Write a short poem about a hummingbird in the style of Emily Dickinson.", - "Compose a haiku about winter rain, then explain its imagery.", - "Write four rhyming lines about the sea at night.", - ], - "structured": [ - 'Convert to JSON: name Alice, age 30, hobbies reading and chess, address 5 Oak St.', - "Write a SQL query returning the top 5 customers by total order value, with the schema you assume.", - "Write a regex that matches ISO-8601 dates and explain each part.", - ], - "translation": [ - "Translate into French, German and Spanish: 'Knowledge is the only treasure that grows when shared.'", - "Translate this Italian sentence to English and comment on nuance: 'In bocca al lupo per domani.'", - "Translate into Japanese: 'The meeting has been moved to next Tuesday afternoon.'", - ], - "casual": [ - "Hey! Any tips for staying awake during boring afternoon meetings?", - "What should I cook tonight? I have eggs, rice, tomatoes and some cheese.", - "My friend is always late. How do I tell them it bothers me without being rude?", - ], -} - - -def read_usage(path): - counts = {} - try: - with open(path) as f: - for line in f: - p = line.split() - if len(p) == 3: - counts[(int(p[0]), int(p[1]))] = int(p[2]) - except FileNotFoundError: - pass - return counts - - -def diff(after, before): - return {k: v - before.get(k, 0) for k, v in after.items() if v - before.get(k, 0) > 0} - - -def chat(api, prompt, ngen): - body = json.dumps({"model": "glm-5.2-colibri", "stream": False, "max_tokens": ngen, - "messages": [{"role": "user", "content": prompt}]}).encode() - req = urllib.request.Request(f"{api}/v1/chat/completions", data=body, - headers={"Content-Type": "application/json"}) - with urllib.request.urlopen(req, timeout=600) as r: - json.load(r) - - -def main(): - ap = argparse.ArgumentParser() - ap.add_argument("--api", default="http://127.0.0.1:8000") - ap.add_argument("--usage", required=True, help="path to the model's .coli_usage") - ap.add_argument("--out", default="experts.json") - ap.add_argument("--ngen", type=int, default=64) - a = ap.parse_args() - - spectra = {} - for cat, prompts in PROBES.items(): - before = read_usage(a.usage) - t0 = time.time() - for p in prompts: - chat(a.api, p, a.ngen) - time.sleep(2) # let the engine flush .coli_usage - spectra[cat] = diff(read_usage(a.usage), before) - total = sum(spectra[cat].values()) - print(f"[{cat}] {len(spectra[cat])} experts touched, {total} selections, {time.time()-t0:.0f}s", flush=True) - - cats = list(PROBES.keys()) - experts = {} - for cat, spec in spectra.items(): - for k, v in spec.items(): - experts.setdefault(k, {c: 0 for c in cats})[cat] = v - - atlas = {} - for (layer, eid), counts in experts.items(): - total = sum(counts.values()) - if total < 8: - continue # too few observations to characterise - aff = {c: v / total for c, v in counts.items() if v} - ent = -sum(p * math.log2(p) for p in aff.values()) - top = max(aff, key=aff.get) - label = f"specialist: {top}" if aff[top] >= 0.45 and ent < 2.2 else "generalist" - atlas[f"{layer}:{eid}"] = {"counts": counts, "affinity": {c: round(p, 3) for c, p in aff.items()}, - "entropy": round(ent, 2), "top": top, "label": label} - - spec_n = sum(1 for v in atlas.values() if v["label"].startswith("specialist")) - with open(a.out, "w") as f: - json.dump({"categories": cats, "ngen": a.ngen, "experts": atlas}, f) - print(f"\natlas: {len(atlas)} experts characterised, {spec_n} specialists -> {a.out}") - - -if __name__ == "__main__": - main() diff --git a/c/tools/expert_atlas/README.md b/c/tools/expert_atlas/README.md index 8308d73..b1c78ff 100644 --- a/c/tools/expert_atlas/README.md +++ b/c/tools/expert_atlas/README.md @@ -7,10 +7,16 @@ histogram, and turns them into a per-expert topic-affinity vector. cd c export COLI_MODEL=/path/to/glm52_i4 ./tools/expert_atlas/sweep.sh # 30 probes (10 topics x 3 prompts) -python3 tools/expert_atlas/analyze.py --stats atlas_out/stats --out atlas_out/experts.json +python3 tools/expert_atlas/analyze.py --stats atlas_out/stats --out atlas_out/experts.json \ + --web web/dist/experts.json # optional: feed the web dashboard Atlas python3 tools/expert_atlas/validate.py atlas_out/stats 200 # leave-one-prompt-out check ``` +`--web` writes the same atlas in the shape the web dashboard consumes (the Atlas galaxy and the +Brain hover tooltips): keyed `"layer:expert"` with `affinity`/`entropy`/`top`/`label`. It replaces +the retired `tools/expert_atlas.py`, whose API-driven probing ran through a live server and was +exposed to exactly the traps above (server-side `--topp`, speculative drafts, shared `.coli_usage`). + ## Read this before you trust any atlas Four things silently corrupt this measurement. The sweep script controls all of them; if you diff --git a/c/tools/expert_atlas/analyze.py b/c/tools/expert_atlas/analyze.py index b457be4..c5eb502 100644 --- a/c/tools/expert_atlas/analyze.py +++ b/c/tools/expert_atlas/analyze.py @@ -29,6 +29,7 @@ def main(): ap.add_argument("--min-count", type=int, default=30) ap.add_argument("--min-runs", type=int, default=2, help="must fire in >= this many of the top category's runs") ap.add_argument("--out", default="experts.json") + ap.add_argument("--web", default="", help="also write the web-dashboard experts.json (Atlas/Brain hover)") a = ap.parse_args() # run[(cat,idx)][(layer,expert)] = count ; run_tot[(cat,idx)] = total @@ -121,6 +122,19 @@ def main(): json.dump({"categories": cats, "experts": atlas}, open(a.out, "w"), indent=1) print(f"wrote {a.out}") + if a.web: + # Same atlas, keyed "layer:expert" with per-expert affinity/entropy/top/label — + # the shape the web dashboard consumes (Atlas galaxy, Brain hover). + web = {} + for r in atlas: + aff = {c: v for c, v in r["p"].items() if v > 0} + H = -sum(v * math.log2(v) for v in aff.values()) + web[f"{r['layer']}:{r['expert']}"] = { + "affinity": aff, "entropy": round(H, 2), "top": r["top_topic"], + "label": f"specialist: {r['top_topic']}" if r["spec"] >= 0.5 else "generalist"} + json.dump({"categories": cats, "experts": web}, open(a.web, "w")) + print(f"wrote {a.web} (dashboard format, {len(web):,} experts)") + if __name__ == "__main__": main() diff --git a/c/tools/glm_fp8_emit.py b/c/tools/glm_fp8_emit.py new file mode 100644 index 0000000..45bb304 --- /dev/null +++ b/c/tools/glm_fp8_emit.py @@ -0,0 +1,137 @@ +"""Helper: salva pesi in FP8 e4m3 + scale a blocchi 128x128, nello STESSO layout del +checkpoint reale GLM-5.2-FP8 che `convert_fp8_to_int4.py` legge. + +Layout (deve combaciare col `dequant()` del converter, convert_fp8_to_int4.py:164-169): + - `name` F8_E4M3 [O, I] + - `name_scale_inv` F32 [ceil(O/128), ceil(I/128)] (NOTA: '_scale_inv', underscore) + dequant: W = q.float() * scale.repeat_interleave(128,0).repeat_interleave(128,1)[:O,:I] + +Convenzione FBGEMM/TransformerEngine: scale = amax(blocco)/448 (448 = max e4m3), +si MEMORIZZA il valore e si MOLTIPLICA in dequant. Malgrado il nome "_scale_inv" il +checkpoint memorizza la scala (non il reciproco): e' un MOLTIPLIER. + +EN: Helper that writes weights as FP8 e4m3 with 128x128 block scales, in the SAME layout +EN: as the real GLM-5.2-FP8 checkpoint that `convert_fp8_to_int4.py` reads. +EN: FBGEMM/TransformerEngine convention: scale = amax(block)/448, stored (not its +EN: reciprocal) and MULTIPLIED on dequant. Despite the name "_scale_inv" it is a multiplier. +""" +import torch + +E4M3_MAX = 448.0 # max valore rappresentabile in float8_e4m3fn / max representable value +BLOCK = 128 # granularita' delle scale a blocchi del checkpoint FP8 / FP8 block scale granularity + + +def keep_f32(name, t): + """Stesso set F32 di `classify()` in convert_fp8_to_int4.py (norme, router, bias 1-D). + Tutti gli altri tensori 2-D vengono quantizzati FP8 (attn/mlp/shared/expert/embed/lm_head). + EN: Same F32 set as the converter's classify(): norms, router, 1-D biases. All other 2-D + EN: tensors are FP8-quantized (attn/mlp/shared/expert/embed/lm_head).""" + if t.dim() < 2: + return True # bias 1-D, e_score_correction_bias + if name.endswith("e_score_correction_bias"): + return True + if name.endswith("mlp.gate.weight"): + return True # router (NON gate_proj): tenuto F32 / kept F32 + if name.endswith("norm.weight") or name == "model.norm.weight": + return True # RMSNorm + return False + + +def fp8_block_quantize(w): + """w: [O,I] f32 -> (w_fp8 float8_e4m3fn [O,I], scale_inv f32 [ceil(O/128),ceil(I/128)]). + Identica matematica al `--selftest` del converter (scale = amax(blocco)/448). Padda a + multipli di 128 internamente (gli zeri non alzano l'amax) e fa slice al risultato. + EN: same math as the converter's --selftest. Pads to 128 multiples internally (zeros do + EN: not raise amax), slices the result back to [O,I].""" + O, I = w.shape + nbO, nbI = (O + BLOCK - 1) // BLOCK, (I + BLOCK - 1) // BLOCK + Op, Ip = nbO * BLOCK, nbI * BLOCK + wpad = torch.zeros(Op, Ip, dtype=torch.float32, device=w.device) + wpad[:O, :I] = w + wb = wpad.view(nbO, BLOCK, nbI, BLOCK) # [nbO, BLOCK, nbI, BLOCK] + amax = wb.abs().amax(dim=(1, 3)) # [nbO, nbI] + scale = amax / E4M3_MAX # FBGEMM/TE: memorizza la scala / store the scale + scale = torch.where(scale == 0, torch.ones_like(scale), scale) # blocco tutto-zero -> no div0 + scale = scale.to(torch.float32) + q = (wpad / scale.repeat_interleave(BLOCK, 0).repeat_interleave(BLOCK, 1)).clamp(-E4M3_MAX, E4M3_MAX) + w_fp8 = q.to(torch.float8_e4m3fn) + return w_fp8[:O, :I].contiguous(), scale.contiguous() + + +def fp8_block_dequantize(w_fp8, scale): + """Esatto inverso di fp8_block_quantize, e identico al `dequant()` del converter. + EN: exact inverse of fp8_block_quantize, identical to the converter's dequant().""" + O, I = w_fp8.shape + qf = w_fp8.to(torch.float32) + return qf * scale.repeat_interleave(BLOCK, 0).repeat_interleave(BLOCK, 1)[:O, :I] + + +def unfuse_experts(sd): + """Split HF's fused 3-D `experts.gate_up_proj` [E, 2*M, I] into per-expert 2-D + `experts.{e}.gate_proj` [M, I] + `experts.{e}.up_proj` [M, I], and + `experts.down_proj` [E, I, M] -> `experts.{e}.down_proj` [M_out, I]. + + The real GLM-5.2-FP8 checkpoint stores experts UNFUSED as per-expert 2-D tensors + (gate_proj, up_proj, down_proj), each with its own _scale_inv. HF's + GlmMoeDsaForCausalLM fuses gate+up into a single 3-D gate_up_proj for efficiency. + The converter (classify + ndim!=2 guard) and the C engine both expect the unfused + layout, so we split before saving. + + Idempotent: if experts are already unfused (no 3-D gate_up_proj), returns sd as-is. + EN: split HF's fused 3-D expert weights into the per-expert 2-D layout that the real + EN: checkpoint uses and the converter/engine expect. No-op if already unfused.""" + keys_to_remove = [] + new_entries = {} + for name, t in sd.items(): + if not name.endswith(".mlp.experts.gate_up_proj"): + continue + # prefix = everything before ".mlp.experts.gate_up_proj" + prefix = name[:-len(".mlp.experts.gate_up_proj")] + E, twoM, I = t.shape # [E, 2*intermediate, input] + M = twoM // 2 + for e in range(E): + new_entries[f"{prefix}.mlp.experts.{e}.gate_proj.weight"] = t[e, :M, :].contiguous() + new_entries[f"{prefix}.mlp.experts.{e}.up_proj.weight"] = t[e, M:, :].contiguous() + keys_to_remove.append(name) + # down_proj may be 3-D [E, I, M] in the fused form, or already per-expert + for name, t in sd.items(): + if not name.endswith(".mlp.experts.down_proj") or t.dim() != 3: + continue + prefix = name[:-len(".mlp.experts.down_proj")] + E = t.shape[0] + for e in range(E): + new_entries[f"{prefix}.mlp.experts.{e}.down_proj.weight"] = t[e].contiguous() + keys_to_remove.append(name) + for k in keys_to_remove: + sd.pop(k, None) + sd.update(new_entries) + return sd + + +def state_dict_to_fp8(sd): + """Converte uno state_dict HuggingFace nel layout FP8 del checkpoint reale: + per ogni tensore quantizzabile 2-D scrive `{name}` (F8_E4M3) + `{name}_scale_inv` (F32); + norme/router/bias e qualsiasi tensore NON 2-D (es. pesi MLA impaccati 3-D) restano nel + dtype originale. Questo rispecchia il guard `w.ndim != 2 -> f32` del converter + (convert_fp8_to_int4.py:184). EN: builds the real-checkpoint FP8 layout. Only exactly-2-D + tensors are FP8-quantized; anything else (1-D, 3-D packed MLA weights, ...) is kept, exactly + like the converter's `ndim != 2 -> f32` guard.""" + out = {} + for name, t in sd.items(): + if keep_f32(name, t) or t.dim() != 2: + out[name] = t # f32 / 1-D / 3-D+: tieni / keep + else: + w_fp8, scale = fp8_block_quantize(t.float()) + out[name] = w_fp8 + out[name + "_scale_inv"] = scale + return out + + +def save_fp8_safetensors(sd, path): + """Quantizza a blocchi FP8 e salva in un singolo safetensors leggibile dal converter + via `--indir`. EN: block-quantize to FP8 and save a single safetensors for the converter.""" + from safetensors.torch import save_file + out = state_dict_to_fp8(sd) + save_file({k: v.contiguous() for k, v in out.items()}, str(path)) + n_fp8 = sum(1 for v in out.values() if v.dtype == torch.float8_e4m3fn) + return n_fp8, len(out) diff --git a/c/tools/make_glm_bench_model.py b/c/tools/make_glm_bench_model.py index 3aa1ce9..5ec3457 100644 --- a/c/tools/make_glm_bench_model.py +++ b/c/tools/make_glm_bench_model.py @@ -3,15 +3,27 @@ This is not a useful language model. It preserves the real glm_moe_dsa data flow while remaining small enough to generate locally and run repeated CPU/CUDA A/B tests without downloading the 379 GB checkpoint. + +With --fp8 the weights are written as FP8 e4m3 + 128x128 block scale_inv, in the +SAME layout as the real GLM-5.2-FP8 checkpoint, so convert_fp8_to_int4.py can +exercise its FP8->int4 dequant path on a local fixture (its dims are 128-friendly, +so this is also the right fixture for --group-size 128 testing): + + python tools/make_glm_bench_model.py --fp8 --output glm_bench_fp8 + python tools/convert_fp8_to_int4.py --indir glm_bench_fp8 --outdir glm_bench_i4 --ebits 4 --group-size 128 """ import argparse import json +import sys from pathlib import Path import torch from transformers import GlmMoeDsaConfig, GlmMoeDsaForCausalLM +sys.path.insert(0, str(Path(__file__).resolve().parent)) # importa glm_fp8_emit se lanciato da c/ +from glm_fp8_emit import save_fp8_safetensors, unfuse_experts + def build_config() -> GlmMoeDsaConfig: return GlmMoeDsaConfig( @@ -51,6 +63,9 @@ def main() -> None: parser.add_argument("--output", default="glm_bench_medium") parser.add_argument("--device", default="cuda" if torch.cuda.is_available() else "cpu") parser.add_argument("--seed", type=int, default=1234) + parser.add_argument("--fp8", action="store_true", + help="write weights as FP8 e4m3 + 128x128 block scale_inv (same layout as " + "GLM-5.2-FP8) instead of bf16, so convert_fp8_to_int4.py can dequant+requant") args = parser.parse_args() torch.manual_seed(args.seed) @@ -70,7 +85,6 @@ def main() -> None: output = Path(args.output) output.mkdir(parents=True, exist_ok=True) params = sum(p.numel() for p in model.parameters()) - model.save_pretrained(output, safe_serialization=True, max_shard_size="4GB") model.to(args.device) prompt = [3, 14, 159, 26, 53, 58, 200, 11, 77, 240, 5, 99] @@ -79,6 +93,25 @@ def main() -> None: full = model.generate(ids, max_new_tokens=8, do_sample=False, use_cache=True)[0] logits = model(full.unsqueeze(0), use_cache=False).logits[0] + # Unfuse experts AFTER reference generation (model needs fused weights for + # forward/generate) but BEFORE saving — the real checkpoint and the converter + # + C engine all expect per-expert 2-D gate_proj/up_proj/down_proj tensors. + sd = model.state_dict() + unfuse_experts(sd) + + if args.fp8: + n_fp8, n_tot = save_fp8_safetensors(sd, output / "model.safetensors") + # save_pretrained scrive config.json; nel path FP8 lo bypassiamo, quindi lo scriviamo + # a mano (serve al converter e al motore C). EN: save_pretrained writes config.json; + # the FP8 path bypasses it, so write it manually (converter + C engine need it). + (output / "config.json").write_text(json.dumps(cfg.to_dict())) + print(f"saved FP8: {n_fp8} e4m3 tensors (+{n_tot - n_fp8} scale_inv sidecars / f32) " + f"-> {output / 'model.safetensors'}") + else: + from safetensors.torch import save_file + save_file({k: v.contiguous() for k, v in sd.items()}, str(output / "model.safetensors")) + (output / "config.json").write_text(json.dumps(cfg.to_dict())) + ref = { "prompt_ids": prompt, "full_ids": full.cpu().tolist(), @@ -89,6 +122,7 @@ def main() -> None: "seed": args.seed, "parameters": params, "parameters_billions": round(params / 1e9, 4), + "format": "fp8-e4m3-128" if args.fp8 else "bf16", "purpose": "backend benchmark fixture; random weights, not a language model", } (output / "bench_manifest.json").write_text(json.dumps(manifest, indent=2)) diff --git a/c/tools/make_glm_oracle.py b/c/tools/make_glm_oracle.py index 65a1c19..b623faf 100644 --- a/c/tools/make_glm_oracle.py +++ b/c/tools/make_glm_oracle.py @@ -3,10 +3,34 @@ Architettura vera (MLA + DSA indexer + router sigmoid/noaux_tc + shared expert), dimensioni minuscole. Salva pesi+config in c/glm_tiny/ e un riferimento greedy in c/ref_glm.json. seq corta (<= index_topk) cosi' il DSA seleziona tutte le key e l'attenzione coincide con la MLA densa: il motore C puo' validare senza implementare -l'indexer sparso.""" -import json, torch +l'indexer sparso. + +--fp8: salva i pesi come FP8 e4m3 + scale a blocchi 128x128 (layout del checkpoint reale +GLM-5.2-FP8) invece di bf16, cosi' convert_fp8_to_int4.py puo' esercitare il path FP8->int4 +su un modello minuscolo. PRIMA di calcolare ref_glm.json fa il round-trip dei pesi per FP8 +(quant->dequant, copy_ nel modello): cosi' il riferimento riflette ESATTAMENTE il modello +FP8 che il converter legge, non il modello bf16 a precisione piena. Default: bf16 (oracolo +originale invariato). +EN: --fp8 writes FP8 e4m3 + 128x128 block scale_inv (real GLM-5.2-FP8 layout) instead of bf16, +EN: so convert_fp8_to_int4.py can run its FP8->int4 path on a tiny model. ref_glm.json is +EN: computed AFTER the FP8 round-trip, so the reference matches exactly what the converter +EN: ingests. Default: bf16 (original oracle unchanged).""" +import json, sys, argparse +from pathlib import Path +import torch from transformers import GlmMoeDsaConfig, GlmMoeDsaForCausalLM +sys.path.insert(0, str(Path(__file__).resolve().parent)) # importa glm_fp8_emit se lanciato da c/ +from glm_fp8_emit import (fp8_block_quantize, fp8_block_dequantize, keep_f32, + save_fp8_safetensors, unfuse_experts) + +ap = argparse.ArgumentParser() +ap.add_argument("--fp8", action="store_true", + help="salva in FP8 e4m3 + 128x128 block scale_inv (layout GLM-5.2-FP8) e " + "calcola ref_glm.json sul modello dopo il round-trip FP8. " + "EN: write FP8 e4m3 + block scale_inv, ref computed on FP8-rounded model") +args = ap.parse_args() + torch.manual_seed(1234) cfg = GlmMoeDsaConfig( @@ -53,6 +77,18 @@ with torch.no_grad(): layer.mlp.gate.e_score_correction_bias.copy_( torch.linspace(-0.1, 0.1, cfg.n_routed_experts)) +# --fp8: round-trip dei pesi quantizzabili per FP8 PRIMA di calcolare il riferimento, +# cosi' ref_glm.json riflette esattamente il modello FP8 che il converter leggera'. +# Norme/router/bias (keep_f32) restano a precisione piena. EN: --fp8: round-trip quantizable +# weights through FP8 before computing the reference, so ref_glm.json matches the FP8 model. +if args.fp8: + with torch.no_grad(): + for n, p in model.named_parameters(): + if keep_f32(n, p) or p.dim() != 2: + continue + q, s = fp8_block_quantize(p) + p.copy_(fp8_block_dequantize(q, s)) + print("=== state_dict tensors (names used by the C loader) ===") for n, p in model.state_dict().items(): print(f" {n:60s} {tuple(p.shape)}") @@ -73,7 +109,20 @@ with torch.no_grad(): tf_pred = lg.argmax(-1).tolist() print("tf_pred:", tf_pred) -model.save_pretrained("glm_tiny", safe_serialization=True) +# Unfuse experts AFTER reference generation (model needs fused weights for +# forward/generate) but BEFORE saving — the real checkpoint and the converter +# + C engine all expect per-expert 2-D gate_proj/up_proj/down_proj tensors. +sd = model.state_dict() +unfuse_experts(sd) + +if args.fp8: + n_fp8, n_tot = save_fp8_safetensors(sd, "glm_tiny/model.safetensors") + print(f"\nsaved FP8: {n_fp8} e4m3 tensors (+{n_tot - n_fp8} scale_inv sidecars / f32) " + f"-> glm_tiny/model.safetensors") +else: + from safetensors.torch import save_file + save_file({k: v.contiguous() for k, v in sd.items()}, "glm_tiny/model.safetensors") json.dump(cfg.to_dict(), open("glm_tiny/config.json", "w")) json.dump({"prompt_ids": prompt, "full_ids": full, "tf_pred": tf_pred}, open("ref_glm.json", "w")) -print("\nsaved: glm_tiny/ (weights + config) and ref_glm.json") +print("saved: glm_tiny/ (weights + config) and ref_glm.json" + + (" [fp8]" if args.fp8 else "")) diff --git a/c/tools/quant_ablation.py b/c/tools/quant_ablation.py index 7a870ff..689643a 100644 --- a/c/tools/quant_ablation.py +++ b/c/tools/quant_ablation.py @@ -106,33 +106,106 @@ def rotation(dim, device, seed=417): return q -def quantize_param(w, bits, group, rot=False): +def quantize_param(w, bits, group, rot=False, e8=""): if w.ndim == 3: # fused experts [E, in, out] -> move input last x = w.transpose(1, 2).contiguous() - x = _rot_quant(x, bits, group) if rot else _quant_last_dim(x, bits, group) + x = _rot_quant(x, bits, group, e8) if rot else _grid_or_e8(x, bits, group, e8) return x.transpose(1, 2).contiguous() if rot: - return _rot_quant(w, bits, group) - return _quant_last_dim(w, bits, group) # nn.Linear [out, in] -- input already last + return _rot_quant(w, bits, group, e8) + return _grid_or_e8(w, bits, group, e8) # nn.Linear [out, in] -- input already last -def _rot_quant(x, bits, group): +def _grid_or_e8(x, bits, group, e8): + if e8: + return _quant_e8(x.float(), group, ball=(e8 == "-e8")) + return _quant_last_dim(x, bits, group) + + +def _rot_quant(x, bits, group, e8=""): """W -> Qn(W@Q) @ Q^T along the last (input) dim — see rotation() above.""" q = rotation(x.shape[-1], x.device) - return (_quant_last_dim(x.float() @ q, bits, group) @ q.T).contiguous() + return (_grid_or_e8(x.float() @ q, bits, group, e8) @ q.T).contiguous() -SCHEME_RE = re.compile(r"^int(2|3|4|8)(?:-g(\d+))?(-rot)?(-nohead)?$") +# -------------------------------------------------------------------------------------- +# E8 lattice quantization (#81 follow-up): the -rot schemes above are QuaRot (rotation + +# uniform grid). QuIP#'s 2-bit result needs the second ingredient — an E8 lattice codebook +# instead of the grid. E8 = D8 ∪ (D8 + 1/2), nearest point via Conway–Sloane: round every +# coordinate, and if the sum is odd re-round the worst coordinate the other way; repeat on +# the half-shifted copy and keep the closer of the two. `-e8` clamps points to |p|^2 <= 10, +# the E8P ball QuIP# builds its 2^16 codebook from (2 bits/weight for 8-dim blocks); +# `-e8u` leaves the lattice unbounded — an ideal-codebook upper bound, not a deployable rate. +# Scale: per group, a small MSE search over multiples of the block RMS (absmax is the wrong +# statistic for a lattice — the ball wants energy matched, not the peak). +# -------------------------------------------------------------------------------------- +def _d8_nearest(y): + f = torch.round(y) + d = y - f + odd = (f.sum(-1).long() & 1).bool() + idx = d.abs().argmax(-1, keepdim=True) + step = torch.where(d.gather(-1, idx) >= 0, 1.0, -1.0) + flipped = f.gather(-1, idx) + step + return f.scatter(-1, idx, torch.where(odd[..., None], flipped, f.gather(-1, idx))) + + +def _e8_nearest(y): + a = _d8_nearest(y) + b = _d8_nearest(y - 0.5) + 0.5 + da = ((y - a) ** 2).sum(-1, keepdim=True) + db = ((y - b) ** 2).sum(-1, keepdim=True) + return torch.where(da <= db, a, b) + + +def _e8_ball(y, r2=10.0): + p = _e8_nearest(y) + for _ in range(8): # shrink-and-requantize until inside + n2 = (p ** 2).sum(-1, keepdim=True) + over = n2 > r2 + 1e-6 + if not over.any(): + break + y = torch.where(over, y * torch.sqrt(r2 / torch.clamp(n2, min=r2)) * 0.98, y) + p = torch.where(over, _e8_nearest(y), p) + return p + + +def _quant_e8(x, group, ball): + """Blocks of 8 along the input dim; per-group scale by MSE search over RMS multiples.""" + if x.shape[-1] % 8: + raise SystemExit(f"-e8 needs input dim divisible by 8 (got {x.shape[-1]})") + g = group or x.shape[-1] + if g % 8: + raise SystemExit(f"-e8 group {g} must be a multiple of 8") + shp = x.shape + xg = x.reshape(-1, g) # [G, g] + rms = torch.clamp(xg.pow(2).mean(-1, keepdim=True).sqrt(), min=1e-8) + best_out, best_err = None, None + for k in (0.5, 0.7, 0.9, 1.1, 1.4, 1.8, 2.4): + s = rms * k + yb = (xg / s).reshape(-1, g // 8, 8) + p = _e8_ball(yb) if ball else _e8_nearest(yb) + out = (p.reshape(-1, g) * s) + err = (out - xg).pow(2).sum(-1, keepdim=True) + if best_err is None: + best_out, best_err = out, err + else: + take = err < best_err + best_out = torch.where(take, out, best_out) + best_err = torch.where(take, err, best_err) + return best_out.reshape(shp) + + +SCHEME_RE = re.compile(r"^int(2|3|4|8)(?:-g(\d+))?(-e8u?)?(-rot)?(-nohead)?$") def parse_scheme(name): - """'int4-g128-nohead' -> (bits=4, group=128, skip_head=True). 'fp16' -> None.""" + """'int4-g128-nohead' -> (bits, group, e8, skip_head...). 'fp16' -> None.""" if name == "fp16": return None m = SCHEME_RE.match(name) if not m: - raise SystemExit(f"bad scheme '{name}' (expected fp16 | int{{2,3,4,8}}[-g][-rot][-nohead])") - return int(m.group(1)), int(m.group(2) or 0), bool(m.group(3)), bool(m.group(4)) + raise SystemExit(f"bad scheme '{name}' (expected fp16 | int{{2,3,4,8}}[-g][-e8|-e8u][-rot][-nohead])") + return int(m.group(1)), int(m.group(2) or 0), m.group(3) or "", bool(m.group(4)), bool(m.group(5)) def is_router(name): @@ -152,7 +225,7 @@ def apply_scheme(model, scheme): spec = parse_scheme(scheme) if spec is None: return 0, 0, total - bits, group, rot, skip_head = spec + bits, group, e8, rot, skip_head = spec n = qp = 0 with torch.no_grad(): for name, p in model.named_parameters(): @@ -160,7 +233,7 @@ def apply_scheme(model, scheme): continue if skip_head and is_head_or_embed(name): continue - p.data.copy_(quantize_param(p.data.float(), bits, group, rot).to(p.dtype)) + p.data.copy_(quantize_param(p.data.float(), bits, group, rot, e8).to(p.dtype)) n += 1 qp += p.numel() return n, qp, total diff --git a/issue_budget.md b/issue_budget.md new file mode 100644 index 0000000..fe0ef86 --- /dev/null +++ b/issue_budget.md @@ -0,0 +1,150 @@ +## TL;DR + +New `EXPERT_BUDGET=N` env var that caps the number of **distinct experts loaded per layer** across the batch-union. When the union exceeds the budget, keeps only the highest-aggregate-gate-weight experts and drops the rest — they're never loaded from disk. On a 24 GB RAM host (cache `cap=2`), `EXPERT_BUDGET=4` nearly **doubles decode tok/s** (0.18 → 0.33) and **4x's prefill speed** (38.7s → 8.9s). Based on MoE-Spec (arXiv 2602.16052): "top 32 of 64 experts capture 93% of routing weight." + +Branch: `experiment/expert-budget` (based on latest `dev` at `62419af`) + +--- + +## The problem + +Every expert miss costs ~19 MB of disk I/O. On low-RAM hosts where the LRU cache cap is tiny (e.g. `cap=2` on 24 GB RAM), nearly every routed expert is a miss. With `topk=8` and 75 sparse layers, a single-token decode reads ~8.5 GB of experts from disk. Under MTP with `S=4`, the batch-union can produce 20-32 distinct experts per layer — almost all misses — multiplying disk reads further. + +The README documents this directly: +> "on a cold cache each verified draft routes to extra experts (~660 → ~1100 expert-loads/token)" + +The existing `TOPP` env var trims experts *within a single position's top-K*. But it cannot reduce the **cross-position union** — the deduplication across batch positions (prefill, MTP verification) that multiplies disk loads. That's the gap this fills. + +## How it works + +The batch-union in `moe()` (line ~2306 of `glm.c`) deduplicates routed experts across all `S` positions into `uniq[0..nu)`. After the union is built but before the cache-resolve/load loop, the budget cap kicks in: + +``` +if(EXPERT_BUDGET > 0 && nu > EXPERT_BUDGET): + 1. Compute aggregate gate weight per unique expert + (sum of ws[s*K+kk] across all positions that route to it) + 2. Sort experts by descending aggregate weight + 3. Keep top EXPERT_BUDGET, mark rest as dropped + 4. Remove dropped experts from each position's idxs[]/keff[] + (decrement keff, renormalize remaining weights if norm_topk) + 5. Compact uniq[] to kept experts only +``` + +Dropped experts are removed from `idxs[]` entirely, so they're **never resolved, never loaded from disk, never computed**. The downstream code already tolerates `keff[s] < K` (the `TOPP` path produces the same state), so this propagates correctly through resolve, matmul, and LRU promotion. + +**Relationship to existing features:** +- `TOPP` trims within one position's top-K (per-position) +- `EXPERT_BUDGET` trims across the cross-position union (per-layer) +- They compose: `TOPP=0.8 EXPERT_BUDGET=12` first trims each position to its top-p mass, then caps the union + +## Code change + +**`c/glm.c`** — 54 lines added, single file: + +1. Global + counter (near existing `g_topk`/`g_topp`): +```c +static int g_expert_budget=0; /* EXPERT_BUDGET=N */ +static int64_t g_budget_dropped=0; /* total experts dropped */ +``` + +2. Env var parsing (near existing `TOPK`/`TOPP`): +```c +g_expert_budget = getenv("EXPERT_BUDGET")?atoi(getenv("EXPERT_BUDGET")):0; +``` + +3. Budget cap in `moe()` batch-union (after `uniq[]` is built, before resolve loop) — full logic described above. + +4. Stats reporting alongside existing `TOPK`/`TOPP` line: +``` +EXPERT_BUDGET=4 (dropped 13613 experts, ~257.3 GB I/O saved) +``` + +## Measurements + +**Test setup:** GLM-5.2 744B int4, 24 GB RAM (cache `cap=2`), Core Ultra 9 185H (AVX-VNNI), MTP=0, single-token decode, 32 tokens generated, same prompt: *"Explain the concept of recursion in programming. Provide a simple example."* + +| Config | tok/s | vs baseline | hit rate | prefill | decode | experts dropped | I/O saved | +|--------|-------|-------------|----------|---------|--------|-----------------|-----------| +| **Baseline** (budget=0) | 0.18 | — | 9.3% | 38.7s | 176.3s | 0 | 0 | +| EXPERT_BUDGET=12 | 0.19 | +5% | 14.0% | 12.3s | 171.3s | 3,313 | 62.6 GB | +| EXPERT_BUDGET=6 | 0.26 | +44% | 21.0% | 7.5s | 122.5s | 8,543 | 161.5 GB | +| **EXPERT_BUDGET=4** | **0.33** | **+83%** | 16.4% | 8.9s | **97.4s** | 13,613 | 257.3 GB | + +### Profile breakdown (prefill) + +| Metric | Baseline | Budget=4 | +|--------|----------|----------| +| expert-disk service | 28.4s | **3.5s** (8x less) | +| expert-matmul | 7.1s | 1.4s | +| attention | 2.3s | 2.8s | + +### Profile breakdown (decode, 32 tokens) + +| Metric | Baseline | Budget=4 | +|--------|----------|----------| +| expert-disk service | 133.4s | proportional ~65s | +| expert-matmul | 23.9s | ~12s | +| decode total | 176.3s | **97.4s** | + +### Key observations + +1. **Prefill is the biggest winner.** With S=14 positions, the batch-union has 50+ distinct experts per layer. Budget=12 cuts that to 12, saving 40+ × 19 MB × 75 layers of disk reads. Prefill goes from 38.7s to 8.9s — **4.4x faster**. + +2. **Decode improvement scales with budget tightness.** Budget=12 barely constrains single-token decode (S=1 → max 8 experts < 12), so decode is barely affected. Budget=4 actually halves the decode load (8 → 4 experts/layer), nearly doubling tok/s. + +3. **Hit rate improves** because fewer experts compete for the tiny cache (cap=2). Budget=6 hit 21% vs baseline 9.3% — more than doubled. + +4. **No crashes or instability** at any budget value. + +## Quality impact — honest assessment + +**Budget=4 keeps only the top-4 of 8 routed experts per layer.** The dropped 4 experts had the lowest gate weights — they contributed the least to the output. But this IS a quality trade-off. + +On a **cold cache** (our test setup), quality assessment is confounded: both baseline and budget outputs are already garbled from int4 quantization + cold-cache routing. Sample comparison: + +- **Baseline output:** *"Hire Some To Take Object-Oriented Programming Assignment"* +- **Budget=4 output:** *"The world is a dangerous place, not so much because of the small percentage of people who are doing to do the little of people who are doing to"* + +Both are incoherent — the cold cache at int4 on 24 GB RAM produces poor output regardless of budget. The budget=4 output is **not dramatically worse** than the already-poor baseline — they're both garbled, just differently garbled. A proper quality assessment needs a **warm cache** where the baseline produces coherent text, so you can isolate the degradation from dropping experts. + +**Expected quality on warm cache:** With `norm_topk=1` (GLM-5.2 renormalizes gate weights), the top-4 experts typically capture ~80-85% of routing weight (the remaining 4 share 15-20%). Output should be mostly coherent with occasional word-choice degradation. Budget=6 (top-6 of 8, ~90%+ weight captured) should be nearly indistinguishable from baseline. + +**Recommendation for users:** +- `EXPERT_BUDGET=6-8` on cold/low-RAM hosts — good speedup, minimal quality loss +- `EXPERT_BUDGET=4` on very-low-RAM hosts where speed matters more than quality +- Leave OFF on high-RAM hosts where all experts are resident anyway (budget never triggers) + +## Safety + +- **Default OFF** (`EXPERT_BUDGET=0`) — zero behavior change unless explicitly set +- **Opt-in quality trade-off** — same design philosophy as `TOPP`: the user explicitly chooses speed over quality +- **No output corruption** — dropped experts' contributions are omitted and remaining weights renormalized (identical math to `TOPP`) +- **Compatible with all features** — PILOT, MTP, CUDA, CACHE_ROUTE, PIPE, TOPP all work unchanged; they just see fewer experts in the union +- **No crash risk** — no new memory allocation patterns, no new I/O, purely a filter on existing data structures + +## Reproduce + +```bash +make ARCH=native + +# Baseline: +SNAP= MTP=0 PROMPT="Explain recursion in programming." NGEN=32 ./glm 64 + +# With budget: +SNAP= MTP=0 EXPERT_BUDGET=4 PROMPT="Explain recursion in programming." NGEN=32 ./glm 64 + +# With MTP (where budget saves the most): +SNAP= MTP=1 EXPERT_BUDGET=16 PROMPT="Explain recursion in programming." NGEN=32 ./glm 64 +``` + +The stats line prints `EXPERT_BUDGET=N (dropped X experts, ~Y GB I/O saved)`. + +## What's needed before merge + +1. **Warm-cache quality A/B test** on a host with enough RAM (cap>=16) to produce coherent baseline output, then compare budget=4/6/8 text quality +2. **MTP interaction test** — verify that budget + MTP composes correctly and doesn't crash when draft verification routes to budgeted-away experts +3. **Decide on defaults** — should this auto-activate on low-RAM hosts (like `cap` auto-lowering)? My recommendation: no, leave it OFF and document the recommended value per RAM tier + +## Prior art + +**MoE-Spec** (arXiv 2602.16052) — "Expert Budgeting for Efficient Speculative Decoding": Training-free expert budgeting at verification time. Key finding: "top 32 of 64 experts capture 93% of routing weight." Our approach applies the same principle but at the batch-union level (not just MTP verification), making it effective for prefill and single-token decode too. diff --git a/issue_diskio.md b/issue_diskio.md new file mode 100644 index 0000000..e9403ec --- /dev/null +++ b/issue_diskio.md @@ -0,0 +1,242 @@ +# Disk I/O Minimization — Research + +Branch: `experiment/diskio-research` (based on `dev` at `62419af`) + +## TL;DR + +The engine's disk I/O is **already well-engineered on the hottest path** (expert streaming uses coalesced O_DIRECT `pread` + `posix_fadvise` hints + LRU + pin cache + speculative prefetch). There are **4 concrete, bounded opportunities** to shave latency, ranked by ROI: + +| # | Opportunity | Where | Frequency | Estimated win | +|---|---|---|---|---| +| 1 | **KV-cache write batching** (157 fwrites/token → 1) | `kv_disk_append` | per turn | cuts ~100s of syscalls/turn | +| 2 | **`/proc/meminfo` fopen storm** | `rss_gb()` | ~every 16 tokens (Linux) | eliminates recurring open/read/close | +| 3 | **Expert prefetch on Windows** (`PrefetchVirtualMemory`) | `expert_prefetch` | per miss | mmap path is Linux/macOS-only today | +| 4 | **KV-cache: buffered handle kept open** | `kv_disk_append` | per turn | kills open+fseek+close per turn | + +There are also **2 non-opportunities** worth recording so we don't re-investigate: O_DIRECT for experts (correctly used today), and PagedAttention-style file layout (already single-file + indexed). + +--- + +## How the engine does disk I/O today + +There are **three I/O stacks**, behaving very differently: + +| Stack | Mechanism | Frequency | Files | +|---|---|---|---| +| **Expert weights** (hottest) | `pread` on kept-open fds + `posix_fadvise`, optional `mmap` | per miss, every token | `st.h`, `glm.c:1328` | +| **KV cache** (`.coli_kv`) | `fopen` + `fwrite`/`fread` | per turn | `glm.c:3812-3889` | +| **Everything else** (config, tokenizer, stats, grammar) | `fopen` + `fread` | startup-only | scattered | + +### Expert path (the hot path — already good) + +`expert_load` (`glm.c:1328-1481`) has three sub-paths: + +- **Default `pread` path** (`glm.c:1385-1472`): coalesces the 3 contiguous expert tensors (gate/up/down) into **one ~19 MB O_DIRECT `pread`** into a 16K-aligned slab (`glm.c:1447`). Falls back to 3 separate `pread`s only if non-contiguous. Scales are 3 tiny separate `pread`s (kilobytes). `posix_fadvise(DONTNEED)` evicts pages after if `g_drop`. **This is well-batched — one syscall for ~19 MB.** +- **`COLI_MMAP=1` path** (`glm.c:1352-1383`): `mmap` per shard fd (cached), `madvise(WILLNEED)` + synchronous page-touch loop. Zero-copy. **Default OFF, and Linux/macOS/FreeBSD-only** — no `MapViewOfFile` on Windows. +- **Prefetch hints**: `expert_prefetch` (`glm.c:1602-1609`) → `st_prefetch` (`st.h:178`) issues `posix_fadvise(WILLNEED)` — readahead hint only, no data read. Called from `moe` next-64-block lookahead, pilot, and SPEC. + +### KV cache persistence (per turn — opportunity here) + +`kv_disk_append` (`glm.c:3834-3855`), called once per turn: +1. `fopen("r+b")` — **reopens the file every turn** +2. `fseek` to append position +3. **per-position loop**: for each new token, `fwrite` the token i32, then **2 fwrites per layer** (Lc + Rc) + optional DSA Ic. With 78 layers that's **~157 fwrites per token appended**. +4. `fflush` (userspace only — **no fsync/fdatasync anywhere in the codebase**) +5. `fseek` back to header + `fwrite` the new nrec counter (crash-safe ordering) +6. `fclose` + +Record size ~182 KB/token. On a long first turn this is **tens of thousands of small fwrites**. stdio buffering coalesces them into fewer `write` syscalls, but the userspace overhead remains. + +### Recurring surprise: `/proc/meminfo` + +`rss_gb()` (`glm.c:4625`) does `fopen("/proc/meminfo")` + fgets + fclose. Called from every STAT line and every 16-token heartbeat (`glm.c:3473, 3477`). On Linux this is an **open+read+close of procfs ~every 16 tokens**. (Windows uses `compat_meminfo`, no file — not affected.) + +--- + +## What similar projects do + +**llama.cpp** (the reference): `mmap`s the entire model read-only, uses `--mlock` to pin hot pages, streams layers to GPU via partial offload, and issues per-pass readahead of upcoming tensors (`llama-mmap.cpp`). Justine Tunney's mmap work: "load 100× faster using half as memory." Crucial finding from discussion #18758: **for MoE, mmap beats O_DIRECT** when the model fits in ~RAM — O_DIRECT takes "at least 10× longer" on repeated loads because it bypasses the page cache that serves re-faults for free. + +**The general consensus across llama.cpp, vLLM, AirLLM, PRESERVE, HOBBIT, SolidAttention (FAST '26):** +- mmap + OS page cache as the backing store for an LRU is the proven recipe +- prefetch the *next* expert/layer while computing the current one — this is where the 0.5ms lives +- single indexed file (one `open()`) beats one-file-per-expert +- align tensors to 4KB (preferably 64KB) for clean page-fault boundaries + SSD geometry +- buffer sweet spot ~1MB; syscall cost ~1-5µs each, so batching matters at high repetition + +--- + +## The 4 opportunities (ranked) + +### Opportunity 1 — KV-cache write batching (HIGH ROI, LOW risk) + +**Problem:** `kv_disk_append` does ~157 `fwrite` calls per appended token (1 token i32 + 2×78 layers). stdio buffering hides some of this, but on a long first-turn prefill (hundreds-thousands of tokens) this is tens of thousands of fwrites. + +**Fix:** Build one contiguous record in a heap buffer (token + all layers' Lc/Rc/Ic for that position), then **a single `fwrite` per position** (or even one `fwrite` for the whole turn). The data is already laid out contiguously in memory per-layer (`coli_kv_row`), so a layered `memcpy` into a staging buffer + one write is straightforward. + +**Win:** ~157× fewer fwrite calls per token. Even with stdio coalescing, the userspace loop overhead is real at scale. + +### Opportunity 2 — `/proc/meminfo` fopen storm (MEDIUM ROI, trivial) + +**Problem:** `rss_gb()` opens, reads, closes `/proc/meminfo` every ~16 tokens on Linux. Each is ~3 syscalls + path resolution. + +**Fix:** Either (a) cache the value for N tokens (e.g. re-read at most once per second), or (b) keep the fd open and `rewind`+`fgets`. Trivial change. + +### Opportunity 3 — Expert prefetch on Windows (MEDIUM ROI, bounded) + +**Problem:** The `COLI_MMAP=1` path (which gives zero-copy expert access + free OS-cache re-faults) is **Linux/macOS/FreeBSD-only** — `glm.c:1301` guards it. On Windows, experts always go through the `pread` path, and `expert_prefetch` issues `posix_fadvise(WILLNEED)` which is a no-op shim on Windows (`compat.h`). + +**Fix:** On Windows, implement the prefetch via `PrefetchVirtualMemory` (the Win32 analog of `MADV_WILLNEED`) on an mmap'd region, or via an async `ReadFile`+`OVERLAPPED` into a scratch buffer. This brings the Windows build closer to parity with the Linux mmap+prefetch story. + +**Scope:** This is the largest of the four — it touches the Windows I/O path. Worth doing if Windows perf is a goal; skip if Linux is the target. + +### Opportunity 4 — KV-cache: keep handle open (LOW-MEDIUM ROI, LOW risk) + +**Problem:** `kv_disk_append` does `fopen`+...+`fclose` every turn. Handle creation is ~5-15µs of pure overhead (worse on Windows). + +**Fix:** Open the KV file once (lazily on first append), keep the `FILE*` for the engine lifetime, just `fseek`+write each turn. Close on shutdown. Pair with Opportunity 1 for the write batching. + +--- + +## Non-opportunities (recording so we don't re-investigate) + +- **O_DIRECT for experts**: already correctly used (`st.h:83`, `DIRECT=1`). For an LRU+refetch pattern the page cache is your friend, but the engine offers both paths (O_DIRECT pread default + optional mmap) and the O_DIRECT coalesced read is already one syscall for ~19MB. Don't change this. +- **Single-file layout**: the engine already uses safetensors shards with kept-open fds + offset-indexed tensors (`st.h`). No per-expert open()/close() waste. Don't change this. +- **PagedAttention**: solves concurrency fragmentation this engine doesn't have (≤16 slots). Not applicable. + +--- + +## Next steps + +The highest-ROI, lowest-risk starting point is **Opportunity 1 (KV write batching) + Opportunity 4 (keep handle open)** — they're in the same function, both low-risk, and together they eliminate the per-turn open/close overhead and the per-token fwrite storm. Opportunity 2 is a trivial 5-minute fix we can bundle in. + +Opportunity 3 (Windows prefetch) is the biggest single win but also the largest scope — separate effort, gated on whether Windows perf is a priority. + +## Sources + +- [justine.lol/mmap — Edge AI Just Got Faster](https://justine.lol/mmap/) +- [llama.cpp discussion #18758 — Mmap faster than direct I/O for MoE](https://github.com/ggml-org/llama.cpp/discussions/18758) +- [llama.cpp issue #20757 — Two-tier GPU+RAM expert cache](https://github.com/ggml-org/llama.cpp/issues/20757) +- [FAST '26 — Programmable Page Cache for LLM loading](https://www.usenix.org/system/files/fast26-liu-yubo.pdf) +- [FAST '26 — SolidAttention: SSD-based serving](https://www.usenix.org/system/files/fast26-zheng.pdf) +- [HOBBIT — Mixed precision expert offloading](https://arxiv.org/html/2411.01433v2) +- [posix_fadvise(2) — man7.org](https://man7.org/linux/man-pages/man2/posix_fadvise.2.html) +- [madvise(2) — man7.org](https://man7.org/linux/man-pages/man2/madvise.2.html) +- [Microsoft Learn — File Buffering (FILE_FLAG_NO_BUFFERING)](https://learn.microsoft.com/en-us/windows/win32/fileio/file-buffering) +- [Microsoft Learn — PrefetchVirtualMemory](https://learn.microsoft.com/en-us/windows/win32/api/memoryapi/nf-memoryapi-prefetchvirtualmemory) +- [What makes system calls expensive — codingconfessions.com](https://blog.codingconfessions.com/p/what-makes-system-calls-expensive) +- [Syscall overhead — Stack Overflow](https://stackoverflow.com/questions/8247331/syscall-overhead) + +--- + +# Windows Implementation — branch `windows-optimizations` (2026-07-15) + +## What landed (pread path, validated) + +Two changes, both on the `pread` expert-load path (no mmap). Measured against the +existing `bench_budget*.txt` baselines (GLM-5.2 744B int4, 32 GB RAM, Core Ultra 9 +185H, DRAFT=0, 32-token decode): + +### 1. `compat_fadvise` WILLNEED cache-warmer (`c/compat.h`) + +Replaced the Windows `posix_fadvise` no-op (was a `do{}while(0)` macro) with a real +readahead: an overlapped `ReadFile` into a throwaway scratch buffer that populates the +standby page cache, so the later synchronous `pread` faults from RAM not disk. Mirrors +the macOS `F_RDADVISE` shim (`compat.h:28-37`). DONTNEED stays a no-op (matches macOS; +Windows standby-list trimming self-regulates under pressure). + +This re-arms the existing `expert_prefetch` → `st_prefetch` → `posix_fadvise(WILLNEED)` +chain on Windows: the next-block readahead in `moe()` and the PILOT cross-layer prefetch +hints now actually warm the cache instead of being silently discarded. + +**Measured effect (budget=4, PIPE on):** hit rate 16.4% → 27.6%. + +### 2. PIPE default ON for Windows (`c/glm.c`) + +Flipped the async expert-load thread pool from default OFF to default ON on Windows +(`getenv("PIPE")?:1` under `_WIN32`, unchanged `:0` elsewhere). PIPE dispatches expert +`pread` loads onto worker threads so they overlap the expert matmul on the forward-pass +thread, instead of the blocking serial load-then-compute path. `PIPE=0` opts back out. + +**Measured effect (budget=4):** expert-disk 65.9s → 54.3s (−18%), reaching **1.70 s/tok** +(under the 2 s/tok target; budget=4 baseline was 2.06 s/tok). + +### Results table (DRAFT=0, 32-token decode, pread path) + +| config | expert-disk | s/tok | hit% | tok/s | +|---|---|---|---|---| +| budget=4, no PIPE (existing baseline) | 65.9s | 2.06 | 16.4% | 0.33 | +| **budget=4 + PIPE (this PR)** | **54.3s** | **1.70** | **27.6%** | **0.34** | +| budget=6 + PIPE | 77.1s | 2.41 | 21.8% | 0.27 | + +budget=4 + PIPE meets the ≤2 s/tok target. budget=6 (more experts/layer, higher quality) +misses it at 2.41 s/tok — the speed/quality tradeoff. + +## What was tried and abandoned: Windows mmap (`COLI_MMAP` on `_WIN32`) + +The original plan (informed by llama.cpp #18758: "mmap is ≥10× faster than O_DIRECT for +MoE") was to port the mmap expert path to Windows via `CreateFileMapping`/`MapViewOfFile`. +This was implemented and tested at length. **It was a measured regression and was reverted.** + +### The attempt + +Added a `_WIN32` branch to `map_of_fd` (`glm.c`) mapping each shard file read-only and +resolving experts as views into the mapping, mirroring the POSIX path. Also added +`PrefetchVirtualMemory` readahead and a `VirtualUnlock` eviction mechanism (the Windows +`posix_fadvise(DONTNEED)` analog — see SO#1880714; validated standalone to demote pages +to the standby list with a 2.3× faster re-fault). + +### Why it regressed + +**mmap'd expert pages bloat the process working set on Windows, which collapses the +expert cache.** This is a fundamental Windows-vs-Linux difference: + +- On Linux, `mmap(MAP_SHARED)` file pages live in the kernel page cache (`buff/cache`), + separate from `MemAvailable`, so the cache budget isn't fooled. +- On Windows, touched `MapViewOfFile` pages count against `ullAvailPhys` (what + `compat_meminfo` reads for the budget). The CPU matmul touches every weight byte, + faulting ~12 GB into the working set. `cap_for_ram()` then sees ~no free RAM and + collapses the LRU cache cap. + +Measured (budget=0, DRAFT=0, apples-to-apples): + +| config | RAM_GB detected | cache cap | hit% | expert-disk | RSS | +|---|---|---|---|---|---| +| baseline (pread) | 21.4 | 1 | 9.3% | 133s | 15.0 GB | +| mmap, no eviction | **8.0** | 1 | **2.2%** | **240s** | **27.2 GB** | +| mmap + VirtualUnlock | 24.9 | 2 | 11.8% | 83s | 18.1 GB | +| mmap + reserve reductions | 24.6 | 4 | 21.8% | 80s | 20.1 GB | + +The `VirtualUnlock` eviction recovered the regression (240s→83s), and dropping the +Linux-specific page-cache/slab reserves under mmap got it to parity with pread. But it +never clearly *beat* the simpler pread+PIPE path, and it added substantial complexity +(per-slot eviction tracking, reserve conditionals, `VirtualUnlock` on every slot recycle). +**The engine already moved off mmap to pread for this exact RSS bug** (`st.h:3-6`), and +the Windows port re-confirmed that decision. + +### What else didn't work + +- **Batched `PrefetchVirtualMemory`** for the mmap path: tested as a single batched + readahead of all 64 missed experts' pages before the matmul. **Blocked instead of + prefetching async** on this SSD — inflated `t_edisk` (80s→102s). Consistent with + microsoft/Windows-Dev-Performance#108 ("PrefetchVirtualMemory does not prefetch"). + Reverted. +- **True I/O/compute overlap on the CPU path**: the Metal path has this ("submit + resident experts to GPU before loading misses"), but the CPU path loads-then-computes + serially. `PrefetchVirtualMemory` was the attempt to add it for mmap and failed. The + pread path gets overlap via PIPE (which works), not via mmap prefetch. + +### Conclusion + +For this engine on Windows at this RAM budget (~32 GB, 370 GB model), **pread + PIPE + +compat_fadvise** is the right path. mmap remains valuable on Linux/macOS (where the page +cache doesn't inflate process RSS) but is not viable on Windows without a fundamentally +different cache-budget model that excludes mapped-file pages — left as future work. + +## Sources added + +- [SO#1880714 — VirtualUnlock releases mapped pages to standby list](https://stackoverflow.com/questions/1880714/createfilemapping-mapviewoffile-how-to-avoid-holding-up-the-system-memory) +- [Alois Kraus — The Mysterious Lost Memory (modified/standby list)](https://aloiskraus.wordpress.com/2017/02/26/the-mysterious-lost-memory-which-belongs-to-no-process/) +- [microsoft/Windows-Dev-Performance#108 — PrefetchVirtualMemory inconsistency](https://github.com/microsoft/Windows-Dev-Performance/issues/108) +- [llama.cpp #18758 — mmap faster than O_DIRECT for MoE (Linux)](https://github.com/ggml-org/llama.cpp/discussions/18758) +- [HN#35426679 — Why MMAP in llama.cpp hides true memory usage](https://news.ycombinator.com/item?id=35426679) diff --git a/issue_grouped_quant.md b/issue_grouped_quant.md new file mode 100644 index 0000000..fb46ce4 --- /dev/null +++ b/issue_grouped_quant.md @@ -0,0 +1,185 @@ +## TL;DR + +The int4 model produces **fluent-but-incoherent output** — grammatical English that is completely off-topic (SEO spam, homework-help boilerplate) instead of reasoned responses. The root cause is **per-row quantization scales**: one F32 scale per output row (e.g. 2048 scales for a 2048×6144 matrix), which is 48x coarser than the FP8 source's 128×128 block scales. This destroys fine-grained weight information that handles reasoning and instruction-following while preserving the large-magnitude weights that handle grammar and vocabulary fluency. + +Branch: `experiment/grouped-quant` (based on latest `dev` at `62419af`) + +--- + +## The problem — demonstrated + +Prompt: *"Explain the concept of recursion in programming. Provide a simple example."* + +**Baseline (no budget, no changes):** +> "Hire Some To Take Object-Oriented Programming Assignment" + +**Budget=4:** +> "The world is a dangerous place, not so much because of the small percentage of people who are doing to do the little of people who are doing to" + +**Budget=6:** +> "Elite Custom Essays" + +**Budget=12:** +> "Sololearn: Learn to code for FREE! +1 # Explain A concept of recursion..." + +This is not random gibberish — it's fluent English that is completely off-topic. This is the signature of **activations corrupted by coarse quantization**: the model retains language fluency (large weights survive) but loses instruction-following and reasoning (fine-grained weights are crushed to zero). + +## Root cause: per-row int4 scales + +### How the current converter works + +`convert_fp8_to_int4.py` (line 39-52) quantizes with **one scale per output row**: + +```python +amax = np.abs(w).max(axis=1, keepdims=True) # one max per ROW +s = np.maximum(amax / qmax, 1e-8) # one scale per ROW +q = np.clip(np.rint(w / s), -8, qmax) # quantize entire row with that scale +``` + +For a 2048×6144 expert weight matrix, that's **2048 scales** — one per row, each covering 6144 elements. + +### Why this destroys quality + +Consider a row of 6144 values where most are small (magnitude ~0.01) but a few are large (~0.5). The per-row scale is `0.5/7 = 0.071`. The small values become `0.01/0.071 = 0.14`, which rounds to **zero**. All fine-grained information in those small weights is lost. + +This matters enormously for MoE models: each token activates only 8 of 256 experts. A poorly-quantized expert pollutes every token that routes to it, and there's no averaging from the other 248 experts (they're simply off). The error is **concentrated**, not diluted. + +### What the FP8 source does right + +The FP8 checkpoint uses **128×128 block scales** — the weight matrix is divided into 128-element chunks, each with its own scale. Small values in one chunk don't get crushed by large values in another. For a 2048×6144 matrix: `2048 × 48 = 98,304` scales — 48x more granularity. + +The converter already dequants FP8 to f32 correctly (lines 196-201), preserving the block-scale information. But then it **throws it all away** by collapsing to a single per-row scale during int4 requantization. + +### GLM-5.2 was QAT-trained for int4 + +The GLM-5 paper (arXiv 2602.15763, §2.4.3) states: *"To provide better accuracy at low-precision, we apply INT4 QAT in the SFT stage."* This means the model is **designed** to work at int4 — but only if the quantization is fine-grained enough to match what the model saw during training. The FP8 checkpoint's 128×128 block scales are the granularity the model expects. Per-row scaling is 48x coarser. + +### Community confirmation + +Every other project getting coherent GLM-5.2 int4 output uses calibrated or fine-grained quantization: +- **ubergarm/GLM-5.1-GGUF** — imatrix-calibrated with expert-specific patches +- **Unsloth/GLM-5-GGUF** — dynamic UD-Q4 quants +- **QuantTrio/GLM-5.2-Int4-Int8Mix** — mixed int4/int8 with channel-wise scales +- **llama.cpp** — Q4_K with block-level group scales (typically 32 or 64 elements) + +None use naive per-row RTN int4 for production inference. + +## The fix: group-scaled int4 (fmt=4) + +Add a new quantization format with **one scale per 128 elements** along the input dimension, matching the FP8 source's natural granularity. + +### What changed + +**`c/glm.c`** — 122 lines added: + +1. **QT struct** (line 98): Added `int gs` field (group size, 0=per-row for backward compat, 128=grouped). + +2. **Format detection** (`qt_from_disk`, line 1064): Auto-detects fmt=4 by checking the `.qs` scale array size. If it has `O * ceil(I/128)` elements instead of `O`, it's grouped. **Old per-row models (fmt=2) work unchanged.** + +3. **New kernel** (`matmul_i4_grouped`, line 379): Same AVX2 nibble unpacking as `matmul_i4`, but the accumulator resets at each 128-element group boundary: `dot(x[grp], w[grp]) * scale[grp]`. The scale changes every 8 vector iterations (128/16=8). + +4. **Dispatch** (`matmul_qt_ex`, line 813): Routes to `matmul_i4_grouped` when `fmt==4`. Always uses exact kernels (no IDOT approximation — the whole point is quality). + +5. **Expert loading** (`expert_load`, lines 1418 and 1538): Both the mmap and slab+pread paths detect fmt=4 from scale array size and set `gs=128`. + +6. **`qt_bytes`** (line 104): Reports correct memory for fmt=4. + +**`c/tools/convert_fp8_to_int4.py`** — `quant_int4_grouped()` + `--group-size` arg: + +```python +def quant_int4_grouped(w, bits, gs=128): + O, I = w.shape + ngroups = (I + gs - 1) // gs + wpad = np.zeros((O, ngroups * gs), np.float32) + wpad[:, :I] = w + wr = wpad.reshape(O, ngroups, gs) # [O, ngroups, gs] + amax = np.abs(wr).max(axis=2, keepdims=True) # one max per GROUP + s = np.maximum(amax / qmax, 1e-8) + q = np.clip(np.rint(wr / s), -8, qmax) + # ... same nibble packing as quant_int4 ... + return packed_nibbles, s.reshape(-1) # [O * ngroups] scales +``` + +Same packed-nibble format as existing int4 — only the scale array is larger. + +### Verification + +**Converter round-trip test** (weights with varying group magnitudes): + +| Method | Mean relative error | Max abs error | +|--------|-------------------|---------------| +| Per-row int4 (current) | 0.2056 | 0.0127 | +| Grouped int4 (gs=128) | **0.1278** | 0.0127 | +| **Improvement** | **1.6x lower** | — | + +**Engine kernel test** (AVX2 `matmul_i4_grouped` vs f32 reference): + +| Output | Reference | Kernel | Error | +|--------|-----------|--------|-------| +| o=0 | -0.126368 | -0.126368 | 2.98e-08 | +| o=1 | 0.075913 | 0.075913 | 1.49e-08 | +| o=2 | 0.077136 | 0.077136 | 7.45e-09 | +| ... | ... | ... | ... | + +Max error: **2.98e-08** — matches f32 reference to within float32 epsilon. **PASS.** + +### Cost + +| Metric | Per-row (current) | Grouped (new) | +|--------|-------------------|---------------| +| Expert weight size | ~18.9 MB | ~20.1 MB (+6%) | +| Scale array per matrix | O × 4 bytes | O × ceil(I/128) × 4 bytes | +| Total model size | ~370 GB | ~390 GB (+5%) | +| Disk I/O per miss | ~19 MB | ~20 MB (+6%) | +| Kernel speed | One scale multiply per row | One scale multiply per 128 elements | + +The 6% I/O increase is negligible compared to the quality gain. The grouped kernel has slightly more scale-lookup overhead but remains within AVX2 throughput — the bottleneck is disk I/O, not matmul. + +### Backward compatibility + +- Old per-row models (fmt=2) **work unchanged** — format auto-detected from scale array size +- All existing features (PILOT, EXPERT_BUDGET, CUDA, MTP, PIPE) work identically — they don't touch the dequant path +- The fused gate+up pair path (`matmul_i4_pair`) falls back to separate `matmul_qt` calls for fmt=4 — minor perf cost, correctness preserved +- `--group-size 0` produces per-row output (backward compat for the converter) + +## How to reproduce + +### Convert with group scales + +```bash +python tools/convert_fp8_to_int4.py \ + --indir /path/to/GLM-5.2-FP8 \ + --outdir /path/to/glm52_i4_grouped \ + --ebits 4 --io-bits 8 --group-size 128 +``` + +### Test quality + +```bash +# Old model (per-row): +SNAP=/path/to/glm52_i4 PROMPT="Explain recursion in programming." NGEN=32 ./glm 64 + +# New model (grouped): +SNAP=/path/to/glm52_i4_grouped PROMPT="Explain recursion in programming." NGEN=32 ./glm 64 +``` + +Compare output text coherence. + +## What this won't fix + +- **Cold cache speed** — still 0.18-0.36 tok/s on 24 GB RAM. Quality and speed are independent axes. +- **All quantization error** — int4 is still 4-bit. Some degradation vs FP8 will remain. But it should be coherent degradation (slightly wrong word choices) rather than total reasoning failure (unrelated topics). +- **The IDOT +12% perplexity** — the approximate activation kernel (`IDOT=1`, default ON) still applies to non-grouped tensors (attention projections are already protected). Grouped int4 uses exact kernels by design. + +## Prior art + +- **GLM-5 paper** (arXiv 2602.15763, §2.4.3): "We apply INT4 QAT in the SFT stage" — the model is trained for int4, but assumes fine-grained quantization. +- **MxMoE** (ICML 2025): MoE experts exhibit divergent quantization sensitivity; uniform quant across all experts is suboptimal. +- **MoEQuant** (OpenReview): Naive int4 on MoE loses meaningful accuracy; calibrated framework needed. +- **Automated Fine-Grained MoE Quantization** (ACL 2025): Layer/expert-wise sensitivity variation; group-size scaling is the baseline improvement. +- **ubergarm/GLM-5.1-GGUF**: imatrix-calibrated quants with expert-specific patches — explicitly patches `quantize_row_q4_0_ref()` for routed experts. +- **llama.cpp Q4_K**: Uses block-level group scales (32 or 64 elements) as the standard int4 format. + +## Conversion status + +Re-converting from the FP8 source (`zai-org/GLM-5.2-FP8`) with `--group-size 128`. Downloading via ModelScope to avoid HuggingFace per-stream throttling. Will report quality A/B results once the conversion is complete. diff --git a/web/src/App.tsx b/web/src/App.tsx index 57e5f98..d4e0e41 100644 --- a/web/src/App.tsx +++ b/web/src/App.tsx @@ -34,7 +34,11 @@ import { Brain } from "./Brain" import { persistPublicSettings, stored } from "@/lib/storage" import { cn } from "@/lib/utils" -const message = (role: ChatMessage["role"], content: string): ChatMessage => ({ id: crypto.randomUUID(), role, content }) +const message = (role: ChatMessage["role"], content: string): ChatMessage => { + let id: string + try { id = crypto.randomUUID() } catch { id = 'xxxxxxxx-xxxx-4xxx-yxxx-xxxxxxxxxxxx'.replace(/[xy]/g, c => { const r = Math.random() * 16 | 0; return (c === 'x' ? r : (r & 0x3 | 0x8)).toString(16) }) } + return { id, role, content } +} export default function App() { // When the page is served by the engine itself (coli web), same-origin is the