Merge remote-tracking branch 'origin/dev' into pr259

# Conflicts:
#	README.md

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=<n>` 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
 

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%

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."""
+    """legge fino alla sentinella; on_bytes riceve i chunk della risposta. Poi legge la riga STAT.
-    pend=b""
+    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:
+        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}")

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

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")

c/glm.c

@@ -35,6 +35,7 @@
 #include <sys/resource.h>
 #include <sys/mman.h>                             /* mlock: inchioda le pagine in RAM / wire pages into RAM */
 #include <sys/stat.h>                             /* fstat per mmap degli shard (COLI_MMAP) */
+#include <signal.h>                               /* 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;j<nu;j++) wsum[j]=0;
+        for(int s=0;s<S;s++) for(int kk=0;kk<keff[s];kk++){
+            int e=idxs[(int64_t)s*K+kk];
+            for(int j=0;j<nu;j++) if(uniq[j]==e){ wsum[j]+=ws[(int64_t)s*K+kk]; break; }
+        }
+        /* residency pre-scan: which experts are already in pin or ecache (hits)? */
+        unsigned char *is_hit=calloc(nu,1); int nhits=0;
+        for(int j=0;j<nu;j++){ int eid=uniq[j];
+            int found=0;
+            ESlot *P=m->pin[layer];
+            for(int z=0;z<m->npin[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;z<nn;z++) if(Sl[z].eid==eid){ found=1; break; } }
+            if(found){ is_hit[j]=1; nhits++; }
+        }
+        /* budget for misses = total budget - hits already kept (min 0) */
+        int miss_budget = g_expert_budget - nhits; if(miss_budget<0) miss_budget=0;
+        /* mark which unique experts to keep (1) or drop (0): keep all hits, fill rest
+         * with top-weight misses up to miss_budget */
+        unsigned char *keep=calloc(nu,1); int nkeep=0;
+        for(int j=0;j<nu;j++) if(is_hit[j]){ keep[j]=1; nkeep++; }
+        for(int rank=0;rank<miss_budget;rank++){
+            int best=-1; float bv=-1e30f;
+            for(int j=0;j<nu;j++) if(!keep[j] && wsum[j]>bv){ 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;j<nu;j++) if(keep[j]) seen[uniq[j]]=1;
+        int dropped=nu-nkeep; g_budget_dropped+=dropped;
+        /* remove dropped experts from each position's routing list */
+        for(int s=0;s<S;s++){
+            int w=0;
+            for(int kk=0;kk<keff[s];kk++){
+                int e=idxs[(int64_t)s*K+kk];
+                if(seen[e]){ idxs[(int64_t)s*K+w]=e; ws[(int64_t)s*K+w]=ws[(int64_t)s*K+kk]; w++; }
+            }
+            if(w<keff[s]){
+                keff[s]=w;
+                /* renormalize remaining weights per position */
+                if(c->norm_topk && w>0){
+                    float sm=0; for(int kk=0;kk<w;kk++) sm+=ws[(int64_t)s*K+kk]; sm+=1e-20f;
+                    for(int kk=0;kk<w;kk++) ws[(int64_t)s*K+kk]/=sm;
+                    for(int kk=0;kk<w;kk++) ws[(int64_t)s*K+kk]*=c->routed_scale;
+                }
+            }
+        }
+        /* compact uniq[] to kept experts only */
+        int nu2=0;
+        for(int j=0;j<nu;j++) if(keep[j]) uniq[nu2++]=uniq[j];
+        nu=nu2;
+        free(wsum); free(is_hit); free(keep);
+    }
     /* ---- FASE C/D: risolvi (pin/cache/disco) e calcola, a blocchi di 64 unici ---- */
     float *xg=falloc((int64_t)S*D), *gg=falloc((int64_t)S*I), *uu=falloc((int64_t)S*I), *hh=falloc((int64_t)S*D);
     int *rows=malloc(S*sizeof(int)); float *rw=malloc(S*sizeof(float));
@@ -3795,12 +3899,34 @@ static void stops_arm(const Cfg *c, int tok_eos){
  * all: storia token (capacita' >= 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<n_new && !done){
+    while(emitted<n_new && !done && !g_intr){   /* g_intr: stessa uscita del tetto n_new */
         int next=pick_tok(logit,V,carry_ban); carry_ban=-1; free(logit); logit=NULL;
         if((eos>=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;i<c->n_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(!kv_disk_open(m)) return;
-    if(!f){ f=fopen(k->disk_path,"wb"); if(!f) return;
+    FILE *f=k->disk_fp;
-        int32_t h[8]; kv_hdr(m,h,0); fwrite(KV_MAGIC,1,8,f); fwrite(h,4,8,f); }
+    int64_t rec = kv_rec_bytes(m);
-    int64_t rec = 4 + (int64_t)c->n_layers*(c->kv_lora+c->qk_rope)*4;
+    /* grow the contiguous staging buffer if the record is larger (#1 batching) */
-    if(m->has_dsa) for(int i=0;i<c->n_layers;i++) if(m->Ic[i]) rec+=(int64_t)c->index_hd*4;
+    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;p<len;p++){
-        int32_t tk=hist[p]; fwrite(&tk,4,1,f);
+        uint8_t *b=k->disk_buf;            /* pack token + every layer into one record */
+        *(int32_t*)b = hist[p]; b+=4;
         for(int i=0;i<c->n_layers;i++){
-            fwrite(m->Lc[i]+(int64_t)p*c->kv_lora, 4, c->kv_lora, f);
+            memcpy(b, m->Lc[i]+(int64_t)p*c->kv_lora, (size_t)c->kv_lora*4); b+=c->kv_lora*4;
-            fwrite(m->Rc[i]+(int64_t)p*c->qk_rope, 4, c->qk_rope, f);
+            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;i<c->n_layers;i++) if(m->Ic[i])
+        if(m->has_dsa) for(int i=0;i<c->n_layers;i++) if(m->Ic[i]){
-            fwrite(m->Ic[i]+(int64_t)p*c->index_hd, 4, c->index_hd, f);
+            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;i<NR;i++){ free(k->Lc[i]); free(k->Rc[i]); }
     if(k->Ic) for(int i=0;i<m->c.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;i<nctx;i++) if(req[i].active) mux_done(m,&ctx[i],&req[i]);
+        }
         int active=0; for(int i=0;i<nctx;i++) active+=req[i].active;
         /* Poll stdin for available input without blocking. On POSIX this is
          * select(); on Windows, select() on a pipe handle routes to winsock
@@ -4563,9 +4735,11 @@ static void run_serve(Model *m, const char *snap){
     #define len   (sc->len)
     #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;

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:
+    if isinstance(maximum, bool) or not isinstance(maximum, int) or maximum < 1:
-        raise APIError(400, f"`{maximum_param}` must be an integer between 1 and {limit}.", maximum_param)
+        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")

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");

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):

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())
@@ -212,8 +225,16 @@ def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits,
             w = dequant(f, name, keys)
             if kind == "f32":
                 out_dict[name] = w.astype(np.float32)
+            else:
+                # 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):

c/tools/expert_atlas.py

@@ -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()

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

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()

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)

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))

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."""
+l'indexer sparso.
-import json, torch
+
+--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 ""))

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 _rot_quant(w, bits, group, e8)
-    return _quant_last_dim(w, bits, group)  # nn.Linear [out, in] -- input already last
+    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<N>][-rot][-nohead])")
+        raise SystemExit(f"bad scheme '{name}' (expected fp16 | int{{2,3,4,8}}[-g<N>][-e8|-e8u][-rot][-nohead])")
-    return int(m.group(1)), int(m.group(2) or 0), bool(m.group(3)), bool(m.group(4))
+    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

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=<model_dir> MTP=0 PROMPT="Explain recursion in programming." NGEN=32 ./glm 64
+
+# With budget:
+SNAP=<model_dir> MTP=0 EXPERT_BUDGET=4 PROMPT="Explain recursion in programming." NGEN=32 ./glm 64
+
+# With MTP (where budget saves the most):
+SNAP=<model_dir> 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.

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)

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.

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